X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/book_gpu.git/blobdiff_plain/b0cfcc742771497c83313352b59170ead2f99f40..69fa38e834150067ce61e4b3e1412f35f05149b7:/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu?ds=sidebyside diff --git a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu old mode 100644 new mode 100755 index 5c79c82..f4ad2c6 --- a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu +++ b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu @@ -6,17 +6,14 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r) //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) ; - // coordonnees absolues du point 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) @@ -29,10 +26,8 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r) } __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++ )