main.cu

   1 // libs C
   2 #include <stdlib.h>
   3 #include <stdio.h>
   4 #include <string.h>
   5 #include <math.h>
   6
   7 #include "lib_lniv.h"
   8
   9 // libs NV
  10 #include <cuda_runtime.h>
  11 #include <cutil_inline.h>
  12
  13 // lib spec
  14 #include "defines.h"
  15 #include "levelines_common.h"
  16 #include "levelines_kernels.cu"
  17
  18
  19 int main(int argc, char **argv){
  20
  21   // use device with highest Gflops/s
  22   cudaSetDevice( 0 );
  23
  24   unsigned int timer ;
  25   cutilCheckError( cutCreateTimer(&timer) );
  26   cutilCheckError( cutResetTimer(timer) );
  27   cutilCheckError( cutStartTimer(timer) );
  28
  29
  30   //alloc bidon pour anticiper l'initialisation du GPU
  31   short * d_bidon ;
  32   cutilSafeCall( cudaMalloc( (void**) &d_bidon, sizeof(short))) ;
  33
  34   //float coef_regul = atof( argv[1] ) ;
  35
  36   /*****************************
  37    *    CHARGEMENT IMAGE
  38    *****************************/
  39   char* image_path = argv[argc-1];
  40   unsigned int r = atoi(argv[1]) ;
  41   bool seq_out = atoi(argv[2]) ;
  42   unsigned int iter , nb_iter = atoi(argv[3]) ;
  43   unsigned int  poids = 15 ;
  44   char * image_out_base = "./image_out" ;
  45   char * pgm_ext = ".pgm" ;
  46   char image_out[80] ;
  47   unsigned int * h_data = NULL ;
  48   unsigned int * h_data_out = NULL ;
  49   unsigned int H, L, size, psize ;
  50
  51   // chargt image
  52   cutilCheckError( cutLoadPGMi(image_path, &h_data, &L, &H));
  53   cutilCheckError( cutStopTimer(timer) );
  54
  55   size = H * L * sizeof( unsigned int );
  56   psize = (r-1)*PSIZE_I*sizeof(int2) ;
  57
  58   printf("Longueur des chemins = %d pixels\n", r);
  59   printf("Init GPU + Image %s  %d x %d = %d pixels en %f ms,\n", image_path, L, H, size, cutGetTimerValue(timer));
  60
  61   /*****************************
  62    *     FIN CHARGEMENT IMAGE
  63    *****************************/
  64
  65   cutilCheckError( cutResetTimer(timer) );
  66   cutilCheckError( cutStartTimer(timer) );
  67   // allocation mem GPU
  68   unsigned int * d_directions =NULL ;
  69   unsigned int * d_lniv, * d_estim = NULL ;
  70   int2 * d_paths ;
  71
  72   cutilSafeCall( cudaMalloc( (void**) &d_directions, size)) ;
  73   cutilSafeCall( cudaMalloc( (void**) &d_lniv, size ) );
  74   cutilSafeCall( cudaMalloc( (void**) &d_estim, size ) );
  75   cutilSafeCall( cudaMalloc( (void**) &d_paths, psize ) );
  76
  77
  78   // allocate array and copy image data
  79   cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindUnsigned);
  80   cudaChannelFormatDesc channelDescP = cudaCreateChannelDesc(32, 32, 0, 0, cudaChannelFormatKindSigned);
  81
  82   cudaArray * array_img_in, *array_img_estim, *array_img_lniv, *array_paths;
  83   cutilSafeCall( cudaMallocArray( &array_img_in, &channelDesc, L, H ));
  84   cutilSafeCall( cudaMemcpyToArray( array_img_in, 0, 0, h_data, size, cudaMemcpyHostToDevice)) ;
  85   cutilSafeCall( cudaBindTextureToArray( tex_img_in, array_img_in, channelDesc));
  86
  87   cutilSafeCall( cudaMallocArray( &array_img_estim, &channelDesc, L, H ));
  88   cutilSafeCall( cudaBindTextureToArray( tex_img_estim, array_img_estim, channelDesc));
  89
  90   cutilSafeCall( cudaMallocArray( &array_img_lniv, &channelDesc, L, H ));
  91   cutilSafeCall( cudaBindTextureToArray( tex_img_lniv, array_img_lniv, channelDesc));
  92
  93   cutilSafeCall( cudaMallocArray( &array_paths, &channelDescP, (r-1), PSIZE_I ));
  94   cutilSafeCall( cudaBindTextureToArray( tex_paths, array_paths, channelDescP));
  95
  96   cutilCheckError( cutStopTimer(timer) );
  97   printf("Temps alloc + transferts en Textures : %f ms\n", cutGetTimerValue(timer)) ;
  98
  99   /*****************************
 100    * GENERATION des CHEMINS
 101    *****************************/
 102   cutilCheckError( cutResetTimer(timer) );
 103   cutilCheckError( cutStartTimer(timer) );
 104
 105   dim3 dimBlock(1,1,1) ;
 106   dim3 dimGrid(1,1,1) ;
 107   // calcul des chemins
 108   kernel_calcul_paths<<< dimGrid, dimBlock, 0 >>>(d_paths, r);
 109
 110   // copie du tableau en texture
 111   cutilSafeCall( cudaMemcpyToArray( array_paths, 0, 0, d_paths, psize, cudaMemcpyDeviceToDevice)) ;
 112
 113   cutilCheckError( cutStopTimer(timer) );
 114   printf("Temps generation chemin + transfert en texture : %f ms\n", cutGetTimerValue(timer)) ;
 115
 116   /*****************************
 117    * APPELS KERNELS et chronos
 118    *****************************/
 119         dimBlock = dim3(16,16,1) ;
 120         dimGrid = dim3( H / dimBlock.x,L / dimBlock.y, 1 ) ;
 121
 122         // pour enregistrement image lniv GPU
 123         free(h_data_out) ;
 124         h_data_out = new unsigned int[H*L] ;
 125
 126         //init image estimee avec image_in
 127         cutilCheckError( cutResetTimer(timer) );
 128         cutilCheckError( cutStartTimer(timer) );
 129         kernel_init_estim_from_img_in<<< dimGrid, dimBlock, 0 >>>(d_estim, L, H, 7);
 130
 131         // kernel de copie
 132         /*
 133         kernel_neutre_img2estim<<< dimGrid, dimBlock, 0>>>(d_estim, L, H);
 134         */
 135         cudaThreadSynchronize() ;
 136         cutilCheckError( cutStopTimer(timer) );
 137         printf("Temps kernel init : %f ms\n", cutGetTimerValue(timer)) ;
 138         // a remplacer par
 139         /*
 140         cutilCheckError( cutResetTimer(timer) );
 141         cutilCheckError( cutStartTimer(timer) );
 142         cutilSafeCall( cudaMemcpyFromArray( d_estim, array_img_in, 0, 0, size, cudaMemcpyDeviceToDevice)) ;
 143         cutilCheckError( cutStopTimer(timer) );
 144         printf("Temps memcpyFromArray : %f ms\n", cutGetTimerValue(timer)) ;
 145         */
 146         printf("Grille : %d x %d de Blocs : %d x %d \n", dimGrid.x, dimGrid.y, dimBlock.x, dimBlock.y) ;
 147
 148         cutilCheckError( cutResetTimer(timer) );
 149         cutilCheckError( cutStartTimer(timer) );
 150
 151         for ( iter =0 ; iter < nb_iter ; iter++ )
 152           {
 153                 cutilSafeCall( cudaMemcpyToArray( array_img_estim, 0, 0, d_estim, size, cudaMemcpyDeviceToDevice)) ;
 154                 kernel_levelines_texture_smem<<< dimGrid, dimBlock, 0 >>>( d_lniv, L, H, r );
 155                 cutilSafeCall( cudaMemcpyToArray( array_img_lniv, 0, 0, d_lniv, size, cudaMemcpyDeviceToDevice)) ;
 156                 if (seq_out){
 157                   sprintf(image_out, "%s%d%s", image_out_base, iter+1, pgm_ext) ;
 158                   printf("chaine : %s\n", image_out);
 159                   cutilSafeCall( cudaMemcpy(h_data_out , d_estim, size, cudaMemcpyDeviceToHost) );
 160                   cutilCheckError( cutSavePGMi(image_out, h_data_out, L, H) ) ;
 161                   }
 162                 kernel_estim_next_step_texture<<< dimGrid, dimBlock, 0 >>>(d_estim, L, H, poids) ;
 163         }
 164
 165         cudaThreadSynchronize();
 166
 167         cutilCheckError( cutStopTimer(timer) );
 168         printf("Execution moy par kernel : %f ms\n", cutGetTimerValue(timer)/(float)nb_iter) ;
 169         printf("Total pour %d kernels : %f ms\n", nb_iter, cutGetTimerValue(timer)) ;
 170
 171         /**************************
 172          * VERIFS
 173          **************************/
 174
 175         /**************************************************
 176          * recuperation matrice des chemins pour affichage
 177          **************************************************/
 178
 179         int2 * h_paths = new int2[(r-1)*PSIZE_I] ;
 180         cutilSafeCall( cudaMemcpyFromArray(h_paths , array_paths, 0, 0, psize, cudaMemcpyDeviceToHost) );
 181         /*
 182         //verif Di
 183         printf("matrice  Di");
 184         for(int idpath=0; idpath< PSIZE_I; idpath++){b
 185
 186           printf("\n");
 187           for(int idpix=0; idpix< r-1; idpix++){
 188                 printf(" % d ", h_paths[idpath*(r-1) + idpix].x );
 189           }
 190           printf("\t// %d°", idpath*15) ;
 191         }
 192         //verif Dj
 193         printf("\nmatrice  Dj");
 194         for(int idpath=0; idpath< PSIZE_I; idpath++){
 195           printf("\n");
 196           for(int idpix=0; idpix< r-1; idpix++){
 197                 printf(" % d ", h_paths[idpath*(r-1) + idpix].y);
 198           }
 199           printf("\t// %d°", idpath*15) ;
 200         }
 201         printf("\n");
 202         */
 203         /***************************************************
 204          *   fin verif visuelle matrices des chemins
 205          ***************************************************/
 206
 207         /***************************************************
 208          *       execution sequentielle pour comparaison
 209          *  la comparaison n'est pertinente que
 210          * si d_lniv contient les lniv de l'image se départ
 211          **************************************************/
 212         /*
 213         // calcul sequentiel
 214         cutilCheckError( cutResetTimer(timer) );
 215         cutilCheckError( cutStartTimer(timer) );
 216
 217         h_data_out = new unsigned int[H*L] ;
 218         int * dout  = new int[H*L] ;
 219         for ( iter =0 ; iter < nb_iter ; iter++ ){
 220           for (int i=r; i<= H-r; i++){
 221                 for (int j=r; j<= L-r; j++){
 222                   h_data_out[i*L+j] = lniv4_value(h_data, h_paths, i, j, H, L, &dout[i*L+j], r) ;
 223                 }
 224           }
 225         }
 226         cutilCheckError( cutStopTimer(timer) );
 227         printf("Execution sequentielle CPU : %f ms\n", cutGetTimerValue(timer)) ;
 228
 229         // comparaison
 230         unsigned int * h_lniv = new unsigned int[H*L] ;
 231         int pos, cpt_err=0, cpt_pix=0 ;
 232         cutilSafeCall( cudaMemcpy(h_lniv , d_lniv, size, cudaMemcpyDeviceToHost) );
 233         for ( iter =0 ; iter < nb_iter ; iter++ ){
 234           for (int i=r; i<= H-r; i++){
 235                 for (int j=r; j<= L-r; j++){
 236                   pos = i*L + j ;
 237                   if ( h_data_out[ pos ] != h_lniv[ pos ] ) {
 238                         cpt_err++ ;
 239                         printf(" pixel ( %d , %d ) -> GPU= %d  CPU= %d \n", i, j, h_lniv[pos], h_data_out[pos]);
 240                   }
 241                   cpt_pix++ ;
 242                 }
 243           }
 244         }
 245         printf("TAUX ERREUR GPU/CPU : %d / %d \n", cpt_err, cpt_pix );
 246         */
 247         //trace des lniv sur grille de 'pas x pas'
 248         //kernel_trace_levelines<<< dimGrid, dimBlock, 0 >>>(d_data, d_directions, d_data2, L, H, 16, 255, r) ;
 249
 250         if (!seq_out){
 251           if ( h_data_out != NULL)
 252                 cutilSafeCall( cudaMemcpy(h_data_out , d_estim, size, cudaMemcpyDeviceToHost) );
 253           else
 254                 printf("Echec allocation mem CPU\n");
 255           sprintf(image_out, "%s%d%s", image_out_base, iter+1, pgm_ext) ;
 256           cutilCheckError( cutSavePGMi(image_out, h_data_out, L, H) ) ;
 257         }
 258
 259         // TODO verifier pourquoi les deux lignes suivantes produisent une erreur
 260         //cutilExit(argc, argv);
 261         //cudaThreadExit();
 262         return EXIT_SUCCESS ;
 263 }