BookGPU/Chapters/chapter3/code/kernMedianForgetGeneric.cu

   1 __global__ void kernel_medianForgetR( short *output, int i_dim, int j_dim, int r)
   2 {
   3   // coordonnees absolues du point
   4   int j = __mul24(blockIdx.x,blockDim.x) + threadIdx.x ;
   5   int i = __mul24(blockIdx.y,blockDim.y) + threadIdx.y ;
   6   short ic, jc ;
   7   short Nreg = ((2*r+1)*(2*r+1))/2 + 2 ;
   8
   9   // remplissage du vecteur de tri minmax
  10   short vect[8066] ;
  11   int Freg=Nreg ;
  12   for (ic=0; ic<2*r+1; ic++)
  13         {
  14           for (jc=0; jc<2*r+1; jc++)
  15                 {
  16                   if ( ic*(2*r+1)+jc < Nreg )
  17                         {
  18                           vect[ ic*(2*r+1)+jc ] = tex2D(tex_img_ins, j-r+jc, i-r+ic) ;
  19                         } else
  20                         {
  21                           minmaxN(vect, Freg--) ;
  22                           vect[ Nreg-1 ] = tex2D(tex_img_ins, j-r+jc, i-r+ic) ;
  23                         }
  24                 }
  25         }
  26   minmax3(&vect[Nreg-3], &vect[Nreg-2], &vect[Nreg-1])
  27
  28   //medRoi[ (threadIdx.y+ic)*(bdimX+L-1)+ (tidX+jc) ]
  29
  30         output[ __mul24(i, j_dim) +j ] = vect[ Nreg-2 ];
  31 }