THESE/Chapters/chapter5/code/median3-2pix.cu

   1 __device__ inline void s(unsigned char * a, unsigned char * b)
   2 {
   3   unsigned char tmp ;
   4   if (*a > *b)
   5         {
   6           tmp = *b ;
   7           *b = *a ;
   8           *a = tmp ;
   9           }
  10 }
  11
  12 #define minmax3(a, b, c)          s(a, b); s(a, c); s(b,c)
  13 #define minmax4(a, b, c, d)       s(a, b); s(c, d); s(a, c); s(b, d);
  14 #define minmax5(a, b, c, d, e)    s(a, b); s(c, d); s(a, c); s(b, d); s(b,d); s(a,e); s(d,e)
  15 #define minmax6(a, b, c, d, e, f) s(a, b); s(c, d); s(e, f); s(a, c); s(d, f); s(a, d); s(c ,f)
  16
  17 __global__ void kernel_median3_2pix( unsigned short*output, int i_dim, int j_dim)
  18 {
  19   // coordonnées absolues du point
  20   int j = __mul24(__mul24(blockIdx.x,blockDim.x) + threadIdx.x,2) ;
  21   int i = __mul24(blockIdx.y,blockDim.y) + threadIdx.y ;
  22
  23   int a0, a1, a2, a3, a4, a5 ;   // les 2 listes de sélection
  24   int b0, b1, b2, b3, b4, b5 ;   //
  25
  26   a0 = tex2D(tex_img_ins, j  , i-1) ;   //
  27   a1 = tex2D(tex_img_ins, j+1, i-1) ;   //
  28   a2 = tex2D(tex_img_ins, j  , i  ) ;   // les 6 premiers éléments
  29   a3 = tex2D(tex_img_ins, j+1, i  ) ;   //
  30   a4 = tex2D(tex_img_ins, j  , i+1) ;   //
  31   a5 = tex2D(tex_img_ins, j+1, i+1) ;   //
  32
  33   minmax6(&a0, &a1, &a2, &a3, &a4, &a5) ; // identification des extrema
  34
  35   b0=a0; b1=a1; b2=a2; b3=a3; b4=a4; b5=a5; // séparation des listes
  36
  37   a5 = tex2D(tex_img_ins, j-1, i  ) ;  // ajout éléments suivants
  38   b5 = tex2D(tex_img_ins, j+2, i  ) ;  //
  39
  40   minmax5(&a1, &a2, &a3, &a4, &a5) ;   // identification des extrema
  41   minmax5(&b1, &b2, &b3, &b4, &b5) ;   //
  42
  43   a5 = tex2D(tex_img_ins, j-1, i-1) ;
  44   b5 = tex2D(tex_img_ins, j+2, i-1) ;
  45   minmax4(&a2, &a3, &a4, &a5) ;
  46   minmax4(&b2, &b3, &b4, &b5) ;
  47   a5 = tex2D(tex_img_ins, j-1, i+1) ;
  48   b5 = tex2D(tex_img_ins, j+2, i+1) ;
  49   minmax3(&a3, &a4, &a5) ;
  50   minmax3(&b3, &b4, &b5) ;
  51
  52   output[ __mul24(i, j_dim) +j ] = a4 ;  // écriture des 2 valeurs
  53   output[ __mul24(i, j_dim) +j+1 ] = b4 ;// en mémoire globale
  54
  55 }
  56