paper_fast_median/code/kernMedianForget1pix3.cu

   1 __device__ inline void s(int* a, int* b)
   2 {
   3   int tmp ;
   4   if (*a > *b)  { tmp = *b; *b = *a; *a = tmp;}
   5 }
   6
   7 #define min3(a, b, c) s(a, b); s(a, c);
   8 #define max3(a, b, c) s(b, c); s(a, c);
   9 #define minmax3(a, b, c) max3(a, b, c); s(a, b);
  10 #define minmax4(a, b, c, d) s(a, b); s(c, d); s(a, c); s(b, d);
  11 #define minmax5(a, b, c, d, e) s(a, b); s(c, d); min3(a, c, e); max3(b, d, e);
  12 #define minmax6(a, b, c, d, e, f) s(a,d); s(b, e); s(c, f); min3(a, b, c); max3(d, e, f);
  13
  14 __global__ void kernel_medianForget1pix3( short *output,
  15                                                                                   int i_dim, int j_dim)
  16 {
  17   int j = __mul24(blockIdx.x,blockDim.x) + threadIdx.x ;
  18   int i = __mul24(blockIdx.y,blockDim.y) + threadIdx.y ;
  19   int a0, a1, a2, a3, a4, a5 ;
  20
  21   a0 = tex2D(tex_img_ins, j-1, i-1) ; // first 6 values
  22   a1 = tex2D(tex_img_ins, j, i-1) ;
  23   a2 = tex2D(tex_img_ins, j+1, i-1) ;
  24   a3 = tex2D(tex_img_ins, j-1, i) ;
  25   a4 = tex2D(tex_img_ins, j, i) ;
  26   a5 = tex2D(tex_img_ins, j+1, i) ;
  27
  28   minmax6(&a0, &a1, &a2, &a3, &a4, &a5);//min->a0 max->a5
  29   a5 = tex2D(tex_img_in, j-1, i+1) ;    //next value in a5
  30   minmax5(&a1, &a2, &a3, &a4, &a5) ;    //min->a1 max->a5
  31   a5 = tex2D(tex_img_ins, j, i+1) ;     //next value in a5
  32   minmax4(&a2, &a3, &a4, &a5) ;         //min->a1 max->a5
  33   a5 = tex2D(tex_img_ins, j+1, i+1) ;   //next value in a5
  34   minmax3(&a3, &a4, &a5) ;              //min->a1 max->a5
  35
  36   output[ __mul24(i, j_dim) +j ] = a4 ; //middle value
  37
  38 }