X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/snake_gpu.git/blobdiff_plain/935bdd1c8b99ce5d70b6e4a53ca29f06353e8baa..48b1bac747f398161b53a67ad80c4596f531c88a:/src/lib_gpu.cu

diff --git a/src/lib_gpu.cu b/src/lib_gpu.cu
index d6df5ec..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,10 +231,53 @@ void cuda_init_img_cumul(unsigned short ** img_in, int H, int L, int nb_nodes,
    * generation snake en mem GPU
    */
   int dist = 140 ;
+
+  /* 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
+
+  //allocations
+  cudaMalloc((void**) &d_all_crit, Nperm*sizeof(t_rectangle_snake));
+  cudaMalloc((void**) &d_best_crit, sizeof(t_rectangle_snake));
   
   tic(&chrono, NULL);
-  if (nb_nodes == 4)  genere_snake_rectangle_4nodes_gpu<<< 1, 1>>>(*d_snake, 140, H, L) ;
-  else if (nb_nodes == 40) genere_snake_rectangle_Nnodes_gpu<<< 1, 1>>>(*d_snake, (H+L)/20, H, L) ;
+
+  //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");
+
+  //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);
+  /*fin test snake rectangle initial optimal*/
+ 
+  //genere_snake_rectangle_4nodes_gpu<<< 1, 1>>>(*d_snake, 140, H, L) ;
+ 
 
   int nnodes = nb_nodes ;
   snake_node_gpu * h_snake = new snake_node_gpu[nnodes];
@@ -261,15 +304,19 @@ void cuda_init_img_cumul(unsigned short ** img_in, int H, int L, int nb_nodes,
   cudaMalloc((void**) &d_stats_ref, 3*nnodes*sizeof(int64));
 
   //DEBUG : pour forcer la mise à zero du tableau intermediaire d_stats_ref
+  /*
   int64 h_stats_ref[3*nnodes] ;
   for (int a=0; a<3*nnodes ; a++) h_stats_ref[a] = 0 ;
   cudaMemcpy( h_stats_ref, d_stats_ref, sizeof(int64), cudaMemcpyHostToDevice) ;
+  */
   //fin forçage a 0
   
   //DEBUG : pour forcer la mise à zero du tableau intermediaire d_sompart
+  /*
      t_sum_x2 h_sompart[ 3*nnodes*bps ] ;
      for (int a=0; a<3*nnodes*bps ; a++) h_sompart[a] = 0 ;
      cudaMemcpy( h_sompart, d_sompart, sizeof(t_sum_x2), cudaMemcpyHostToDevice) ;
+  */
   //fin forçage a 0
   
   calcul_contribs_segments_snake<<< nnodes*bps, tpb, (CFI(tpb))*(3*sizeof(t_sum_x2))>>>