diff --git a/modules/imgproc/src/opencl/medianFilter.cl b/modules/imgproc/src/opencl/medianFilter.cl index 6aed88ef70..c1ab045455 100644 --- a/modules/imgproc/src/opencl/medianFilter.cl +++ b/modules/imgproc/src/opencl/medianFilter.cl @@ -29,52 +29,52 @@ // or tort (including negligence or otherwise) arising in any way out of // the use of this software, even if advised of the possibility of such damage. -#define DATA_TYPE type - -#define scnbytes ((int)sizeof(type)) - -#define op(a,b) { mid=a; a=min(a,b); b=max(mid,b);} - -__kernel void medianFilter3(__global const uchar* srcptr, int srcStep, int srcOffset, - __global uchar* dstptr, int dstStep, int dstOffset, - int rows, int cols) +#if cn != 3 +#define loadpix(addr) *(__global const T *)(addr) +#define storepix(val, addr) *(__global T *)(addr) = val +#define TSIZE (int)sizeof(T) +#else +#define loadpix(addr) vload3(0, (__global const T1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr)) +#define TSIZE (int)sizeof(T1) * cn +#endif + +#define op(a, b) { mid = a; a = min(a, b); b = max(mid, b); } + +__kernel void medianFilter3(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols) { - __local DATA_TYPE data[18][18]; + __local T data[18][18]; int x = get_local_id(0); int y = get_local_id(1); - int gx= get_global_id(0); - int gy= get_global_id(1); + int gx = get_global_id(0); + int gy = get_global_id(1); int dx = gx - x - 1; int dy = gy - y - 1; - const int id = min((int)(x*16+y), 9*18-1); + int id = min(mad24(x, 16, y), 9*18-1); int dr = id / 18; int dc = id % 18; - int c = clamp(dx+dc, 0, cols-1); - - int r = clamp(dy+dr, 0, rows-1); - int index1 = mad24(r, srcStep, srcOffset + c*scnbytes); - - r = clamp(dy+dr+9, 0, rows-1); - int index9 = mad24(r, srcStep, srcOffset + c*scnbytes); + int c = clamp(dx + dc, 0, dst_cols - 1); - __global DATA_TYPE * src = (__global DATA_TYPE *)(srcptr + index1); - data[dr][dc] = src[0]; - - src = (__global DATA_TYPE *)(srcptr + index9); - data[dr+9][dc] = src[0]; + int r = clamp(dy + dr, 0, dst_rows - 1); + int index1 = mad24(r, src_step, mad24(c, TSIZE, src_offset)); + r = clamp(dy + dr + 9, 0, dst_rows - 1); + int index9 = mad24(r, src_step, mad24(c, TSIZE, src_offset)); + data[dr][dc] = loadpix(srcptr + index1); + data[dr+9][dc] = loadpix(srcptr + index9); barrier(CLK_LOCAL_MEM_FENCE); - DATA_TYPE p0=data[y][x], p1=data[y][(x+1)], p2=data[y][(x+2)]; - DATA_TYPE p3=data[y+1][x], p4=data[y+1][(x+1)], p5=data[y+1][(x+2)]; - DATA_TYPE p6=data[y+2][x], p7=data[y+2][(x+1)], p8=data[y+2][(x+2)]; - DATA_TYPE mid; + T p0 = data[y][x], p1 = data[y][(x+1)], p2 = data[y][(x+2)]; + T p3 = data[y+1][x], p4 = data[y+1][(x+1)], p5 = data[y+1][(x+2)]; + T p6 = data[y+2][x], p7 = data[y+2][(x+1)], p8 = data[y+2][(x+2)]; + T mid; op(p1, p2); op(p4, p5); op(p7, p8); op(p0, p1); op(p3, p4); op(p6, p7); op(p1, p2); op(p4, p5); @@ -82,56 +82,48 @@ __kernel void medianFilter3(__global const uchar* srcptr, int srcStep, int srcOf op(p3, p6); op(p1, p4); op(p2, p5); op(p4, p7); op(p4, p2); op(p6, p4); op(p4, p2); - int dst_index = mad24( gy, dstStep, dstOffset + gx * scnbytes); + int dst_index = mad24( gy, dst_step, mad24(gx, TSIZE, dst_offset)); - if( gy < rows && gx < cols) - { - __global DATA_TYPE* dst = (__global DATA_TYPE *)(dstptr + dst_index); - dst[0] = p4; - } + if (gy < dst_rows && gx < dst_cols) + storepix(p4, dstptr + dst_index); } -__kernel void medianFilter5(__global const uchar* srcptr, int srcStep, int srcOffset, - __global uchar* dstptr, int dstStep, int dstOffset, - int rows, int cols) +__kernel void medianFilter5(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols) { - __local DATA_TYPE data[20][20]; + __local T data[20][20]; - int x =get_local_id(0); - int y =get_local_id(1); + int x = get_local_id(0); + int y = get_local_id(1); - int gx=get_global_id(0); - int gy=get_global_id(1); + int gx = get_global_id(0); + int gy = get_global_id(1); int dx = gx - x - 2; int dy = gy - y - 2; - const int id = min((int)(x*16+y), 10*20-1); - - int dr=id/20; - int dc=id%20; - - int c=clamp(dx+dc, 0, cols-1); + int id = min(mad24(x, 16, y), 10*20-1); - int r = clamp(dy+dr, 0, rows-1); - int index1 = mad24(r, srcStep, srcOffset + c*scnbytes); + int dr = id / 20; + int dc = id % 20; - r = clamp(dy+dr+10, 0, rows-1); - int index10 = mad24(r, srcStep, srcOffset + c*scnbytes); + int c = clamp(dx + dc, 0, dst_cols - 1); + int r = clamp(dy + dr, 0, dst_rows - 1); + int index1 = mad24(r, src_step, mad24(c, TSIZE, src_offset)); - __global DATA_TYPE * src = (__global DATA_TYPE *)(srcptr + index1); - data[dr][dc] = src[0]; - src = (__global DATA_TYPE *)(srcptr + index10); - data[dr+10][dc] = src[0]; + r = clamp(dy + dr + 10, 0, dst_rows - 1); + int index10 = mad24(r, src_step, mad24(c, TSIZE, src_offset)); + data[dr][dc] = loadpix(srcptr + index1); + data[dr+10][dc] = loadpix(srcptr + index10); barrier(CLK_LOCAL_MEM_FENCE); - DATA_TYPE p0=data[y][x], p1=data[y][x+1], p2=data[y][x+2], p3=data[y][x+3], p4=data[y][x+4]; - DATA_TYPE p5=data[y+1][x], p6=data[y+1][x+1], p7=data[y+1][x+2], p8=data[y+1][x+3], p9=data[y+1][x+4]; - DATA_TYPE p10=data[y+2][x], p11=data[y+2][x+1], p12=data[y+2][x+2], p13=data[y+2][x+3], p14=data[y+2][x+4]; - DATA_TYPE p15=data[y+3][x], p16=data[y+3][x+1], p17=data[y+3][x+2], p18=data[y+3][x+3], p19=data[y+3][x+4]; - DATA_TYPE p20=data[y+4][x], p21=data[y+4][x+1], p22=data[y+4][x+2], p23=data[y+4][x+3], p24=data[y+4][x+4]; - DATA_TYPE mid; + T p0 = data[y][x], p1 = data[y][x+1], p2 = data[y][x+2], p3 = data[y][x+3], p4 = data[y][x+4]; + T p5 = data[y+1][x], p6 = data[y+1][x+1], p7 = data[y+1][x+2], p8 = data[y+1][x+3], p9 = data[y+1][x+4]; + T p10 = data[y+2][x], p11 = data[y+2][x+1], p12 = data[y+2][x+2], p13 = data[y+2][x+3], p14 = data[y+2][x+4]; + T p15 = data[y+3][x], p16 = data[y+3][x+1], p17 = data[y+3][x+2], p18 = data[y+3][x+3], p19 = data[y+3][x+4]; + T p20 = data[y+4][x], p21 = data[y+4][x+1], p22 = data[y+4][x+2], p23 = data[y+4][x+3], p24 = data[y+4][x+4]; + T mid; op(p1, p2); op(p0, p1); op(p1, p2); op(p4, p5); op(p3, p4); op(p4, p5); op(p0, p3); op(p2, p5); op(p2, p3); op(p1, p4); @@ -157,11 +149,8 @@ __kernel void medianFilter5(__global const uchar* srcptr, int srcStep, int srcOf op(p13, p17); op(p3, p15); op(p11, p23); op(p11, p15); op(p7, p19); op(p7, p11); op(p11, p13); op(p11, p12); - int dst_index = mad24( gy, dstStep, dstOffset + gx * scnbytes); + int dst_index = mad24(gy, dst_step, mad24(gx, TSIZE, dst_offset)); - if( gy < rows && gx < cols) - { - __global DATA_TYPE* dst = (__global DATA_TYPE *)(dstptr + dst_index); - dst[0] = p12; - } + if (gy < dst_rows && gx < dst_cols) + storepix(p12, dstptr + dst_index); } \ No newline at end of file diff --git a/modules/imgproc/src/smooth.cpp b/modules/imgproc/src/smooth.cpp index b7a5d81e5d..6a18af5c2d 100644 --- a/modules/imgproc/src/smooth.cpp +++ b/modules/imgproc/src/smooth.cpp @@ -1904,35 +1904,27 @@ medianBlur_SortNet( const Mat& _src, Mat& _dst, int m ) #ifdef HAVE_OPENCL -static bool ocl_medianFilter ( InputArray _src, OutputArray _dst, int m) +static bool ocl_medianFilter(InputArray _src, OutputArray _dst, int m) { - int type = _src.type(); - int depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); - - if (!((depth == CV_8U || depth == CV_16U || depth == CV_16S || depth == CV_32F) && (cn != 3 && cn <= 4))) - return false; - - const char * kernelName; + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); - if (m == 3) - kernelName = "medianFilter3"; - else if (m == 5) - kernelName = "medianFilter5"; - else + if ( !((depth == CV_8U || depth == CV_16U || depth == CV_16S || depth == CV_32F) && cn <= 4 && (m == 3 || m == 5)) ) return false; - ocl::Kernel k(kernelName,ocl::imgproc::medianFilter_oclsrc,format("-D type=%s",ocl::typeToStr(type))); + ocl::Kernel k(format("medianFilter%d", m).c_str(), ocl::imgproc::medianFilter_oclsrc, + format("-D T=%s -D T1=%s -D cn=%d", ocl::typeToStr(type), + ocl::typeToStr(depth), cn)); if (k.empty()) return false; UMat src = _src.getUMat(); - _dst.create(_src.size(),type); + _dst.create(src.size(), type); UMat dst = _dst.getUMat(); - size_t globalsize[2] = {(src.cols + 18) / 16 * 16, (src.rows + 15) / 16 * 16}; - size_t localsize[2] = {16, 16}; + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst)); - return k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst)).run(2,globalsize,localsize,false); + size_t globalsize[2] = { (src.cols + 18) / 16 * 16, (src.rows + 15) / 16 * 16}, localsize[2] = { 16, 16 }; + return k.run(2, globalsize, localsize, false); } #endif diff --git a/modules/imgproc/test/ocl/test_medianfilter.cpp b/modules/imgproc/test/ocl/test_medianfilter.cpp index 63690872a5..6015ed1d77 100644 --- a/modules/imgproc/test/ocl/test_medianfilter.cpp +++ b/modules/imgproc/test/ocl/test_medianfilter.cpp @@ -102,7 +102,7 @@ OCL_TEST_P(MedianFilter, Mat) OCL_INSTANTIATE_TEST_CASE_P(ImageProc, MedianFilter, Combine( Values(CV_8U, CV_16U, CV_16S, CV_32F), - Values(1, 2, 4), + OCL_ALL_CHANNELS, Values(3, 5), Bool()) );