version opérationnelle

[lniv_gpu.git] / main.cu

diff --git a/main.cu b/main.cu
index 0bb7c3d43d83f8b5c55cb37551253ef3e859037e..63bf259d42b679536d30a44871be75f65c70b6ef 100644 (file)
--- 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