X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/snake_gpu.git/blobdiff_plain/935bdd1c8b99ce5d70b6e4a53ca29f06353e8baa..af4b787ce73a80f23e9e2b1ef9ac52660e8ab754:/src/lib_kernel_snake_2_gpu.cu diff --git a/src/lib_kernel_snake_2_gpu.cu b/src/lib_kernel_snake_2_gpu.cu index 8d819a8..5293308 100644 --- a/src/lib_kernel_snake_2_gpu.cu +++ b/src/lib_kernel_snake_2_gpu.cu @@ -33,6 +33,28 @@ __global__ void genere_snake_rectangle_4nodes_gpu(snake_node_gpu * d_snake, int } } +__global__ void genere_diagos_rectangle(uint4 * d_diagos, int h, int l, int q, int * n_diagos){ + int inci = h/q; + int incj = l/q; + int iM,jM, iN, jN ; + int idxDiago = 0 ; + // boucle double pour les positions du point NO de la diagonale + for ( iM = 0; iM < q-1; iM++){ + for (jM = 0 ; jM < q-1 ; jM++){ + //boucle double pour les positions du point SE de la diagonale + for (iN = iM+1; iN < q; iN++){ + for (jN = jM+1; jN < q; jN++){ + d_diagos[idxDiago].x = iM*inci; + d_diagos[idxDiago].y = jM*incj; + d_diagos[idxDiago].z = iN*inci; + d_diagos[idxDiago].w = jN*incj; + idxDiago++; + } + } + } + } + *n_diagos = --idxDiago ; +} __global__ void genere_snake_rectangle_Nnodes_gpu(snake_node_gpu * d_snake, int dist_bords, int i_dim, int j_dim){ int nb_node_seg = 9 ; @@ -445,3 +467,63 @@ __global__ void calcul_stats_snake(snake_node_gpu * d_snake, int nnodes, int64 *vrais_min = codage_gl_gauss(s_stats_snake[0], s_stats_snake[1], s_stats_snake[2], d_stats_snake[3], d_stats_snake[4], d_stats_snake[5]); } + + +__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) +{ + // nb de diagonales testees par bloc (ie. par point de base NO) + 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 ; + OPi += BPi ; + OPj += BPj ; + //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 + extern __shared__ tcontribs scumuls[]; + + //calcul contribs du snake + 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 ]; + } + + //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; + + /* variance des valeurs des niveaux de gris a l'interieur du snake */ + sigi2 = + ((double)scumuls[CFI(OPib)].cx2/(double)C1) - + ((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 ; + sige2 = + ((double)SUM_X2-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)) ; + + //tri meilleur snake du bloc ( necessite de passer SUM_1, SUM_X et SUM_X2 ) + +}