BookGPU/Chapters/chapter4/code/convoGeneSh1~

   1 __global__ void kernel_convoNonSepSh_8p(unsigned char *output, int i_dim, int j_dim, int r)
   2 {
   3   int idb, ic, jc;
   4   int L = 2*r+1 ;
   5   float outval0=0.0, outval1=0.0, outval2=0.0, outval3=0.0, outval4=0.0, outval5=0.0, outval6=0.0, outval7=0.0 ;
   6   int bdimX = blockDim.x<<3 ;
   7   int tidX = threadIdx.x<<3 ;
   8
   9   // coordonnees absolues du point de base
  10   int j = (__umul24(blockIdx.x,blockDim.x) + threadIdx.x)<<3 ;
  11   int i = __umul24( blockIdx.y, blockDim.y) + threadIdx.y ;
  12   int j0= __umul24(blockIdx.x,blockDim.x)<<3 ;
  13
  14   int idx = __umul24(i,j_dim) + j ;              // indice  dans l'image
  15
  16
  17   // chargement en smem
  18   int idrow = threadIdx.y*(bdimX+L-1) ;
  19
  20   extern __shared__ unsigned char roi8p[];
  21
  22   // bloc 0 (en haut à gauche)
  23   for (int p=0; p<8; p++)
  24   roi8p[ idrow  + tidX +p ] = tex2D(tex_img_inc, j-r+p  , i-r) ;
  25
  26   // bloc 1 (en haut à droite)...
  27   if ( threadIdx.x < r ) //...ou plutot ce qu'il en manque
  28         {
  29           roi8p[ idrow + bdimX + threadIdx.x    ] = tex2D( tex_img_inc, j0-r +bdimX+threadIdx.x  , i-r ) ;
  30           roi8p[ idrow + bdimX + threadIdx.x +r ] = tex2D( tex_img_inc, j0   +bdimX+threadIdx.x  , i-r ) ;
  31         }
  32   // bloc 2 ( en bas à gauche)
  33   if ( threadIdx.y < L-1 )
  34         {
  35           idrow = (threadIdx.y+blockDim.y)*(bdimX+L-1) ;
  36           for (int p=0; p<8; p++)
  37                 roi8p[ idrow + tidX +p  ] = tex2D( tex_img_inc, j-r+p  , i+blockDim.y-r ) ;
  38
  39           //bloc 4 ( en bas à doite )...
  40           if ( threadIdx.x < r ) //...ou ce qu'il en manque
  41                 {
  42                   roi8p[ idrow + bdimX +threadIdx.x   ] = tex2D( tex_img_inc, j0-r +bdimX +threadIdx.x, i+blockDim.y-r ) ;
  43                   roi8p[ idrow + bdimX +threadIdx.x +r] = tex2D( tex_img_inc, j0   +bdimX +threadIdx.x, i+blockDim.y-r ) ;
  44                 }
  45         }
  46   __syncthreads();
  47
  48   // calculs de convolution
  49   for (ic=0 ; ic<L ; ic++)
  50         for( jc=0 ; jc<L ; jc++)
  51           {
  52                 int baseRoi = __umul24(ic+threadIdx.y,(bdimX+L-1)) + jc+tidX ;
  53                 float valMask = masque[ __umul24(ic,L)+jc ] ;
  54                 outval0 += valMask*roi8p[ baseRoi ] ;
  55                 outval1 += valMask*roi8p[ baseRoi +1 ] ;
  56                 outval2 += valMask*roi8p[ baseRoi +2 ] ;
  57                 outval3 += valMask*roi8p[ baseRoi +3 ] ;
  58                 outval4 += valMask*roi8p[ baseRoi +4 ] ;
  59                 outval5 += valMask*roi8p[ baseRoi +5 ] ;
  60                 outval6 += valMask*roi8p[ baseRoi +6 ] ;
  61                 outval7 += valMask*roi8p[ baseRoi +7 ] ;
  62           }
  63
  64   // 1 pixel par thread --> global mem
  65   output[ idx   ] = outval0 ;
  66   output[ idx+1 ] = outval1 ;
  67   output[ idx+2 ] = outval2 ;
  68   output[ idx+3 ] = outval3 ;
  69   output[ idx+4 ] = outval4 ;
  70   output[ idx+5 ] = outval5 ;
  71   output[ idx+6 ] = outval6 ;
  72   output[ idx+7 ] = outval7 ;
  73 }