X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/book_gpu.git/blobdiff_plain/1b27427d56ddc474598ea928b11c31bcaf3f1e12..1ac5b5a535d9154c4f080e94f2f9a49ab6e299b7:/BookGPU/Chapters/chapter3/code/medianGeneric.cu~ diff --git a/BookGPU/Chapters/chapter3/code/medianGeneric.cu~ b/BookGPU/Chapters/chapter3/code/medianGeneric.cu~ index 7280904..bf41842 100755 --- 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 ; }