BookGPU/Chapters/chapter4/code/convoGene8r.cu

   1 __global__ void kernel_convoGene8r( unsigned char *output, int j_dim, int r)
   2 {
   3   int ic, jc ;
   4   int k=2*r+1 ;
   5   float outval0=0.0 ;
   6
   7   // absolute coordinates of base point
   8   int j = __umul24( blockIdx.x, blockDim.x ) + threadIdx.x ;
   9   int i = __umul24( blockIdx.y, blockDim.y) + threadIdx.y ;
  10
  11   // convolution computation
  12   for (ic=-r ; ic<=r ; ic++)
  13         for(jc=-r ; jc<=r ; jc++)
  14           outval0 += mask[ __umul24(ic,k)+jc+r ]
  15                         *tex2D(tex_img_inc, j+jc, i+ic) ;
  16
  17   output[ __umul24(i, j_dim) + j ] = outval0 ;
  18 }