new ch5 reread

[book_gpu.git] / BookGPU / Chapters / chapter3 / code / medianGeneric.cu~

diff --git a/BookGPU/Chapters/chapter3/code/medianGeneric.cu~ b/BookGPU/Chapters/chapter3/code/medianGeneric.cu~
index 7280904e3f1093e15bd1a087730eb72a72be9abd..bf418424274c290615fff244069fb49ea51ed9b9 100755 (executable)
--- a/BookGPU/Chapters/chapter3/code/medianGeneric.cu~
+++ b/BookGPU/Chapters/chapter3/code/medianGeneric.cu~
@@ -1,11 +1,25 @@
-__device__ inline void s(int* a, int* b)
-{
+__global__ void kernel_medianR( short *output,
+                                int i_dim, int j_dim, int r)
+{   
+  // absolute coordinates of the center pixel
+  int j = __mul24(blockIdx.x,blockDim.x) + threadIdx.x ; 
+  int i = __mul24(blockIdx.y,blockDim.y) + threadIdx.y ;
   
-  int tmp ;
-  if (*a > *b)
+  short cpt, ic, jc ;
+  short histogram[256] ; // 8 bit image    
+  // zeroing histogram data
+  for (ic =0; ic<256; ic++) histogram[ ic ]=0 ; 
+  // histogram filling
+  for(ic=i-r; ic<=i+r; ic++ )
+       for(jc=j-r; jc<=j+r; jc++)
+         histogram[ tex2D(tex_img_ins, jc, ic) ]++ ;
+  // histogram parsing
+  cpt = 0 ;
+  for(ic=0; ic<256; ic++)
        {
-         tmp = *b ;
-         *b = *a ;
-         *a = tmp ;
-         }
+         cpt += histogram[ ic ] ;
+         // selection of 50% percentile 
+         if ( cpt > ((2*r+1)*(2*r+1))>>1 ) break ; 
+       } 
+  output[ __mul24(i, j_dim) +j ] =  ic ;  
 }