X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/book_gpu.git/blobdiff_plain/b0cfcc742771497c83313352b59170ead2f99f40..ecd2ddac55172a779e7e26d3bd5b1b2cb95033d6:/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu?ds=inline diff --git a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu old mode 100644 new mode 100755 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) { - 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 ; - 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++ ) @@ -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 ; - inf = sup = egal = 0 ; + inf = sup = equal = 0 ; 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 ; - } else egal++ ; + } else equal++ ; } 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 ; - 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 ;