add ch17

[book_gpu.git] / BookGPU / Chapters / chapter3 / code / kernMedianSeparable.cu

diff --git a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu
old mode 100644 (file)
new mode 100755 (executable)
index 5c79c82..bdf7023
--- a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu
+++ b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu
@@ -1,22 +1,19 @@
 __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r)
 {
   
 __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r)
 {
   

-  int idc, val, min, max, inf, egal, sup, mxinf, minsup, estim ;
+  int idc, val, min, max, inf, equal, sup, mxinf, minsup, estim ;

 
   //coordinates in the block
   int ib = threadIdx.y ;
   int jb = threadIdx.x ;
 
   //coordinates in the block
   int ib = threadIdx.y ;
   int jb = threadIdx.x ;

-  int idx_h = __mul24(ib+r,blockDim.x) + jb ;   // index pixel deans shmem (bloc+halo)
+  int idx_h = __mul24(ib+r,blockDim.x) +jb; // base pixel index

   int offset = __mul24(blockDim.x,r) ;
   
   int offset = __mul24(blockDim.x,r) ;
   

-  // coordonnees absolues du point

   int j = __mul24(blockIdx.x,blockDim.x) + jb ; 
   int i = __mul24(blockIdx.y,blockDim.y) + ib ;
   
   int j = __mul24(blockIdx.x,blockDim.x) + jb ; 
   int i = __mul24(blockIdx.y,blockDim.y) + ib ;
   

-  extern __shared__ int buff[] ;
-  /***********************************************************************************
-   *              CHARGEMENT DATA EN SHARED MEM
-   ***********************************************************************************/
+  //      DATA PREFETCHING INTO SHARED MEM
+  extern __shared__ int buff[] ;              

   buff[ idx_h ] = tex2D(tex_img_ins, j, i) ;
                                  
   if (ib < r)
   buff[ idx_h ] = tex2D(tex_img_ins, j, i) ;
                                  
   if (ib < r)


@@ -29,10 +26,8 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r)
          }
   
   __syncthreads() ;
          }
   
   __syncthreads() ;

-  /**********************************************************************************************
-   *               TRI VERTICAL par algo TORBEN MOGENSEN
-   *          (a little bit slow but saves memory => faster !)
-   **********************************************************************************************/
+
+  //      TORBEN MOGENSEN SORTING

   min = max = buff[ ib*blockDim.x +jb] ;
   
   for (idc= 0 ; idc< 2*r+1 ; idc++ )
   min = max = buff[ ib*blockDim.x +jb] ;
   
   for (idc= 0 ; idc< 2*r+1 ; idc++ )


@@ -45,7 +40,7 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r)
   while (1)
        {  
          estim = (min+max)/2 ;
   while (1)
        {  
          estim = (min+max)/2 ;

-         inf = sup = egal = 0  ;
+         inf = sup = equal = 0  ;

          mxinf = min ;
          minsup= max ;
          for (idc =0; idc< 2*r+1 ; idc++)
          mxinf = min ;
          minsup= max ;
          for (idc =0; idc< 2*r+1 ; idc++)


@@ -59,7 +54,7 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r)
                        {
                          sup++;
                          if( val < minsup) minsup = val ;
                        {
                          sup++;
                          if( val < minsup) minsup = val ;

-                       } else egal++ ;
+                       } else equal++ ;

                }
          if ( (inf <= (r+1))&&(sup <=(r+1)) ) break ;
          else if (inf>sup) max = mxinf ;
                }
          if ( (inf <= (r+1))&&(sup <=(r+1)) ) break ;
          else if (inf>sup) max = mxinf ;


@@ -67,7 +62,7 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r)
          }
   
   if ( inf >= r+1 ) val = mxinf ;
          }
   
   if ( inf >= r+1 ) val = mxinf ;

-  else if (inf+egal >= r+1) val = estim ;
+  else if (inf+equal >= r+1) val = estim ;

   else val = minsup ;
   
   output[ __mul24(j, i_dim)  +i  ] =  val ; 
   else val = minsup ;
   
   output[ __mul24(j, i_dim)  +i  ] =  val ;