X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/lniv_gpu.git/blobdiff_plain/7cc36a5689750ee7d77b959362edc65948a79427..refs/heads/master:/main.cu

diff --git a/main.cu b/main.cu
index 0bb7c3d..63bf259 100644
--- a/main.cu
+++ b/main.cu
@@ -53,7 +53,7 @@ int main(int argc, char **argv){
   cutilCheckError( cutStopTimer(timer) );
   
   size = H * L * sizeof( unsigned int );
-  psize = (r-1)*PSIZE_I*sizeof(int2) ;
+  psize = (r-1)*PSIZE_I*sizeof(ushort) ;
 
   printf("Longueur des chemins = %d pixels\n", r);
   printf("Init GPU + Image %s  %d x %d = %d pixels en %f ms,\n", image_path, L, H, size, cutGetTimerValue(timer));
@@ -67,7 +67,7 @@ int main(int argc, char **argv){
   // allocation mem GPU
   unsigned int * d_directions =NULL ;
   unsigned int * d_lniv, * d_estim = NULL ;
-  int2 * d_paths ;
+  ushort * d_paths ;
 
   cutilSafeCall( cudaMalloc( (void**) &d_directions, size)) ;
   cutilSafeCall( cudaMalloc( (void**) &d_lniv, size ) );
@@ -77,7 +77,8 @@ int main(int argc, char **argv){
   
   // allocate array and copy image data
   cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindUnsigned);
-  cudaChannelFormatDesc channelDescP = cudaCreateChannelDesc(32, 32, 0, 0, cudaChannelFormatKindSigned);
+  //cudaChannelFormatDesc channelDescP = cudaCreateChannelDesc(8, 8, 0, 0, cudaChannelFormatKindSigned);
+  cudaChannelFormatDesc channelDescP = cudaCreateChannelDesc<ushort>();
   
   cudaArray * array_img_in, *array_img_estim, *array_img_lniv, *array_paths;
   cutilSafeCall( cudaMallocArray( &array_img_in, &channelDesc, L, H )); 
@@ -116,7 +117,7 @@ int main(int argc, char **argv){
   /*****************************
    * APPELS KERNELS et chronos
    *****************************/
-	dimBlock = dim3(16,16,1) ;
+	dimBlock = dim3(8,8,1) ;
 	dimGrid = dim3( H / dimBlock.x,L / dimBlock.y, 1 ) ;
 
 	// pour enregistrement image lniv GPU
@@ -151,7 +152,8 @@ int main(int argc, char **argv){
 	for ( iter =0 ; iter < nb_iter ; iter++ )
 	  {
 		cutilSafeCall( cudaMemcpyToArray( array_img_estim, 0, 0, d_estim, size, cudaMemcpyDeviceToDevice)) ;
-		kernel_levelines_texture_smem<<< dimGrid, dimBlock, 0 >>>( d_lniv, L, H, r );
+		//version avec/sans tentative d'utilisation de smem ( pas probante )
+		kernel_levelines_texture<<< dimGrid, dimBlock, 24*(r-1)*sizeof(short) >>>( d_lniv, L, H, r );
 		cutilSafeCall( cudaMemcpyToArray( array_img_lniv, 0, 0, d_lniv, size, cudaMemcpyDeviceToDevice)) ;
 		if (seq_out){
 		  sprintf(image_out, "%s%d%s", image_out_base, iter+1, pgm_ext) ;
@@ -175,10 +177,10 @@ int main(int argc, char **argv){
 	/**************************************************
 	 * recuperation matrice des chemins pour affichage
 	 **************************************************/
-	
+	/*
 	int2 * h_paths = new int2[(r-1)*PSIZE_I] ;
 	cutilSafeCall( cudaMemcpyFromArray(h_paths , array_paths, 0, 0, psize, cudaMemcpyDeviceToHost) );
-	/*
+	
 	//verif Di
 	printf("matrice  Di");
 	for(int idpath=0; idpath< PSIZE_I; idpath++){b