From 36d33dd6c2b247339e634b408a71fee97f79ab1b Mon Sep 17 00:00:00 2001 From: unknown Date: Mon, 14 Apr 2014 16:10:11 -0700 Subject: [PATCH] Review comments --- modules/imgproc/src/opencl/bilateral.cl | 54 +++++++++---------------- modules/imgproc/src/smooth.cpp | 4 -- 2 files changed, 18 insertions(+), 40 deletions(-) diff --git a/modules/imgproc/src/opencl/bilateral.cl b/modules/imgproc/src/opencl/bilateral.cl index e22ebaa37..e4c0e7fe5 100644 --- a/modules/imgproc/src/opencl/bilateral.cl +++ b/modules/imgproc/src/opencl/bilateral.cl @@ -55,6 +55,7 @@ #endif //Read pixels as integers +// Intel Device - Read Pixels as floats __kernel void bilateral(__global const uchar * src, int src_step, int src_offset, __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, __constant float * space_weight, __constant int * space_ofs) @@ -69,51 +70,31 @@ __kernel void bilateral(__global const uchar * src, int src_step, int src_offset float_t sum = (float_t)(0.0f); float wsum = 0.0f; + #ifdef INTEL_DEVICE + float_t val0 = convert_float_t(loadpix(src + src_index)); + #else int_t val0 = convert_int_t(loadpix(src + src_index)); - + #endif #pragma unroll for (int k = 0; k < maxk; k++ ) { - int_t val = convert_int_t(loadpix(src + src_index + space_ofs[k])); - uint diff = (uint)SUM(abs(val - val0)); - float w = space_weight[k] * native_exp((float)(diff * diff * as_float(gauss_color_coeff))); + #ifdef INTEL_DEVICE + float_t val = convert_float_t(loadpix(src + src_index + space_ofs[k])); + float diff = SUM(fabs(val - val0)); + #else + int_t val = convert_int_t(loadpix(src + src_index + space_ofs[k])); + int diff = SUM(abs(val - val0)); + #endif + float w = space_weight[k] * native_exp((float)(diff * diff * gauss_color_coeff)); sum += convert_float_t(val) * (float_t)(w); wsum += w; } - - storepix(convert_uchar_t(sum / (float_t)(wsum)), dst + dst_index); - } -} - -//Read pixels as floats -__kernel void bilateral_float(__global const uchar * src, int src_step, int src_offset, - __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, - __constant float * space_weight, __constant int * space_ofs) -{ - int x = get_global_id(0); - int y = get_global_id(1); - - if (y < dst_rows && x < dst_cols) - { - int src_index = mad24(y + radius, src_step, mad24(x + radius, TSIZE, src_offset)); - int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset)); - - float_t sum = (float_t)(0.0f); - float wsum = 0.0f; - float_t val0 = convert_float_t(loadpix(src + src_index)); - - for (int k = 0; k < maxk; k++ ) - { - float_t val = convert_float_t(loadpix(src + src_index + space_ofs[k])); - float i = SUM(fabs(val - val0)); - float w = space_weight[k] * native_exp(i * i * as_float(gauss_color_coeff)); - sum += val * w; - wsum += w; - } storepix(convert_uchar_t(sum / (float_t)(wsum)), dst + dst_index); } } +#ifdef INTEL_DEVICE +#if cn == 1 //for single channgel x4 sized images. __kernel void bilateral_float4(__global const uchar * src, int src_step, int src_offset, __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, @@ -131,8 +112,7 @@ __kernel void bilateral_float4(__global const uchar * src, int src_step, int src for (int k = 0; k < maxk; k++ ) { float4 val = convert_float4(vload4(0, src + src_index + space_ofs[k])); - float spacew = space_weight[k]; - float4 w = spacew * native_exp((val - val0) * (val - val0) * as_float(gauss_color_coeff)); + float4 w = space_weight[k] * native_exp((val - val0) * (val - val0) * gauss_color_coeff); sum += val * w; wsum += w; } @@ -140,3 +120,5 @@ __kernel void bilateral_float4(__global const uchar * src, int src_step, int src vstore4(convert_uchar4_rtz(sum), 0, dst + dst_index); } } +#endif +#endif \ No newline at end of file diff --git a/modules/imgproc/src/smooth.cpp b/modules/imgproc/src/smooth.cpp index 0f41d472d..8b29a8ea3 100644 --- a/modules/imgproc/src/smooth.cpp +++ b/modules/imgproc/src/smooth.cpp @@ -2370,10 +2370,6 @@ static bool ocl_bilateralFilter_8u(InputArray _src, OutputArray _dst, int d, kernelName = "bilateral_float4"; sizeDiv = 4; } - else - { - kernelName = "bilateral_float"; - } } ocl::Kernel k(kernelName.c_str(), ocl::imgproc::bilateral_oclsrc, format("-D radius=%d -D maxk=%d -D cn=%d -D int_t=%s -D uint_t=uint%s -D convert_int_t=%s"