Fin de test multisnake sur iter 1

Pas performant

diff --git a/exec/SNAKE2D b/exec/SNAKE2D
index 1d1de071a0d2a0dcc1c7f70a2ca453a92fabc7f3..40c66024076844a19714280bbc66d872d0a5a098 100755 (executable)
Binary files a/exec/SNAKE2D and b/exec/SNAKE2D differ

diff --git a/lib/lib_contour.o b/lib/lib_contour.o
index 26fc29db2a5d6849f73dc93f103d4fc005baa0a5..e197c41044bf1f37622f1d7d48a040fb27508838 100644 (file)
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 1d978d71c79b038ef396127e2d289423a7156f74..dc40914efc7094023ffe440c5a153c150a20f8b5 100644 (file)
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 4a3a81f9f538ab9912f4eee5c83e875c41a45c4d..d74fedd6c903794c99b963174e1cd494e99e7515 100644 (file)
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 a7abed2ef461d8a955933849c850726a52adb319..4abaa10263948e0b65d0fafec97c96e3c0d0be65 100644 (file)
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 5a6e207084fa2fad62942bcf51e75e492d46e1d1..676e14d94b4e7d6b37b6e874f1da597a93a36d7b 100644 (file)
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 7be4b887b6719091caa4945ae7b133fa8bad7209..3b6bf11f1860977ebd5ecee462f954e1d3d5542c 100644 (file)
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 1bbd61d5ca8b4b564b1984762ace1cf1c245d7b6..5a90c1ccf8169373919d07eeeedc449488f8dcdd 100644 (file)
--- 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);
   
   
   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
        { 
          
          //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 ;
    * 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<<<grid, bs, CFI(bs)*sizeof(tcontribs) >>>(*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);
   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 ;
  
 
   int nnodes = nb_nodes ;

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.
+    

 }
 }

diff --git a/src/structures.h b/src/structures.h
index b2a9554a159d39a6a0577c69c55bdc6c88979cb4..3977661b3b2ee86de41e8697f662faa068914799 100644 (file)
--- a/src/structures.h
+++ b/src/structures.h
@@ -141,5 +141,13 @@ struct snake_node_gpu
   int last_move ; /* dernier test deplacement accepte */
 };
 
   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
 
 #endif