Fin de test multisnake sur iter 1

[snake_gpu.git] / 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 52933088a07994cc9ffd8fe0b8f263108b860e37..3796e6569c7e152a73e4bc1f96885610a07f06a9 100644 (file)
--- 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]);
 }
 
                                                           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)
 {
   // 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)
   // 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 ;
   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) ; 
   //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
   //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[]; 
   extern __shared__ tcontribs scumuls[]; 

-   
+  

   //calcul contribs du snake
   //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 ;
   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 */ 
   } 
   
   //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 = 
   
   /* 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 */
     ((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 =
   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) ;
   
     ((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.
+    

 }
 }