X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/snake_gpu.git/blobdiff_plain/af4b787ce73a80f23e9e2b1ef9ac52660e8ab754..48b1bac747f398161b53a67ad80c4596f531c88a:/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 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.
+    
 }