X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/snake_gpu.git/blobdiff_plain/935bdd1c8b99ce5d70b6e4a53ca29f06353e8baa..refs/heads/master:/src/lib_kernels_contribs.cu

diff --git a/src/lib_kernels_contribs.cu b/src/lib_kernels_contribs.cu
index 2440f0c..43dc923 100644
--- a/src/lib_kernels_contribs.cu
+++ b/src/lib_kernels_contribs.cu
@@ -103,7 +103,7 @@ __global__ void calcul_contribs_segments_blocs_full(snake_node_gpu * d_snake, in
   // indices des elements
   int blockSize = blockDim.x ;					                        // nb threads par bloc
   int tib = threadIdx.x ;						                        // position du thread dans le bloc
-  int nblocs_noeud =  gridDim.x / (nb_nodes/2 + pairs*(nb_nodes%2)) ;   // nb de blocs dédié à chaque noeud 
+  int nblocs_noeud =  gridDim.x / nb_nodes ;                            // nb de blocs dédié à chaque noeud 
   int nblocs_seg = nblocs_noeud / 16 ;                                  // nb de blocs dédiés à un segment de test
   int idx = blockDim.x*blockIdx.x + threadIdx.x ;                       // position absolue du thread ds la grille
   int id_interval = blockIdx.x / nblocs_noeud ;                         // indice de l'intervalle du noeud dans la grille
@@ -125,18 +125,10 @@ __global__ void calcul_contribs_segments_blocs_full(snake_node_gpu * d_snake, in
   int xprec, xsuiv ; 
 
   // determine les indices des noeuds prec, courant, suiv
-  if (pairs)
-	{
-	  n1 = 2*id_interval -1 ;
-	  n2 = 2*id_interval ;
-	  n3 = 2*id_interval +1 ;
-	}
-  else
-	{
-	  n1 = 2*id_interval ;
-	  n2 = 2*id_interval +1 ;
-	  n3 = 2*id_interval +2 ;
-	}
+  n1 = id_interval -1 ;
+  n2 = id_interval ;
+  n3 = id_interval +1 ;
+
   //gestion du bouclage du snake
   if (n1 < 0) n1 = nb_nodes-1 ;
   if (n3 >= nb_nodes) n3 = 0 ;
@@ -293,6 +285,8 @@ __global__ void calcul_contribs_segments_blocs_full(snake_node_gpu * d_snake, in
  }
 }
 
+
+
 /*
  calcul des freeman et du centre de chaque segment de test
  a executer sur 'n_interval' blocs de 16 threads
@@ -528,7 +522,7 @@ __global__ void soustrait_aux_stats_2N_segments_noeud(snake_node_gpu * d_snake,
 
   EXEC : sur n_interval blocs de 8 threads
 */
-__global__ void calcul_stats_full(snake_node_gpu * d_snake, int nnodes, bool pairs, int64 * d_stats_snake,
+__global__ void calcul_stats_full(snake_node_gpu * d_snake, snake_node_gpu * d_snake_tmp, int nnodes, bool pairs, int64 * d_stats_snake,
 								  int64 * d_stats_ref, int64 * d_stats, uint64 * d_contribs,
 								  uint4 * d_liste_points, int * code_segment, uint4 * d_freemans,
 								  int * d_table_codes, t_cumul_x * cumul_x, t_cumul_x2 * cumul_x2,
@@ -541,7 +535,7 @@ __global__ void calcul_stats_full(snake_node_gpu * d_snake, int nnodes, bool pai
   int code_noeud; 
   __shared__ int64 s_stats_snake[3*8] ;
 
-  id_nx = 2*interval + !pairs ;
+  id_nx = interval ;
   if (id_nx == 0) id_nprec = nnodes - 1 ;
   else id_nprec = id_nx - 1 ;
   if (id_nprec == 0) id_nprecprec = nnodes -1 ;
@@ -672,44 +666,19 @@ __global__ void calcul_stats_full(snake_node_gpu * d_snake, int nnodes, bool pai
 		pos_optim = v;
 	  }
 	}
-	if (pos_optim >-1){
-	  if ( !croisement(d_snake, id_nx, d_liste_points[8*id_nx + pos_optim].x, d_liste_points[8*id_nx + pos_optim].y, nnodes) )
+	if ( (pos_optim >-1) && !croisement(d_snake, id_nx, d_liste_points[8*id_nx + pos_optim].x, d_liste_points[8*id_nx + pos_optim].y, nnodes) )
 		{
 		  /*maj data snake*/
 		  move[id_nx] = true ;
 		  //new position 
-		  d_snake[id_nx].posi = d_liste_points[8*id_nx + pos_optim].x ;
-		  d_snake[id_nx].posj = d_liste_points[8*id_nx + pos_optim].y ;
-		  //nb pixels segment precedent
-		  d_snake[id_nprec].nb_pixels = d_liste_points[8*id_nx + pos_optim].z ;
-		  //nb pixels segment suivant
-		  d_snake[id_nx].nb_pixels = d_liste_points[8*id_nx + pos_optim].w ;
-		  //contribs segment precedent
-		  d_snake[id_nprec].sum_1 = d_contribs[3*(16*interval + pos_optim)] ;
-		  d_snake[id_nprec].sum_x = d_contribs[3*(16*interval + pos_optim) + 1] ;
-		  d_snake[id_nprec].sum_x2 = d_contribs[3*(16*interval + pos_optim) + 2] ;
-		  //contribs segment suivant
-		  d_snake[id_nx].sum_1 = d_contribs[3*(16*interval + pos_optim + 8)] ;
-		  d_snake[id_nx].sum_x = d_contribs[3*(16*interval + pos_optim + 8) + 1] ;
-		  d_snake[id_nx].sum_x2 = d_contribs[3*(16*interval + pos_optim + 8) + 2] ;
-		  //freemans segment precedent
-		  d_snake[id_nprec].freeman_out = d_freemans[16*interval + pos_optim].z ;
-		  d_snake[id_nprec].freeman_in = d_freemans[16*interval + pos_optim].w ;
-		  //freemans segment suivant
-		  d_snake[id_nx].freeman_out = d_freemans[16*interval + pos_optim + 8].z ;
-		  d_snake[id_nx].freeman_in = d_freemans[16*interval + pos_optim + 8].w ;
-		  //codes segment precedent
-		  d_snake[id_nprec].code_segment = code_segment[16*interval + pos_optim] ;
-		  //code segment suivant
-		  d_snake[id_nx].code_segment = code_segment[16*interval + pos_optim + 8] ;
-		  //centre segment precedent
-		  d_snake[id_nprec].centre_i = d_freemans[16*interval + pos_optim ].x ;
-		  d_snake[id_nprec].centre_j = d_freemans[16*interval + pos_optim ].y ;
-		  //centre segment suivant 
-		  d_snake[id_nx].centre_i = d_freemans[16*interval + pos_optim + 8].x ;
-		  d_snake[id_nx].centre_j = d_freemans[16*interval + pos_optim + 8].y ;
+		  d_snake_tmp[id_nx].posi = d_liste_points[8*id_nx + pos_optim].x ;
+		  d_snake_tmp[id_nx].posj = d_liste_points[8*id_nx + pos_optim].y ;
+		}
+	else {
+		  //keep old position 
+		  d_snake_tmp[id_nx].posi = d_snake[id_nx].posi ;
+		  d_snake_tmp[id_nx].posj = d_snake[id_nx].posj ;
 		  
-		}			   
 	}
   }
 
@@ -787,32 +756,39 @@ __global__ void recalcul_stats_snake(snake_node_gpu * d_snake, int  nnodes, int6
 
 __global__ void ajoute_noeuds(snake_node_gpu * snake, snake_node_gpu * snake_tmp, int nnodes, int seuil, int * new_nb_nodes){
 
-  volatile snake_node_gpu * st = snake_tmp ;
-
+  
  int id_cpy = 0;  
  for (int id_nx=0; id_nx < nnodes; id_nx++){
 	//position du noeud existant
-	st[id_cpy].posi =  snake[id_nx].posi ;
-    st[id_cpy].posj =  snake[id_nx].posj ;
+	snake_tmp[id_cpy].posi =  snake[id_nx].posi ;
+    snake_tmp[id_cpy].posj =  snake[id_nx].posj ;
         
 	id_cpy++ ;
 	
-	if ( snake[id_nx].nb_pixels > seuil)
+	if ( snake_tmp[id_nx].nb_pixels > seuil)
 	{
 		//position du nouveau noeud
-		st[id_cpy].posi = snake[id_nx].centre_i ;
-		st[id_cpy].posj = snake[id_nx].centre_j ;
+		snake_tmp[id_cpy].posi = snake[id_nx].centre_i ;
+		snake_tmp[id_cpy].posj = snake[id_nx].centre_j ;
 		id_cpy++ ;
 	}	
  }
- for( int node=0; node<id_cpy; node++ ){
-   snake[node].posi = st[node].posi ;
-   snake[node].posj = st[node].posj ;
+
+ for (int id_nx=0; id_nx < id_cpy; id_nx++){
+	snake[id_cpy].posi =  snake_tmp[id_nx].posi ;
+    snake[id_cpy].posj =  snake_tmp[id_nx].posj ;
  }
+ 
 *new_nb_nodes = id_cpy-nnodes ;
 }
-  
 
+// pour copier le contenu d'un snake dans un autre
+// execution sur nnodes blocs de 1 threads
+__global__ void copie_snake(snake_node_gpu * d_source, snake_node_gpu * d_dest){
+  int id = blockIdx.x ; 
+  d_dest[id].posi =  d_source[id].posi ;
+  d_dest[id].posj =  d_source[id].posj ;
+}
 
 
 
@@ -985,7 +961,8 @@ __global__ void recalcul_contribs_segments_snake(snake_node_gpu * d_snake, int n
 
 __global__ void recalcul_freemans_centre(snake_node_gpu * snake, uint2 * liste_pix, int * d_table_freeman){
 	
-  int id_segment = blockIdx.x ;    
+  int id_segment = blockIdx.x ;
+  int nb_segments = gridDim.x ;
   int id_base_pix = 5*id_segment ;    
 
   //calculs freemans, centre et code segment
@@ -1005,9 +982,13 @@ __global__ void recalcul_freemans_centre(snake_node_gpu * snake, uint2 * liste_p
 	//centre 
 	snake[id_segment].centre_i = liste_pix[id_base_pix +2].x ;
 	snake[id_segment].centre_j = liste_pix[id_base_pix +2].y ;
-	   
-}
 
+	//nb_pixels
+	int nd = id_segment ;
+	int nd_suiv = nd + 1  ;
+	if (nd_suiv >= nb_segments ) nd_suiv = 0 ;
+	snake[id_segment].nb_pixels = calcul_nb_pixels(snake[nd].posi, snake[nd].posj, snake[nd_suiv].posi, snake[nd_suiv].posj) ;
+}
 
 /*
   sommme des contribs par bloc -> contribs segment, pour le snake