diff --git a/modules/photo/src/fast_nlmeans_denoising_opencl.hpp b/modules/photo/src/fast_nlmeans_denoising_opencl.hpp index 1cdd8fa494..cd7dde3859 100644 --- a/modules/photo/src/fast_nlmeans_denoising_opencl.hpp +++ b/modules/photo/src/fast_nlmeans_denoising_opencl.hpp @@ -70,11 +70,11 @@ static bool ocl_calcAlmostDist2Weight(UMat & almostDist2Weight, int searchWindow static bool ocl_fastNlMeansDenoising(InputArray _src, OutputArray _dst, float h, int templateWindowSize, int searchWindowSize) { - int type = _src.type(), cn = CV_MAT_CN(type); + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int ctaSize = ocl::Device::getDefault().isIntel() ? CTA_SIZE_INTEL : CTA_SIZE_DEFAULT; Size size = _src.size(); - if ( type != CV_8UC1 && type != CV_8UC2 && type != CV_8UC4 ) + if ( type != CV_8UC1 && type != CV_8UC2 && type != CV_8UC3 ) return false; int templateWindowHalfWize = templateWindowSize / 2; @@ -86,13 +86,15 @@ static bool ocl_fastNlMeansDenoising(InputArray _src, OutputArray _dst, float h, char cvt[2][40]; String opts = format("-D OP_CALC_FASTNLMEANS -D TEMPLATE_SIZE=%d -D SEARCH_SIZE=%d" - " -D uchar_t=%s -D int_t=%s -D BLOCK_COLS=%d -D BLOCK_ROWS=%d" + " -D sample_t=%s -D pixel_t=%s -D int_t=%s" + " -D BLOCK_COLS=%d -D BLOCK_ROWS=%d" " -D CTA_SIZE=%d -D TEMPLATE_SIZE2=%d -D SEARCH_SIZE2=%d" - " -D convert_int_t=%s -D cn=%d -D convert_uchar_t=%s", - templateWindowSize, searchWindowSize, ocl::typeToStr(type), - ocl::typeToStr(CV_32SC(cn)), BLOCK_COLS, BLOCK_ROWS, ctaSize, - templateWindowHalfWize, searchWindowHalfSize, - ocl::convertTypeStr(CV_8U, CV_32S, cn, cvt[0]), cn, + " -D convert_int_t=%s -D cn=%d -D convert_pixel_t=%s", + templateWindowSize, searchWindowSize, + ocl::typeToStr(depth), ocl::typeToStr(type), ocl::typeToStr(CV_32SC(cn)), + BLOCK_COLS, BLOCK_ROWS, + ctaSize, templateWindowHalfWize, searchWindowHalfSize, + ocl::convertTypeStr(CV_8U, CV_32S, cn, cvt[0]), type == CV_8UC3 ? 4 : cn, ocl::convertTypeStr(CV_32S, CV_8U, cn, cvt[1])); ocl::Kernel k("fastNlMeansDenoising", ocl::photo::nlmeans_oclsrc, opts); @@ -107,10 +109,22 @@ static bool ocl_fastNlMeansDenoising(InputArray _src, OutputArray _dst, float h, UMat srcex; int borderSize = searchWindowHalfSize + templateWindowHalfWize; - copyMakeBorder(_src, srcex, borderSize, borderSize, borderSize, borderSize, BORDER_DEFAULT); + if (type == CV_8UC3) { + Mat src_rgb = _src.getMat(), src_rgba(size, CV_8UC4); + int from_to[] = { 0,0, 1,1, 2,2 }; + mixChannels(&src_rgb, 1, &src_rgba, 1, from_to, 3); + copyMakeBorder(src_rgba, srcex, + borderSize, borderSize, borderSize, borderSize, BORDER_DEFAULT); + } + else + copyMakeBorder(_src, srcex, borderSize, borderSize, borderSize, borderSize, BORDER_DEFAULT); _dst.create(size, type); - UMat dst = _dst.getUMat(); + UMat dst; + if (type == CV_8UC3) + dst.create(size, CV_8UC4); + else + dst = _dst.getUMat(); int searchWindowSizeSq = searchWindowSize * searchWindowSize; Size upColSumSize(size.width, searchWindowSizeSq * nblocksy); @@ -123,7 +137,15 @@ static bool ocl_fastNlMeansDenoising(InputArray _src, OutputArray _dst, float h, ocl::KernelArg::PtrReadOnly(buffer), almostTemplateWindowSizeSqBinShift); size_t globalsize[2] = { nblocksx * ctaSize, nblocksy }, localsize[2] = { ctaSize, 1 }; - return k.run(2, globalsize, localsize, false); + if (!k.run(2, globalsize, localsize, false)) return false; + + if (type == CV_8UC3) { + Mat dst_rgba = dst.getMat(ACCESS_READ), dst_rgb = _dst.getMat(); + int from_to[] = { 0,0, 1,1, 2,2 }; + mixChannels(&dst_rgba, 1, &dst_rgb, 1, from_to, 3); + } + + return true; } static bool ocl_fastNlMeansDenoisingColored( InputArray _src, OutputArray _dst, diff --git a/modules/photo/src/opencl/nlmeans.cl b/modules/photo/src/opencl/nlmeans.cl index af3fb1f9b0..c48adda0b2 100644 --- a/modules/photo/src/opencl/nlmeans.cl +++ b/modules/photo/src/opencl/nlmeans.cl @@ -29,8 +29,11 @@ __kernel void calcAlmostDist2Weight(__global int * almostDist2Weight, int almost if (almostDist < almostMaxDist) { FT dist = almostDist * almostDist2ActualDistMultiplier; +#ifdef ABS + int weight = convert_int_sat_rte(fixedPointMult * exp(-dist*dist * den)); +#else int weight = convert_int_sat_rte(fixedPointMult * exp(-dist * den)); - +#endif if (weight < WEIGHT_THRESHOLD * fixedPointMult) weight = 0; @@ -44,21 +47,33 @@ __kernel void calcAlmostDist2Weight(__global int * almostDist2Weight, int almost #define SEARCH_SIZE_SQ (SEARCH_SIZE * SEARCH_SIZE) -inline int calcDist(uchar_t a, uchar_t b) +inline int calcDist(pixel_t a, pixel_t b) { +#ifdef ABS + int_t retval = convert_int_t(abs_diff(a, b)); +#else int_t diff = convert_int_t(a) - convert_int_t(b); int_t retval = diff * diff; +#endif #if cn == 1 return retval; #elif cn == 2 return retval.x + retval.y; +#elif cn == 3 || cn == 4 /* A is ignored */ + return retval.x + retval.y + retval.z; #else -#error "cn should be either 1 or 2" +#error "cn should be either 1, 2, 3 or 4" #endif } -inline int calcDistUpDown(uchar_t down_value, uchar_t down_value_t, uchar_t up_value, uchar_t up_value_t) +#ifdef ABS +inline int calcDistUpDown(pixel_t down_value, pixel_t down_value_t, pixel_t up_value, pixel_t up_value_t) +{ + return calcDist(down_value, down_value_t) - calcDist(up_value, up_value_t); +} +#else +inline int calcDistUpDown(pixel_t down_value, pixel_t down_value_t, pixel_t up_value, pixel_t up_value_t) { int_t A = convert_int_t(down_value) - convert_int_t(down_value_t); int_t B = convert_int_t(up_value) - convert_int_t(up_value_t); @@ -68,14 +83,17 @@ inline int calcDistUpDown(uchar_t down_value, uchar_t down_value_t, uchar_t up_v return retval; #elif cn == 2 return retval.x + retval.y; +#elif cn == 3 || cn == 4 /* A is ignored */ + return retval.x + retval.y + retval.z; #else -#error "cn should be either 1 or 2" +#error "cn should be either 1, 2, 3 or 4" #endif } +#endif #define COND if (x == 0 && y == 0) -inline void calcFirstElementInRow(__global const uchar * src, int src_step, int src_offset, +inline void calcFirstElementInRow(__global const sample_t * src, int src_step, int src_offset, __local int * dists, int y, int x, int id, __global int * col_dists, __global int * up_col_dists) { @@ -87,9 +105,9 @@ inline void calcFirstElementInRow(__global const uchar * src, int src_step, int { int dist = 0, value; - __global const uchar_t * src_template = (__global const uchar_t *)(src + + __global const pixel_t * src_template = (__global const pixel_t *)(src + mad24(sy + i / SEARCH_SIZE, src_step, mad24(cn, sx + i % SEARCH_SIZE, src_offset))); - __global const uchar_t * src_current = (__global const uchar_t *)(src + mad24(y, src_step, mad24(cn, x, src_offset))); + __global const pixel_t * src_current = (__global const pixel_t *)(src + mad24(y, src_step, mad24(cn, x, src_offset))); __global int * col_dists_current = col_dists + i * TEMPLATE_SIZE; #pragma unroll @@ -107,8 +125,8 @@ inline void calcFirstElementInRow(__global const uchar * src, int src_step, int dist += value; } - src_current = (__global const uchar_t *)((__global const uchar *)src_current + src_step); - src_template = (__global const uchar_t *)((__global const uchar *)src_template + src_step); + src_current = (__global const pixel_t *)((__global const sample_t *)src_current + src_step); + src_template = (__global const pixel_t *)((__global const sample_t *)src_template + src_step); } #pragma unroll @@ -120,7 +138,7 @@ inline void calcFirstElementInRow(__global const uchar * src, int src_step, int } } -inline void calcElementInFirstRow(__global const uchar * src, int src_step, int src_offset, +inline void calcElementInFirstRow(__global const sample_t * src, int src_step, int src_offset, __local int * dists, int y, int x0, int x, int id, int first, __global int * col_dists, __global int * up_col_dists) { @@ -130,8 +148,8 @@ inline void calcElementInFirstRow(__global const uchar * src, int src_step, int for (int i = id; i < SEARCH_SIZE_SQ; i += CTA_SIZE) { - __global const uchar_t * src_current = (__global const uchar_t *)(src + mad24(y, src_step, mad24(cn, x, src_offset))); - __global const uchar_t * src_template = (__global const uchar_t *)(src + + __global const pixel_t * src_current = (__global const pixel_t *)(src + mad24(y, src_step, mad24(cn, x, src_offset))); + __global const pixel_t * src_template = (__global const pixel_t *)(src + mad24(sy + i / SEARCH_SIZE, src_step, mad24(cn, sx + i % SEARCH_SIZE, src_offset))); __global int * col_dists_current = col_dists + TEMPLATE_SIZE * i; @@ -142,8 +160,8 @@ inline void calcElementInFirstRow(__global const uchar * src, int src_step, int { col_dist += calcDist(src_current[0], src_template[0]); - src_current = (__global const uchar_t *)((__global const uchar *)src_current + src_step); - src_template = (__global const uchar_t *)((__global const uchar *)src_template + src_step); + src_current = (__global const pixel_t *)((__global const sample_t *)src_current + src_step); + src_template = (__global const pixel_t *)((__global const sample_t *)src_template + src_step); } dists[i] += col_dist - col_dists_current[first]; @@ -152,7 +170,7 @@ inline void calcElementInFirstRow(__global const uchar * src, int src_step, int } } -inline void calcElement(__global const uchar * src, int src_step, int src_offset, +inline void calcElement(__global const sample_t * src, int src_step, int src_offset, __local int * dists, int y, int x0, int x, int id, int first, __global int * col_dists, __global int * up_col_dists) { @@ -160,8 +178,8 @@ inline void calcElement(__global const uchar * src, int src_step, int src_offset int sy_up = y - TEMPLATE_SIZE2 - 1; int sy_down = y + TEMPLATE_SIZE2; - uchar_t up_value = *(__global const uchar_t *)(src + mad24(sy_up, src_step, mad24(cn, sx, src_offset))); - uchar_t down_value = *(__global const uchar_t *)(src + mad24(sy_down, src_step, mad24(cn, sx, src_offset))); + pixel_t up_value = *(__global const pixel_t *)(src + mad24(sy_up, src_step, mad24(cn, sx, src_offset))); + pixel_t down_value = *(__global const pixel_t *)(src + mad24(sy_down, src_step, mad24(cn, sx, src_offset))); sx -= SEARCH_SIZE2; sy_up -= SEARCH_SIZE2; @@ -171,8 +189,8 @@ inline void calcElement(__global const uchar * src, int src_step, int src_offset { int wx = i % SEARCH_SIZE, wy = i / SEARCH_SIZE; - uchar_t up_value_t = *(__global const uchar_t *)(src + mad24(sy_up + wy, src_step, mad24(cn, sx + wx, src_offset))); - uchar_t down_value_t = *(__global const uchar_t *)(src + mad24(sy_down + wy, src_step, mad24(cn, sx + wx, src_offset))); + pixel_t up_value_t = *(__global const pixel_t *)(src + mad24(sy_up + wy, src_step, mad24(cn, sx + wx, src_offset))); + pixel_t down_value_t = *(__global const pixel_t *)(src + mad24(sy_down + wy, src_step, mad24(cn, sx + wx, src_offset))); __global int * col_dists_current = col_dists + mad24(i, TEMPLATE_SIZE, first); __global int * up_col_dists_current = up_col_dists + mad24(x0, SEARCH_SIZE_SQ, i); @@ -185,9 +203,9 @@ inline void calcElement(__global const uchar * src, int src_step, int src_offset } } -inline void convolveWindow(__global const uchar * src, int src_step, int src_offset, +inline void convolveWindow(__global const sample_t * src, int src_step, int src_offset, __local int * dists, __global const int * almostDist2Weight, - __global uchar * dst, int dst_step, int dst_offset, + __global sample_t * dst, int dst_step, int dst_offset, int y, int x, int id, __local int * weights_local, __local int_t * weighted_sum_local, int almostTemplateWindowSizeSqBinShift) { @@ -197,7 +215,7 @@ inline void convolveWindow(__global const uchar * src, int src_step, int src_off for (int i = id; i < SEARCH_SIZE_SQ; i += CTA_SIZE) { int src_index = mad24(sy + i / SEARCH_SIZE, src_step, mad24(i % SEARCH_SIZE + sx, cn, src_offset)); - int_t src_value = convert_int_t(*(__global const uchar_t *)(src + src_index)); + int_t src_value = convert_int_t(*(__global const pixel_t *)(src + src_index)); int almostAvgDist = dists[i] >> almostTemplateWindowSizeSqBinShift; int weight = almostDist2Weight[almostAvgDist]; @@ -228,13 +246,13 @@ inline void convolveWindow(__global const uchar * src, int src_step, int src_off weighted_sum_local[2] + weighted_sum_local[3]; int weights_local_0 = weights_local[0] + weights_local[1] + weights_local[2] + weights_local[3]; - *(__global uchar_t *)(dst + dst_index) = convert_uchar_t(weighted_sum_local_0 / (int_t)(weights_local_0)); + *(__global pixel_t *)(dst + dst_index) = convert_pixel_t(weighted_sum_local_0 / (int_t)(weights_local_0)); } } -__kernel void fastNlMeansDenoising(__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, - __global const int * almostDist2Weight, __global uchar * buffer, +__kernel void fastNlMeansDenoising(__global const sample_t * src, int src_step, int src_offset, + __global sample_t * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, + __global const int * almostDist2Weight, __global sample_t * buffer, int almostTemplateWindowSizeSqBinShift) { int block_x = get_group_id(0), nblocks_x = get_num_groups(0);