BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu

   1 __global__ void kernel_Median3RegSort9( short *output,
   2                                                                            int i_dim, int j_dim)
   3 {
   4   int j = __mul24(blockIdx.x,blockDim.x) + threadIdx.x ;
   5   int i = __mul24(blockIdx.y,blockDim.y) + threadIdx.y ;
   6   int a0, a1, a2, a3, a4, a5, a6, a7, a8 ; // 1 register per pixel
   7
   8   a0 = tex2D(tex_img_ins, j-1, i-1) ;      // fetching values
   9   a1 = tex2D(tex_img_ins, j  , i-1) ;
  10   a2 = tex2D(tex_img_ins, j+1, i-1) ;
  11   a3 = tex2D(tex_img_ins, j-1, i  ) ;
  12   a4 = tex2D(tex_img_ins, j  , i  ) ;
  13   a5 = tex2D(tex_img_ins, j+1, i  ) ;
  14   a6 = tex2D(tex_img_ins, j-1, i+1) ;
  15   a7 = tex2D(tex_img_ins, j  , i+1) ;
  16   a8 = tex2D(tex_img_ins, j+1, i+1) ;
  17
  18   bubReg9(&a0,&a1,&a2,&a3,&a4,&a5,&a6,&a7,&a8);    // bubble sort
  19
  20   output[ __mul24(i, j_dim) +j ] = a4 ;   // median at the middle
  21  }