__device__ inline void s(int* a, int* b)
{
  
  int tmp ;
  if (*a > *b)
	{
	  tmp = *b ;
	  *b = *a ;
	  *a = tmp ;
	  }
}

#define min3(a, b, c) s(a, b); s(a, c);
#define max3(a, b, c) s(b, c); s(a, c);
#define minmax3(a, b, c) max3(a, b, c); s(a, b);
#define minmax4(a, b, c, d) s(a, b); s(c, d); s(a, c); s(b, d);
#define minmax5(a, b, c, d, e) s(a, b); s(c, d); min3(a, c, e); max3(b, d, e);
#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);

__global__ void kernel_median3( short *output, int i_dim, int j_dim)
{  
  
  // coordonnees absolues du point
  int j = __mul24(blockIdx.x,blockDim.x) + threadIdx.x ; 
  int i = __mul24(blockIdx.y,blockDim.y) + threadIdx.y ;
  
  /**************************************************************************
   *             tri(s)
   **************************************************************************/
  int a0, a1, a2, a3, a4, a5 ;

  /********************************************************************************
   * les six premieres valeurs (suffisant pour median 3x3 par forgetfull selection)
   ********************************************************************************/
  a0 = tex2D(tex_img_ins, j-1, i-1) ;
  a1 = tex2D(tex_img_ins, j, i-1) ;
  a2 = tex2D(tex_img_ins, j+1, i-1) ;
  a3 = tex2D(tex_img_ins, j-1, i) ;
  a4 = tex2D(tex_img_ins, j, i) ;
  a5 = tex2D(tex_img_ins, j+1, i) ;

  
  //min max aux extremites
  minmax6(&a0, &a1, &a2, &a3, &a4, &a5) ;

  /********************************************
   * les deux valeurs suivantes aux extremites 
   ********************************************/
  a5 = tex2D(tex_img_in, j-1, i+1) ;

  minmax5(&a1, &a2, &a3, &a4, &a5) ;

  /********************************************
   * la derniere valeur a la fin
   ********************************************/
  
  a5 = tex2D(tex_img_ins, j, i+1) ;

  minmax4(&a2, &a3, &a4, &a5) ;

  a5 = tex2D(tex_img_ins, j+1, i+1) ;
  minmax3(&a3, &a4, &a5) ;
  
  
  //median au milieu !
  output[ __mul24(i, j_dim) +j ] = a4 ;
 
}