From: Gilles Perrot Date: Fri, 11 Mar 2011 15:10:56 +0000 (+0100) Subject: Fin de test multisnake sur iter 1 X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/snake_gpu.git/commitdiff_plain/48b1bac747f398161b53a67ad80c4596f531c88a?hp=af4b787ce73a80f23e9e2b1ef9ac52660e8ab754 Fin de test multisnake sur iter 1 Pas performant --- diff --git a/exec/SNAKE2D b/exec/SNAKE2D index 1d1de07..40c6602 100755 Binary files a/exec/SNAKE2D and b/exec/SNAKE2D differ diff --git a/lib/lib_contour.o b/lib/lib_contour.o index 26fc29d..e197c41 100644 Binary files a/lib/lib_contour.o and b/lib/lib_contour.o differ diff --git a/lib/lib_gpu.o b/lib/lib_gpu.o index 1d978d7..dc40914 100644 Binary files a/lib/lib_gpu.o and b/lib/lib_gpu.o differ diff --git a/lib/lib_snake_2_gpu.o b/lib/lib_snake_2_gpu.o index 4a3a81f..d74fedd 100644 Binary files a/lib/lib_snake_2_gpu.o and b/lib/lib_snake_2_gpu.o differ diff --git a/lib/lib_snake_common.o b/lib/lib_snake_common.o index a7abed2..4abaa10 100644 Binary files a/lib/lib_snake_common.o and b/lib/lib_snake_common.o differ diff --git a/lib/lib_test_gpu.o b/lib/lib_test_gpu.o index 5a6e207..676e14d 100644 Binary files a/lib/lib_test_gpu.o and b/lib/lib_test_gpu.o differ diff --git a/lib/snake2D_gpu.o b/lib/snake2D_gpu.o index 7be4b88..3b6bf11 100644 Binary files a/lib/snake2D_gpu.o and b/lib/snake2D_gpu.o differ diff --git a/src/lib_gpu.cu b/src/lib_gpu.cu index 1bbd61d..5a90c1c 100644 --- a/src/lib_gpu.cu +++ b/src/lib_gpu.cu @@ -153,9 +153,9 @@ void cuda_init_img_cumul(unsigned short ** img_in, int H, int L, int nb_nodes, calcul_stats_image<<<1, 1>>>( *d_img_x, *d_img_x2, H, L, (uint64*)*d_stats_snake); - cudaThreadSynchronize() ; - toc(chrono, "\tTemps GPU"); - if(DEBUG_IMG_CUMUL) + cudaThreadSynchronize() ; + toc(chrono, "\tTemps GPU"); + if(DEBUG_IMG_CUMUL) { //allocation memoire CPU @@ -231,33 +231,52 @@ void cuda_init_img_cumul(unsigned short ** img_in, int H, int L, int nb_nodes, * generation snake en mem GPU */ int dist = 140 ; - - tic(&chrono, NULL); - int MAX_DIAGOS = 1024*65536 ; - int ret, Q = 120 ; - int * d_n_diagos, h_n_diagos; - uint4 * d_diagos_snake ; - uint4 * h_diagos_snake = new uint4[MAX_DIAGOS]; - - - cudaMalloc( (void**) &d_n_diagos, sizeof(int)) ; - cudaMalloc( (void**) &d_diagos_snake, MAX_DIAGOS*sizeof(uint4)) ; - - genere_diagos_rectangle<<<1,1>>>(d_diagos_snake, H,L,Q, d_n_diagos); + /* Test de determination du snake rectangle initial optimal*/ + int div = 100;//nb de divisions de l'image : cela définit le pas. La valeur max découle du nb max de threads possible ds une grille + int Nperm = div*div*bs;//nb total de rectangles a tester. La distribution est ainsi irrégulière, mais plus simple. + double best_crit ; + int ind_best_crit ; + + t_rectangle_snake * d_all_crit, d_best_crit;//tableaux pour les résultats des différents rectangles / le meilleur + t_rectangle_snake * h_all_crit = new t_rectangle_snake[Nperm];//correspondant CPU - cudaMemcpy( &h_n_diagos, d_n_diagos, sizeof(int), cudaMemcpyDeviceToHost) ; - ret = cudaMemcpy( h_diagos_snake, d_diagos_snake, MAX_DIAGOS*sizeof(uint4), cudaMemcpyDeviceToHost) ; + //allocations + cudaMalloc((void**) &d_all_crit, Nperm*sizeof(t_rectangle_snake)); + cudaMalloc((void**) &d_best_crit, sizeof(t_rectangle_snake)); + + tic(&chrono, NULL); - toc(chrono, "\tCalcul diagos"); + //execution kernel + dim3 grid = dim3(H/div, L/div, 1); + calcul_contribs_snake4<<>>(*d_snake, *d_img_x, *d_img_x2, H, L, *d_stats_snake, d_all_crit) ; + cudaThreadSynchronize(); + toc(chrono, "\nCALCULS RECTANGLES"); - printf("COPY : %s, DIAGOS = %d / %d\n", (ret==0)?"OK":"ERREUR", h_n_diagos,MAX_DIAGOS); - for (int d=0; d<200;d++){ - printf("| (%d,%d)-(%d,%d) ", h_diagos_snake[d].x, h_diagos_snake[d].y, h_diagos_snake[d].z, h_diagos_snake[d].w ); + //recup data rectangles + int ret; + ret = cudaMemcpy( h_all_crit, d_all_crit, Nperm*sizeof(t_rectangle_snake), cudaMemcpyDeviceToHost) ; + printf("COPIE DATA = %s\n",(ret==0)?"OK":"ERR"); + + //optimum sur CPU + best_crit = h_all_crit[0].crit ; + ind_best_crit = 0 ; + for (int k=1; k<100; k++){ + if ((h_all_crit[k].crit > 0) && (h_all_crit[k].crit < best_crit)) { + best_crit = h_all_crit[k].crit ; + ind_best_crit = k ; + } + printf("%d -> ( %d, %d )--( %d, %d) CRITERE = %f\n", k, h_all_crit[k].bpi, h_all_crit[k].bpj, + h_all_crit[k].opi, h_all_crit[k].opj, h_all_crit[k].crit ); } + printf("BEST RECTANGLE/%d tests : %d -> ( %d, %d )--( %d, %d) CRITERE = %f\n", Nperm, ind_best_crit, h_all_crit[ind_best_crit].bpi, h_all_crit[ind_best_crit].bpj, + h_all_crit[ind_best_crit].opi, h_all_crit[ind_best_crit].opj, best_crit ); + exit(0); - genere_snake_rectangle_4nodes_gpu<<< 1, 1>>>(*d_snake, 140, H, L) ; + /*fin test snake rectangle initial optimal*/ + + //genere_snake_rectangle_4nodes_gpu<<< 1, 1>>>(*d_snake, 140, H, L) ; int nnodes = nb_nodes ; diff --git a/src/lib_kernel_snake_2_gpu.cu b/src/lib_kernel_snake_2_gpu.cu index 5293308..3796e65 100644 --- a/src/lib_kernel_snake_2_gpu.cu +++ b/src/lib_kernel_snake_2_gpu.cu @@ -468,46 +468,52 @@ __global__ void calcul_stats_snake(snake_node_gpu * d_snake, int nnodes, int64 d_stats_snake[3], d_stats_snake[4], d_stats_snake[5]); } +// kernel d'init rectangulaire au niveau snake +// un block par point de base et un point opposé par thread du bloc -__global__ void calcul_contribs_snake4(t_cumul_x * cumul_x, t_cumul_x2 * cumul_x2, int h, int l, tcontribs * gcontribs, - uint64 SUM_1, uint64 SUM_X, uint64 SUM_X2) +__global__ void calcul_contribs_snake4(snake_node_gpu * snake, t_cumul_x * cumul_x, t_cumul_x2 * cumul_x2, int h, int l, + int64 * sums, t_rectangle_snake * critere) { // nb de diagonales testees par bloc (ie. par point de base NO) - int blockSize = blockDim.x ; + int blockSize = blockDim.x ; // indice du second point de chaque diagonale (=Opposite Point, = point SE) int OPib = threadIdx.x ; // coordonnees de chaque point de base (NO) int BPi = blockIdx.x ; int BPj = blockIdx.y ; //coordonnees de chaque Opposite Point (SE) - int OPi = OPib / (l - BPj) ; - int OPj = OPib - (l - BPj)*OPi ; + double incThread = ((h-BPi)*(l-BPj) + blockSize -1)/(double)blockSize ; + int OPi = (double)(OPib*incThread)/(l - BPj) ; + int OPj = OPib*incThread - OPi*(l-BPj) ; OPi += BPi ; OPj += BPj ; + if (OPi >= h) OPi = h-1 ; + if (OPj >= l) OPj = l-1 ; //indices des pixels dans les images cumulees int posG, posD; //contrib 1 du snake int C1 = (OPi - BPi)*(OPj - BPj) ; - - + //pour stocker contribs de chaque snake d'un block - //TODO on peut utiliser une structure restreinte (sans le c1) = gain d'espace + //TODO on peut utiliser une structure restreinte (sans le c1) extern __shared__ tcontribs scumuls[]; - + //calcul contribs du snake + scumuls[CFI(OPib)].cx = 0; + scumuls[CFI(OPib)].cx2 = 0; for (int k=BPi ; k < OPi ; k++) { posG = (BPi+k)*l + BPj ; posD = posG - BPj + OPj ; - scumuls[CFI(OPib)].cx += cumul_x[ posD ] - cumul_x[ posG ] ; - scumuls[CFI(OPib)].cx2 += cumul_x2[ posD ] - cumul_x2[ posG ]; + scumuls[CFI(OPib)].cx += (cumul_x[ posD ] - cumul_x[ posG ]) ; + scumuls[CFI(OPib)].cx2 += (cumul_x2[ posD ] - cumul_x2[ posG ]); } //calcul de critère pour chaque snake uint64 stat_sum_xe ; /* somme des xn region exterieure */ uint32 ne ; /* nombre de pixel region exterieure */ double sigi2, sige2; /* variance region interieure et exterieure */ - double criterion; + int index_crit; /* variance des valeurs des niveaux de gris a l'interieur du snake */ sigi2 = @@ -515,15 +521,34 @@ __global__ void calcul_contribs_snake4(t_cumul_x * cumul_x, t_cumul_x2 * cumul_x ((double)scumuls[CFI(OPib)].cx/(uint64)C1)*((double)scumuls[CFI(OPib)].cx/(uint64)C1) ; /* variance des valeurs des niveaux de gris a l'exterieur du snake */ - ne = SUM_1 - C1 ; - stat_sum_xe = SUM_X - scumuls[CFI(OPib)].cx ; + + ne = sums[3] - C1 ; + stat_sum_xe = sums[4] - scumuls[CFI(OPib)].cx ; sige2 = - ((double)SUM_X2-scumuls[CFI(OPib)].cx2)/(double)ne - + ((double)sums[5]-scumuls[CFI(OPib)].cx2)/(double)ne - ((double)stat_sum_xe/(uint64)ne)*((double)stat_sum_xe/(uint64)ne) ; - if ((sigi2 > 0)|(sige2 > 0)) - criterion = 0.5*((double)C1*log(sigi2) + (double)ne*log(sige2)) ; + index_crit = blockSize*(BPi*gridDim.y + BPj) + OPib ; + - //tri meilleur snake du bloc ( necessite de passer SUM_1, SUM_X et SUM_X2 ) + if ((sigi2 > 0)|(sige2 > 0)) + { + critere[ index_crit ].bpi = BPi; + critere[ index_crit ].bpj = BPj; + critere[ index_crit ].opi = OPi; + critere[ index_crit ].opj = OPj; + critere[ index_crit ].crit = 0.5*((double)C1*log(sigi2) + (double)ne*log(sige2)) ; + } + else + { + critere[ index_crit ].bpi = BPi; + critere[ index_crit ].bpj = BPj; + critere[ index_crit ].opi = OPi; + critere[ index_crit ].opj = OPj; + critere[ index_crit ].crit = -1 ; + } + // identification meilleur snake du bloc + // laissé au CPU pour test mais le principe de ce kernel n'est pas efficace. + } diff --git a/src/structures.h b/src/structures.h index b2a9554..3977661 100644 --- a/src/structures.h +++ b/src/structures.h @@ -141,5 +141,13 @@ struct snake_node_gpu int last_move ; /* dernier test deplacement accepte */ }; +/* pour mémoriser les coordonnées des diagonales des rectangles en multisnake*/ + struct t_rectangle_snake { + int bpi; + int bpj; + int opi; + int opj; + double crit; + }; #endif