From: Vladimir Bystricky Date: Wed, 11 Dec 2013 10:26:33 +0000 (+0400) Subject: Fix problems with border extrapolation in kernel. Add Isolated/Nonisolated borders. X-Git-Tag: accepted/tizen/6.0/unified/20201030.111113~3640^2~1 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=e7227d3e4b4b0e816b17a57b7ade9e39030c18dc;p=platform%2Fupstream%2Fopencv.git Fix problems with border extrapolation in kernel. Add Isolated/Nonisolated borders. --- diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index 832b0d0..d548168 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3210,6 +3210,13 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, size_t localsize[2] = {0, 1}; ocl::Kernel kernel; + UMat src; Size wholeSize; + if (!isIsolatedBorder) + { + src = _src.getUMat(); + Point ofs; + src.locateROI(wholeSize, ofs); + } size_t maxWorkItemSizes[32]; device.maxWorkItemSizes(maxWorkItemSizes); size_t tryWorkItems = maxWorkItemSizes[0]; @@ -3233,8 +3240,8 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, int requiredLeft = (int)BLOCK_SIZE; // not this: anchor.x; int requiredBottom = ksize.height - 1 - anchor.y; int requiredRight = (int)BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x; - int h = sz.height; - int w = sz.width; + int h = isIsolatedBorder ? sz.height : wholeSize.height; + int w = isIsolatedBorder ? sz.width : wholeSize.width; bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight; if ((w < ksize.width) || (h < ksize.height)) @@ -3268,10 +3275,22 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, _dst.create(sz, CV_MAKETYPE(ddepth, cn)); UMat dst = _dst.getUMat(); - UMat src = _src.getUMat(); + if (src.empty()) + src = _src.getUMat(); int idxArg = 0; - idxArg = kernel.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(src)); + idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(src)); + idxArg = kernel.set(idxArg, (int)src.step); + + int srcOffsetX = (int)((src.offset % src.step) / src.elemSize()); + int srcOffsetY = (int)(src.offset / src.step); + int srcEndX = (isIsolatedBorder ? (srcOffsetX + sz.width) : wholeSize.width); + int srcEndY = (isIsolatedBorder ? (srcOffsetY + sz.height) : wholeSize.height); + idxArg = kernel.set(idxArg, srcOffsetX); + idxArg = kernel.set(idxArg, srcOffsetY); + idxArg = kernel.set(idxArg, srcEndX); + idxArg = kernel.set(idxArg, srcEndY); + idxArg = kernel.set(idxArg, ocl::KernelArg::WriteOnly(dst)); float borderValue[4] = {0, 0, 0, 0}; double borderValueDouble[4] = {0, 0, 0, 0}; diff --git a/modules/imgproc/src/opencl/filter2D.cl b/modules/imgproc/src/opencl/filter2D.cl index 1225be9..d360714 100644 --- a/modules/imgproc/src/opencl/filter2D.cl +++ b/modules/imgproc/src/opencl/filter2D.cl @@ -102,7 +102,7 @@ do \ { \ if (x < minX) \ - x = -(x - minX) - 1 + delta; \ + x = minX - (x - minX) - 1 + delta; \ else \ x = maxX - 1 - (x - maxX) - delta; \ } \ @@ -114,7 +114,7 @@ do \ { \ if (y < minY) \ - y = -(y - minY) - 1 + delta; \ + y = minY - (y - minY) - 1 + delta; \ else \ y = maxY - 1 - (y - maxY) - delta; \ } \ @@ -222,7 +222,7 @@ struct RectCoords #endif -inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, int srcstep, int srcoffset, const struct RectCoords srcCoords +inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, int srcstep, const struct RectCoords srcCoords #ifdef BORDER_CONSTANT , SCALAR_TYPE borderValue #endif @@ -235,7 +235,7 @@ inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, in #endif { //__global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes); - __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + srcoffset + pos.x * sizeof(TYPE)); + __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + pos.x * sizeof(TYPE)); return CONVERT_TO_FPTYPE(*ptr); } else @@ -262,7 +262,7 @@ inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, in if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) { //__global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes); - __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + srcoffset + pos.x * sizeof(TYPE)); + __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + pos.x * sizeof(TYPE)); return CONVERT_TO_FPTYPE(*ptr); } else @@ -279,8 +279,8 @@ inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, in __kernel __attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1))) -void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +void filter2D(__global const uchar* srcptr, int srcstep, int srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY, + __global uchar* dstptr, int dststep, int dstoffset, int rows, int cols, #ifdef BORDER_CONSTANT SCALAR_TYPE borderValue, @@ -288,8 +288,7 @@ void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, __constant FPTYPE* kernelData // transposed: [KERNEL_SIZE_X][KERNEL_SIZE_Y2_ALIGNED] ) { - const struct RectCoords srcCoords = {0, 0, cols, rows}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY - const struct RectCoords dstCoords = {0, 0, cols, rows}; + const struct RectCoords srcCoords = {srcOffsetX, srcOffsetY, srcEndX, srcEndY}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY const int local_id = get_local_id(0); const int x = local_id + (LOCAL_SIZE - (KERNEL_SIZE_X - 1)) * get_group_id(0) - ANCHOR_X; @@ -300,23 +299,23 @@ void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y); - int2 pos = (int2)(dstCoords.x1 + x, dstCoords.y1 + y); + int2 pos = (int2)(x, y); __global TYPE* dstPtr = (__global TYPE*)((__global char*)dstptr + pos.y * dststep + dstoffset + pos.x * sizeof(TYPE)); // Pointer can be out of bounds! - bool writeResult = (local_id >= ANCHOR_X && local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X) && - pos.x >= dstCoords.x1 && pos.x < dstCoords.x2); + bool writeResult = ((local_id >= ANCHOR_X) && (local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X)) && + (pos.x >= 0) && (pos.x < cols)); #if BLOCK_SIZE_Y > 1 bool readAllpixels = true; int sy_index = 0; // current index in data[] array - dstCoords.y2 = min(dstCoords.y2, pos.y + BLOCK_SIZE_Y); + dstRowsMax = min(rows, pos.y + BLOCK_SIZE_Y); for (; - pos.y < dstCoords.y2; + pos.y < dstRowsMax; pos.y++, dstPtr = (__global TYPE*)((__global char*)dstptr + dststep)) #endif { - ASSERT(pos.y < dstCoords.y2); + ASSERT(pos.y < dstRowsMax); for ( #if BLOCK_SIZE_Y > 1 @@ -326,7 +325,7 @@ void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, #endif sy++, srcPos.y++) { - data[sy + sy_index] = readSrcPixel(srcPos, srcptr, srcstep, srcoffset, srcCoords + data[sy + sy_index] = readSrcPixel(srcPos, srcptr, srcstep, srcCoords #ifdef BORDER_CONSTANT , borderValue #endif @@ -361,7 +360,6 @@ void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, if (writeResult) { - ASSERT(pos.y >= dstCoords.y1 && pos.y < dstCoords.y2); *dstPtr = CONVERT_TO_TYPE(total_sum); } diff --git a/modules/imgproc/test/ocl/test_filter2d.cpp b/modules/imgproc/test/ocl/test_filter2d.cpp index ca14509..484080d 100644 --- a/modules/imgproc/test/ocl/test_filter2d.cpp +++ b/modules/imgproc/test/ocl/test_filter2d.cpp @@ -49,17 +49,11 @@ namespace cvtest { namespace ocl { -enum -{ - noType = -1, -}; - - ///////////////////////////////////////////////////////////////////////////////////////////////// // Filter2D -PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool) +PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool, bool) { - static const int kernelMinSize = 1; + static const int kernelMinSize = 2; static const int kernelMaxSize = 10; int type; @@ -75,8 +69,8 @@ PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool) virtual void SetUp() { type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1)); - borderType = GET_PARAM(2); - useRoi = GET_PARAM(3); + borderType = GET_PARAM(2) | (GET_PARAM(3) ? BORDER_ISOLATED : 0); + useRoi = GET_PARAM(4); } void random_roi() @@ -84,16 +78,19 @@ PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool) dsize = randomSize(1, MAX_VALUE); Size ksize = randomSize(kernelMinSize, kernelMaxSize); - kernel = randomMat(ksize, CV_MAKE_TYPE(((CV_64F == CV_MAT_DEPTH(type)) ? CV_64F : CV_32F), 1), -MAX_VALUE, MAX_VALUE); + Mat temp = randomMat(ksize, CV_MAKE_TYPE(((CV_64F == CV_MAT_DEPTH(type)) ? CV_64F : CV_32F), 1), -MAX_VALUE, MAX_VALUE); + cv::normalize(temp, kernel, 1.0, 0.0, NORM_L1); - Size roiSize = randomSize(1, MAX_VALUE); + //Size roiSize = randomSize(ksize.width, MAX_VALUE, ksize.height, MAX_VALUE); + Size roiSize(1024, 1024); Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); randomSubMat(dst, dst_roi, dsize, dstBorder, type, -MAX_VALUE, MAX_VALUE); - anchor.x = anchor.y = -1; + anchor.x = randomInt(-1, ksize.width); + anchor.y = randomInt(-1, ksize.height); UMAT_UPLOAD_INPUT_PARAMETER(src) UMAT_UPLOAD_OUTPUT_PARAMETER(dst) @@ -128,7 +125,9 @@ OCL_INSTANTIATE_TEST_CASE_P(ImageProc, Filter2D, (BorderType)BORDER_REPLICATE, (BorderType)BORDER_REFLECT, (BorderType)BORDER_REFLECT_101), - Bool()) + Bool(), // BORDER_ISOLATED + Bool() // ROI + ) );