BookGPU/Chapters/chapter3/code/kernMedian2pix5.cu

   1 __global__ void kernel_median5_2pix( short *output,
   2                                                                          int i_dim, int j_dim)
   3 {
   4   int j= __mul24(__mul24(blockIdx.x,blockDim.x) + threadIdx.x,2);
   5   int i= __mul24(blockIdx.y,blockDim.y) + threadIdx.y;
   6   int a0,a1,a2,a3,a4,a5,a6,a7,a8,a9,a10,a11,a12,a13;//left window
   7   int b6,b7,b8,b9,b10,b11,b12,b13 ;                 //right window
   8   //first 14 common pixels
   9   a0  = tex2D(tex_img_ins, j-1, i-2) ;  // first line
  10   a1  = tex2D(tex_img_ins, j  , i-2) ;
  11   a2  = tex2D(tex_img_ins, j+1, i-2) ;
  12   a3  = tex2D(tex_img_ins, j+2, i-2) ;
  13   a4  = tex2D(tex_img_ins, j-1, i-1) ;  //seconde line
  14   a5  = tex2D(tex_img_ins, j  , i-1) ;
  15   a6  = tex2D(tex_img_ins, j+1, i-1) ;
  16   a7  = tex2D(tex_img_ins, j+2, i-1) ;
  17   a8  = tex2D(tex_img_ins, j-1, i) ;    // third line
  18   a9  = tex2D(tex_img_ins, j  , i) ;
  19   a10 = tex2D(tex_img_ins, j+1, i) ;
  20   a11 = tex2D(tex_img_ins, j+2, i) ;    // first 2 of fourth line
  21   a12 = tex2D(tex_img_ins, j-1, i+1) ;
  22   a13 = tex2D(tex_img_ins, j  , i+1) ;
  23
  24   //common selection
  25   minmax14(&a0,&a1,&a2,&a3,&a4,&a5,&a6,&a7,&a8,&a9,&a10,&a11,&a12,&a13);
  26   a13 = tex2D(tex_img_ins, j+1, i+1);
  27   minmax13(&a1,&a2,&a3,&a4,&a5,&a6,&a7,&a8,&a9,&a10,&a11,&a12,&a13);
  28   a13 = tex2D(tex_img_ins, j+2, i+1);
  29   minmax12(&a2,&a3,&a4,&a5,&a6,&a7,&a8,&a9,&a10,&a11,&a12,&a13);
  30   a13 = tex2D(tex_img_ins, j-1, i+2);
  31   minmax11(&a3,&a4,&a5,&a6,&a7,&a8,&a9,&a10,&a11,&a12,&a13);
  32   a13 = tex2D(tex_img_ins, j  , i+2);
  33   minmax10(&a4,&a5,&a6,&a7,&a8,&a9,&a10,&a11,&a12,&a13);
  34   a13 = tex2D(tex_img_ins, j+1, i+2);
  35   minmax9(&a5,&a6,&a7,&a8,&a9,&a10,&a11,&a12,&a13);
  36   a13 = tex2D(tex_img_ins, j+2, i+2);
  37   minmax8(&a6,&a7,&a8,&a9,&a10,&a11,&a12,&a13);
  38
  39   // separation
  40   b6=a6; b7=a7; b8=a8; b9=a9; b10=a10; b11=a11; b12=a12; b13=a13;
  41
  42   // separate selections: 5 remaining pixels in both windows
  43   a13 = tex2D(tex_img_ins, j-2, i-2);
  44   b13 = tex2D(tex_img_ins, j+3, i-2);
  45   minmax7(&a7,&a8,&a9,&a10,&a11,&a12,&a13);
  46   minmax7(&b7,&b8,&b9,&b10,&b11,&b12,&b13);
  47   a13 = tex2D(tex_img_ins, j-2, i-1);
  48   b13 = tex2D(tex_img_ins, j+3, i-1);
  49   minmax6(&a8,&a9,&a10,&a11,&a12,&a13);
  50   minmax6(&b8,&b9,&b10,&b11,&b12,&b13);
  51   a13 = tex2D(tex_img_ins, j-2, i  );
  52   b13 = tex2D(tex_img_ins, j+3, i  );
  53   minmax5(&a9,&a10,&a11,&a12,&a13);
  54   minmax5(&b9,&b10,&b11,&b12,&b13);
  55   a13 = tex2D(tex_img_ins, j-2, i+1);
  56   b13 = tex2D(tex_img_ins, j+3, i+1);
  57   minmax4(&a10,&a11,&a12,&a13);
  58   minmax4(&b10,&b11,&b12,&b13);
  59   a13 = tex2D(tex_img_ins, j-2, i+2);
  60   b13 = tex2D(tex_img_ins, j+3, i+2);
  61   minmax3(&a11,&a12,&a13);
  62   minmax3(&b11,&b12,&b13);
  63
  64   output[ __mul24(i, j_dim) +j ] = a12 ; //middle values
  65   output[ __mul24(i, j_dim) +j+1 ] = b12 ;
  66 }