THESE/Chapters/chapter6/code/convoSepShH.cu

   1 __global__ void kernel_convoSepShx8pH(unsigned char *output, int j_dim, int r)
   2 {
   3   int ic, jc, p;
   4   int k = 2*r+1 ;
   5   float outval0=0.0, outval1=0.0, outval2=0.0, outval3=0.0 ;
   6   float outval4=0.0, outval5=0.0, outval6=0.0, outval7=0.0 ;
   7   int bdimX = blockDim.x<<3 ; // nombre de pixels traités par une ligne d'un bloc
   8   int tidX = threadIdx.x<<3 ; // décalage paquet
   9
  10   // coordonnées absolues du point de base
  11   int j = (__umul24(blockIdx.x,blockDim.x) + threadIdx.x)<<3 ;
  12   int i = __umul24( blockIdx.y, blockDim.y) + threadIdx.y ;
  13   int j0= __umul24(blockIdx.x,blockDim.x)<<3 ;
  14   // indice absolu dans l'image
  15   int idx = __umul24(i,j_dim) + j ;
  16
  17   // adresse d'une ligne de la ROI en mémoire partagée
  18   int idrow = threadIdx.y*(bdimX+k-1) ;
  19
  20   extern __shared__ unsigned char roi8p[];
  21
  22   // gauche de la ROI
  23   for (p=0; p<8; p++)
  24         roi8p[  idrow  + tidX +p ] = tex2D(tex_img_inc, j-r+p  , i) ;
  25   // droite de la ROI
  26   if ( threadIdx.x < r )
  27         {
  28           roi8p[  idrow + bdimX + threadIdx.x    ] = tex2D( tex_img_inc, j0-r
  29                                                             +bdimX+threadIdx.x  , i ) ;
  30           roi8p[  idrow + bdimX + threadIdx.x +r ] = tex2D( tex_img_inc, j0
  31                                                             +bdimX+threadIdx.x  , i ) ;
  32         }
  33
  34   __syncthreads();
  35
  36   // convolution horizontale
  37   for (jc=0 ; jc<k ; jc++)
  38           {
  39                 int baseRoi = idrow + tidX +jc ;
  40                 float valMask = maskh[ jc ] ;
  41                 outval0 += valMask*roi8p[  baseRoi    ] ;
  42                 outval1 += valMask*roi8p[  baseRoi +1 ] ;
  43                 outval2 += valMask*roi8p[  baseRoi +2 ] ;
  44                 outval3 += valMask*roi8p[  baseRoi +3 ] ;
  45                 outval4 += valMask*roi8p[  baseRoi +4 ] ;
  46                 outval5 += valMask*roi8p[  baseRoi +5 ] ;
  47                 outval6 += valMask*roi8p[  baseRoi +6 ] ;
  48                 outval7 += valMask*roi8p[  baseRoi +7 ] ;
  49           }
  50
  51   // sortie des 8 pixels
  52   output[ idx++ ] = outval0 ;
  53   output[ idx++ ] = outval1 ;
  54   output[ idx++ ] = outval2 ;
  55   output[ idx++ ] = outval3 ;
  56   output[ idx++ ] = outval4 ;
  57   output[ idx++ ] = outval5 ;
  58   output[ idx++ ] = outval6 ;
  59   output[ idx   ] = outval7 ;
  60 }