diff --git a/modules/ocl/doc/image_filtering.rst b/modules/ocl/doc/image_filtering.rst index bd929b988..cbec29b11 100644 --- a/modules/ocl/doc/image_filtering.rst +++ b/modules/ocl/doc/image_filtering.rst @@ -133,7 +133,7 @@ Creates a normalized 2D box filter. .. ocv:function:: Ptr ocl::getBoxFilter_GPU(int srcType, int dstType, const Size &ksize, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT) - :param srcType: Input image type supporting ``CV_8UC1`` and ``CV_8UC4`` . + :param srcType: Input image type. :param dstType: Output image type. It supports only the same values as the source type. @@ -141,9 +141,7 @@ Creates a normalized 2D box filter. :param anchor: Anchor point. The default value ``Point(-1, -1)`` means that the anchor is at the kernel center. - :param borderType: Supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101,BORDER_WRAP. - -.. note:: This filter does not check out-of-border accesses, so only a proper sub-matrix of a bigger matrix has to be passed to it. + :param borderType: Border type. .. seealso:: :ocv:func:`boxFilter` @@ -153,21 +151,19 @@ Smooths the image using the normalized box filter. .. ocv:function:: void ocl::boxFilter(const oclMat &src, oclMat &dst, int ddepth, Size ksize, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT) - :param src: Input image. ``CV_8UC1`` and ``CV_8UC4`` source types are supported. + :param src: Input image. :param dst: Output image type. The size and type is the same as ``src`` . - :param ddepth: Output image depth. If -1, the output image has the same depth as the input one. The only values allowed here are ``CV_8U`` and -1. + :param ddepth: Desired depth of the destination image. If it is negative, it is the same as ``src.depth()`` . It supports only the same depth as the source image depth. :param ksize: Kernel size. :param anchor: Anchor point. The default value ``Point(-1, -1)`` means that the anchor is at the kernel center. - :param borderType: Supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101,BORDER_WRAP. + :param borderType: Border type. -Smoothes image using box filter.Supports data type: CV_8UC1, CV_8UC4, CV_32FC1 and CV_32FC4. - -.. note:: This filter does not check out-of-border accesses, so only a proper sub-matrix of a bigger matrix has to be passed to it. +Smoothes image using box filter. ocl::blur ------------- @@ -175,7 +171,7 @@ Acts as a synonym for the normalized box filter. .. ocv:function:: void ocl::blur(const oclMat &src, oclMat &dst, Size ksize, Point anchor = Point(-1, -1), int borderType = BORDER_CONSTANT) - :param src: Input image. ``CV_8UC1`` and ``CV_8UC4`` source types are supported. + :param src: Input image. :param dst: Output image type with the same size and type as ``src`` . @@ -183,9 +179,7 @@ Acts as a synonym for the normalized box filter. :param anchor: Anchor point. The default value Point(-1, -1) means that the anchor is at the kernel center. - :param borderType: Supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101,BORDER_WRAP. - -.. note:: This filter does not check out-of-border accesses, so only a proper sub-matrix of a bigger matrix has to be passed to it. + :param borderType: Border type. .. seealso:: :ocv:func:`blur`, :ocv:func:`ocl::boxFilter` @@ -217,11 +211,11 @@ Creates a non-separable linear filter. .. ocv:function:: Ptr ocl::createLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT) - :param srcType: Input image type. Supports ``CV_8U`` , ``CV_16U`` and ``CV_32F`` one and four channel image. + :param srcType: Input image type.. :param dstType: Output image type. The same type as ``src`` is supported. - :param kernel: 2D array of filter coefficients. Floating-point coefficients will be converted to fixed-point representation before the actual processing. Supports size up to 16. For larger kernels use :ocv:func:`ocl::convolve`. + :param kernel: 2D array of filter coefficients. :param anchor: Anchor point. The default value Point(-1, -1) means that the anchor is at the kernel center. @@ -234,9 +228,9 @@ ocl::filter2D ----------------- Applies the non-separable 2D linear filter to an image. -.. ocv:function:: void ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT) +.. ocv:function:: void ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor = Point(-1, -1), double delta = 0.0, int borderType = BORDER_DEFAULT) - :param src: Source image. Supports ``CV_8U`` , ``CV_16U`` and ``CV_32F`` one and four channel image. + :param src: Source image. :param dst: Destination image. The size and the number of channels is the same as ``src`` . @@ -246,9 +240,9 @@ Applies the non-separable 2D linear filter to an image. :param anchor: Anchor of the kernel that indicates the relative position of a filtered point within the kernel. The anchor resides within the kernel. The special default value (-1,-1) means that the anchor is at the kernel center. - :param borderType: Pixel extrapolation method. For details, see :ocv:func:`borderInterpolate` . + :param delta: optional value added to the filtered pixels before storing them in ``dst``. Value '0' is supported only. - :param stream: Stream for the asynchronous version. + :param borderType: Pixel extrapolation method. For details, see :ocv:func:`borderInterpolate` . ocl::getLinearRowFilter_GPU ------------------------------- @@ -447,7 +441,7 @@ ocl::Laplacian ------------------ Returns void -.. ocv:function:: void ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1) +.. ocv:function:: void ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1, double delta = 0, int borderType = BORDER_DEFAULT) :param src: The source image @@ -459,6 +453,10 @@ Returns void :param scale: The optional scale factor for the computed Laplacian values (by default, no scaling is applied + :param delta: Optional delta value that is added to the results prior to storing them in ``dst`` . Supported value is 0 only. + + :param bordertype: Pixel extrapolation method. + The function calculates the Laplacian of the source image by adding up the second x and y derivatives calculated using the Sobel operator. ocl::convolve diff --git a/modules/ocl/include/opencv2/ocl/ocl.hpp b/modules/ocl/include/opencv2/ocl/ocl.hpp index 5ccab64cb..db386952a 100644 --- a/modules/ocl/include/opencv2/ocl/ocl.hpp +++ b/modules/ocl/include/opencv2/ocl/ocl.hpp @@ -718,11 +718,12 @@ namespace cv CV_EXPORTS Ptr createDerivFilter_GPU( int srcType, int dstType, int dx, int dy, int ksize, int borderType = BORDER_DEFAULT ); //! applies Laplacian operator to the image - // supports only ksize = 1 and ksize = 3 8UC1 8UC4 32FC1 32FC4 data type - CV_EXPORTS void Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1); + // supports only ksize = 1 and ksize = 3 + CV_EXPORTS void Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize = 1, double scale = 1, + double delta=0, int borderType=BORDER_DEFAULT); //! returns 2D box filter - // supports CV_8UC1 and CV_8UC4 source type, dst type must be the same as source type + // dst type must be the same as source type CV_EXPORTS Ptr getBoxFilter_GPU(int srcType, int dstType, const Size &ksize, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT); @@ -731,17 +732,16 @@ namespace cv const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT); //! returns 2D filter with the specified kernel - // supports CV_8UC1 and CV_8UC4 types + // supports: dst type must be the same as source type CV_EXPORTS Ptr getLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Size &ksize, const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT); //! returns the non-separable linear filter engine + // supports: dst type must be the same as source type CV_EXPORTS Ptr createLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Point &anchor = Point(-1, -1), int borderType = BORDER_DEFAULT); //! smooths the image using the normalized box filter - // supports data type: CV_8UC1, CV_8UC4, CV_32FC1 and CV_32FC4 - // supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101,BORDER_WRAP CV_EXPORTS void boxFilter(const oclMat &src, oclMat &dst, int ddepth, Size ksize, Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT); @@ -757,8 +757,6 @@ namespace cv const Point &anchor = Point(-1, -1), int iterations = 1); //! a synonym for normalized box filter - // supports data type: CV_8UC1, CV_8UC4, CV_32FC1 and CV_32FC4 - // supports border type: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT,BORDER_REFLECT_101 static inline void blur(const oclMat &src, oclMat &dst, Size ksize, Point anchor = Point(-1, -1), int borderType = BORDER_CONSTANT) { @@ -766,10 +764,8 @@ namespace cv } //! applies non-separable 2D linear filter to the image - // Note, at the moment this function only works when anchor point is in the kernel center - // and kernel size supported is either 3x3 or 5x5; otherwise the function will fail to output valid result CV_EXPORTS void filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, - Point anchor = Point(-1, -1), int borderType = BORDER_DEFAULT); + Point anchor = Point(-1, -1), double delta = 0.0, int borderType = BORDER_DEFAULT); //! applies separable 2D linear filter to the image CV_EXPORTS void sepFilter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernelX, const Mat &kernelY, diff --git a/modules/ocl/src/filtering.cpp b/modules/ocl/src/filtering.cpp index 112205ba7..4a04e2de8 100644 --- a/modules/ocl/src/filtering.cpp +++ b/modules/ocl/src/filtering.cpp @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -69,37 +69,14 @@ inline void normalizeAnchor(Point &anchor, const Size &ksize) normalizeAnchor(anchor.y, ksize.height); } -inline void normalizeROI(Rect &roi, const Size &ksize, const Point &anchor, const Size &src_size) +inline void normalizeROI(Rect &roi, const Size &ksize, const Point &/*anchor*/, const Size &src_size) { if (roi == Rect(0, 0, -1, -1)) roi = Rect(0, 0, src_size.width, src_size.height); CV_Assert(ksize.height > 0 && ksize.width > 0 && ((ksize.height & 1) == 1) && ((ksize.width & 1) == 1)); - CV_Assert((anchor.x == -1 && anchor.y == -1) || (anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1)); CV_Assert(roi.x >= 0 && roi.y >= 0 && roi.width <= src_size.width && roi.height <= src_size.height); } - - -inline void normalizeKernel(const Mat &kernel, oclMat &gpu_krnl, int type = CV_8U, int *nDivisor = 0, bool reverse = false) -{ - int scale = nDivisor && (kernel.depth() == CV_32F || kernel.depth() == CV_64F) ? 256 : 1; - - if (nDivisor) - *nDivisor = scale; - Mat temp(kernel.size(), type); - kernel.convertTo(temp, type, scale); - Mat cont_krnl = temp.reshape(1, 1); - - if (reverse) - { - int count = cont_krnl.cols >> 1; - - for (int i = 0; i < count; ++i) - std::swap(cont_krnl.at(0, i), cont_krnl.at(0, cont_krnl.cols - 1 - i)); - } - - gpu_krnl.upload(cont_krnl); -} } //////////////////////////////////////////////////////////////////////////////////////////////////// @@ -168,7 +145,7 @@ typedef void (*GPUMorfFilter_t)(const oclMat & , oclMat & , oclMat & , Size &, c class MorphFilter_GPU : public BaseFilter_GPU { public: - MorphFilter_GPU(const Size &ksize_, const Point &anchor_, const oclMat &kernel_, GPUMorfFilter_t func_) : + MorphFilter_GPU(const Size &ksize_, const Point &anchor_, const Mat &kernel_, GPUMorfFilter_t func_) : BaseFilter_GPU(ksize_, anchor_, BORDER_CONSTANT), kernel(kernel_), func(func_), rectKernel(false) {} virtual void operator()(const oclMat &src, oclMat &dst) @@ -344,27 +321,22 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel, openCLExecuteKernel(clCxt, &filtering_morph, kernelName, globalThreads, localThreads, args, -1, -1, compile_option); } -Ptr cv::ocl::getMorphologyFilter_GPU(int op, int type, const Mat &kernel, const Size &ksize, Point anchor) +Ptr cv::ocl::getMorphologyFilter_GPU(int op, int type, const Mat &_kernel, const Size &ksize, Point anchor) { - static const GPUMorfFilter_t GPUMorfFilter_callers[2][5] = - { - {0, GPUErode, 0, GPUErode, GPUErode }, - {0, GPUDilate, 0, GPUDilate, GPUDilate} - }; - CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE); CV_Assert(type == CV_8UC1 || type == CV_8UC3 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC3 || type == CV_32FC4); - oclMat gpu_krnl; - normalizeKernel(kernel, gpu_krnl); normalizeAnchor(anchor, ksize); + Mat kernel8U; + _kernel.convertTo(kernel8U, CV_8U); + Mat kernel = kernel8U.reshape(1, 1); bool noZero = true; for(int i = 0; i < kernel.rows * kernel.cols; ++i) - if(kernel.data[i] != 1) + if(kernel.at(i) != 1) noZero = false; - MorphFilter_GPU* mfgpu = new MorphFilter_GPU(ksize, anchor, gpu_krnl, GPUMorfFilter_callers[op][CV_MAT_CN(type)]); + MorphFilter_GPU* mfgpu = new MorphFilter_GPU(ksize, anchor, kernel, op == MORPH_ERODE ? GPUErode : GPUDilate); if(noZero) mfgpu->rectKernel = true; @@ -524,12 +496,12 @@ void cv::ocl::morphologyEx(const oclMat &src, oclMat &dst, int op, const Mat &ke namespace { -typedef void (*GPUFilter2D_t)(const oclMat & , oclMat & , const oclMat & , const Size &, const Point&, const int); +typedef void (*GPUFilter2D_t)(const oclMat & , oclMat & , const Mat & , const Size &, const Point&, const int); class LinearFilter_GPU : public BaseFilter_GPU { public: - LinearFilter_GPU(const Size &ksize_, const Point &anchor_, const oclMat &kernel_, GPUFilter2D_t func_, + LinearFilter_GPU(const Size &ksize_, const Point &anchor_, const Mat &kernel_, GPUFilter2D_t func_, int borderType_) : BaseFilter_GPU(ksize_, anchor_, borderType_), kernel(kernel_), func(func_) {} @@ -543,118 +515,192 @@ public: }; } -static void GPUFilter2D(const oclMat &src, oclMat &dst, const oclMat &mat_kernel, +// prepare kernel: transpose and make double rows (+align). Returns size of aligned row +// Samples: +// a b c +// Input: d e f +// g h i +// Output, last two zeros is the alignment: +// a d g a d g 0 0 +// b e h b e h 0 0 +// c f i c f i 0 0 +template +static int _prepareKernelFilter2D(std::vector& data, const Mat &kernel) +{ + Mat _kernel; kernel.convertTo(_kernel, DataDepth::value); + int size_y_aligned = roundUp(kernel.rows * 2, 4); + data.clear(); data.resize(size_y_aligned * kernel.cols, 0); + for (int x = 0; x < kernel.cols; x++) + { + for (int y = 0; y < kernel.rows; y++) + { + data[x * size_y_aligned + y] = _kernel.at(y, x); + data[x * size_y_aligned + y + kernel.rows] = _kernel.at(y, x); + } + } + return size_y_aligned; +} + +static void GPUFilter2D(const oclMat &src, oclMat &dst, const Mat &kernel, const Size &ksize, const Point& anchor, const int borderType) { CV_Assert(src.clCxt == dst.clCxt); CV_Assert((src.cols == dst.cols) && (src.rows == dst.rows)); - CV_Assert((src.oclchannels() == dst.oclchannels())); - CV_Assert(ksize.height > 0 && ksize.width > 0 && ((ksize.height & 1) == 1) && ((ksize.width & 1) == 1)); - CV_Assert((anchor.x == -1 && anchor.y == -1) || (anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1)); - CV_Assert(ksize.width == ksize.height); - Context *clCxt = src.clCxt; + CV_Assert(src.oclchannels() == dst.oclchannels()); - int filterWidth = ksize.width; - bool ksize_3x3 = filterWidth == 3 && src.type() != CV_32FC4 && src.type() != CV_32FC3; // CV_32FC4 is not tuned up with filter2d_3x3 kernel + CV_Assert(kernel.cols == ksize.width && kernel.rows == ksize.height); + CV_Assert(kernel.channels() == 1); - string kernelName = ksize_3x3 ? "filter2D_3x3" : "filter2D"; + CV_Assert(anchor.x >= 0 && anchor.x < kernel.cols); + CV_Assert(anchor.y >= 0 && anchor.y < kernel.rows); - size_t src_offset_x = (src.offset % src.step) / src.elemSize(); - size_t src_offset_y = src.offset / src.step; + bool useDouble = src.depth() == CV_64F; - size_t dst_offset_x = (dst.offset % dst.step) / dst.elemSize(); - size_t dst_offset_y = dst.offset / dst.step; - - int paddingPixels = filterWidth & (-2); - - size_t localThreads[3] = {ksize_3x3 ? 256 : 16, ksize_3x3 ? 1 : 16, 1}; - size_t globalThreads[3] = {src.wholecols, src.wholerows, 1}; - - int cn = src.oclchannels(); - int src_step = (int)(src.step/src.elemSize()); - int dst_step = (int)(dst.step/src.elemSize()); - - int localWidth = localThreads[0] + paddingPixels; - int localHeight = localThreads[1] + paddingPixels; - - size_t localMemSize = ksize_3x3 ? 260 * 6 * src.elemSize() : (localWidth * localHeight) * src.elemSize(); - - int vector_lengths[4][7] = {{4, 4, 4, 4, 4, 4, 4}, - {4, 4, 1, 1, 1, 1, 1}, - {1, 1, 1, 1, 1, 1, 1}, - {4, 4, 4, 4, 1, 1, 4} - }; - int cols = dst.cols + ((dst_offset_x) & (vector_lengths[cn - 1][src.depth()] - 1)); - - vector< pair > args; - args.push_back(make_pair(sizeof(cl_mem), (void *)&src.data)); - args.push_back(make_pair(sizeof(cl_mem), (void *)&dst.data)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src_step)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst_step)); - args.push_back(make_pair(sizeof(cl_mem), (void *)&mat_kernel.data)); - args.push_back(make_pair(localMemSize, (void *)NULL)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src_offset_x)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src_offset_y)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst_offset_x)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst_offset_y)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.cols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.rows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&cols)); - char btype[30]; - switch (borderType) + std::vector kernelDataFloat; + std::vector kernelDataDouble; + int kernel_size_y2_aligned = useDouble ? + _prepareKernelFilter2D(kernelDataDouble, kernel) + : _prepareKernelFilter2D(kernelDataFloat, kernel); + oclMat oclKernelParameter; + if (useDouble) { - case 0: - sprintf(btype, "BORDER_CONSTANT"); + oclKernelParameter.createEx(1, kernelDataDouble.size(), CV_64FC1, DEVICE_MEM_R_ONLY, DEVICE_MEM_DEFAULT); + openCLMemcpy2D(src.clCxt, oclKernelParameter.data, kernelDataDouble.size()*sizeof(double), + &kernelDataDouble[0], kernelDataDouble.size()*sizeof(double), + kernelDataDouble.size()*sizeof(double), 1, clMemcpyHostToDevice); + } + else + { + oclKernelParameter.createEx(1, kernelDataFloat.size(), CV_32FC1, DEVICE_MEM_R_ONLY, DEVICE_MEM_DEFAULT); + openCLMemcpy2D(src.clCxt, oclKernelParameter.data, kernelDataFloat.size()*sizeof(float), + &kernelDataFloat[0], kernelDataFloat.size()*sizeof(float), + kernelDataFloat.size()*sizeof(float), 1, clMemcpyHostToDevice); + } + + size_t BLOCK_SIZE = src.clCxt->getDeviceInfo().maxWorkItemSizes[0]; +#if 1 // TODO Mode with several blocks requires a much more VGPRs, so this optimization is not actual for the current devices + size_t BLOCK_SIZE_Y = 1; +#else + size_t BLOCK_SIZE_Y = 8; // TODO Check heuristic value on devices + while (BLOCK_SIZE_Y < BLOCK_SIZE / 8 && BLOCK_SIZE_Y * src.clCxt->getDeviceInfo().maxComputeUnits * 32 < (size_t)src.rows) + BLOCK_SIZE_Y *= 2; +#endif + + CV_Assert((size_t)ksize.width <= BLOCK_SIZE); + + bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0; + + vector > args; + + args.push_back( make_pair( sizeof(cl_mem), (void *)&src.data)); + cl_uint stepBytes = src.step; + args.push_back( make_pair( sizeof(cl_uint), (void *)&stepBytes)); + int offsetXBytes = src.offset % src.step; + int offsetX = offsetXBytes / src.elemSize(); + CV_Assert((int)(offsetX * src.elemSize()) == offsetXBytes); + int offsetY = src.offset / src.step; + int endX = (offsetX + src.cols); + int endY = (offsetY + src.rows); + cl_int rect[4] = {offsetX, offsetY, endX, endY}; + if (!isIsolatedBorder) + { + rect[2] = src.wholecols; + rect[3] = src.wholerows; + } + args.push_back( make_pair( sizeof(cl_int)*4, (void *)&rect[0])); + + args.push_back( make_pair( sizeof(cl_mem), (void *)&dst.data)); + cl_uint _stepBytes = dst.step; + args.push_back( make_pair( sizeof(cl_uint), (void *)&_stepBytes)); + int _offsetXBytes = dst.offset % dst.step; + int _offsetX = _offsetXBytes / dst.elemSize(); + CV_Assert((int)(_offsetX * dst.elemSize()) == _offsetXBytes); + int _offsetY = dst.offset / dst.step; + int _endX = (_offsetX + dst.cols); + int _endY = (_offsetY + dst.rows); + cl_int _rect[4] = {_offsetX, _offsetY, _endX, _endY}; + args.push_back( make_pair( sizeof(cl_int)*4, (void *)&_rect[0])); + + float borderValue[4] = {0, 0, 0, 0}; // DON'T move into 'if' body + double borderValueDouble[4] = {0, 0, 0, 0}; // DON'T move into 'if' body + if ((borderType & ~BORDER_ISOLATED) == BORDER_CONSTANT) + { + if (useDouble) + args.push_back( make_pair( sizeof(double) * src.oclchannels(), (void *)&borderValue[0])); + else + args.push_back( make_pair( sizeof(float) * src.oclchannels(), (void *)&borderValueDouble[0])); + } + + args.push_back( make_pair( sizeof(cl_mem), (void *)&oclKernelParameter.data)); + + const char* btype = NULL; + + switch (borderType & ~BORDER_ISOLATED) + { + case BORDER_CONSTANT: + btype = "BORDER_CONSTANT"; break; - case 1: - sprintf(btype, "BORDER_REPLICATE"); + case BORDER_REPLICATE: + btype = "BORDER_REPLICATE"; break; - case 2: - sprintf(btype, "BORDER_REFLECT"); + case BORDER_REFLECT: + btype = "BORDER_REFLECT"; break; - case 3: + case BORDER_WRAP: CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!"); return; - case 4: - sprintf(btype, "BORDER_REFLECT_101"); + case BORDER_REFLECT101: + btype = "BORDER_REFLECT_101"; break; } - int type = src.depth(); - char build_options[150]; - sprintf(build_options, "-D %s -D IMG_C_%d_%d -D CN=%d -D FILTER_SIZE=%d", btype, cn, type, cn, ksize.width); - openCLExecuteKernel(clCxt, &filtering_laplacian, kernelName, globalThreads, localThreads, args, -1, -1, build_options); + + int requiredTop = anchor.y; + int requiredLeft = BLOCK_SIZE; // not this: anchor.x; + int requiredBottom = ksize.height - 1 - anchor.y; + int requiredRight = BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x; + int h = isIsolatedBorder ? src.rows : src.wholerows; + int w = isIsolatedBorder ? src.cols : src.wholecols; + bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight; + + char build_options[1024]; + sprintf(build_options, "-D LOCAL_SIZE=%d -D BLOCK_SIZE_Y=%d -D DATA_DEPTH=%d -D DATA_CHAN=%d -D USE_DOUBLE=%d " + "-D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d -D KERNEL_SIZE_Y2_ALIGNED=%d " + "-D %s -D %s -D %s", + (int)BLOCK_SIZE, (int)BLOCK_SIZE_Y, + src.depth(), src.oclchannels(), useDouble ? 1 : 0, + anchor.x, anchor.y, ksize.width, ksize.height, kernel_size_y2_aligned, + btype, + extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION", + isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED"); + + size_t gt[3] = {divUp(dst.cols, BLOCK_SIZE - (ksize.width - 1)) * BLOCK_SIZE, divUp(dst.rows, BLOCK_SIZE_Y), 1}, lt[3] = {BLOCK_SIZE, 1, 1}; + openCLExecuteKernel(src.clCxt, &filtering_filter2D, "filter2D", gt, lt, args, -1, -1, build_options); } -Ptr cv::ocl::getLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Size &ksize, +Ptr cv::ocl::getLinearFilter_GPU(int /*srcType*/, int /*dstType*/, const Mat &kernel, const Size &ksize, const Point &anchor, int borderType) { - static const GPUFilter2D_t GPUFilter2D_callers[] = {0, GPUFilter2D, 0, GPUFilter2D, GPUFilter2D}; - - CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC3 || srcType == CV_8UC4 || srcType == CV_32FC1 || srcType == CV_32FC3 || srcType == CV_32FC4) && dstType == srcType); - - oclMat gpu_krnl; Point norm_archor = anchor; - normalizeKernel(kernel, gpu_krnl, CV_32FC1); normalizeAnchor(norm_archor, ksize); - return Ptr(new LinearFilter_GPU(ksize, anchor, gpu_krnl, GPUFilter2D_callers[CV_MAT_CN(srcType)], + return Ptr(new LinearFilter_GPU(ksize, norm_archor, kernel, GPUFilter2D, borderType)); } Ptr cv::ocl::createLinearFilter_GPU(int srcType, int dstType, const Mat &kernel, const Point &anchor, int borderType) { - Size ksize = kernel.size(); + Size ksize = kernel.size(); // TODO remove duplicated parameter Ptr linearFilter = getLinearFilter_GPU(srcType, dstType, kernel, ksize, anchor, borderType); return createFilter2D_GPU(linearFilter); } -void cv::ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor, int borderType) +void cv::ocl::filter2D(const oclMat &src, oclMat &dst, int ddepth, const Mat &kernel, Point anchor, double delta, int borderType) { + CV_Assert(delta == 0); + if (ddepth < 0) ddepth = src.depth(); @@ -713,276 +759,126 @@ Ptr cv::ocl::createSeparableFilter_GPU(const Ptr(new SeparableFilterEngine_GPU(rowFilter, columnFilter)); } -/* -**data type supported: CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4 -**support four border types: BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, BORDER_REFLECT_101 -*/ - -static void GPUFilterBox_8u_C1R(const oclMat &src, oclMat &dst, +static void GPUFilterBox(const oclMat &src, oclMat &dst, Size &ksize, const Point anchor, const int borderType) { //Normalize the result by default - float alpha = ksize.height * ksize.width; + float alpha = 1.0f / (ksize.height * ksize.width); CV_Assert(src.clCxt == dst.clCxt); CV_Assert((src.cols == dst.cols) && (src.rows == dst.rows)); - Context *clCxt = src.clCxt; + CV_Assert(src.oclchannels() == dst.oclchannels()); - string kernelName = "boxFilter_C1_D0"; + size_t BLOCK_SIZE = src.clCxt->getDeviceInfo().maxWorkItemSizes[0]; + size_t BLOCK_SIZE_Y = 8; // TODO Check heuristic value on devices + while (BLOCK_SIZE_Y < BLOCK_SIZE / 8 && BLOCK_SIZE_Y * src.clCxt->getDeviceInfo().maxComputeUnits * 32 < (size_t)src.rows) + BLOCK_SIZE_Y *= 2; - char btype[30]; + CV_Assert((size_t)ksize.width <= BLOCK_SIZE); - switch (borderType) + bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0; + + vector > args; + + args.push_back( make_pair( sizeof(cl_mem), (void *)&src.data)); + cl_uint stepBytes = src.step; + args.push_back( make_pair( sizeof(cl_uint), (void *)&stepBytes)); + int offsetXBytes = src.offset % src.step; + int offsetX = offsetXBytes / src.elemSize(); + CV_Assert((int)(offsetX * src.elemSize()) == offsetXBytes); + int offsetY = src.offset / src.step; + int endX = (offsetX + src.cols); + int endY = (offsetY + src.rows); + cl_int rect[4] = {offsetX, offsetY, endX, endY}; + if (!isIsolatedBorder) { - case 0: - sprintf(btype, "BORDER_CONSTANT"); + rect[2] = src.wholecols; + rect[3] = src.wholerows; + } + args.push_back( make_pair( sizeof(cl_int)*4, (void *)&rect[0])); + + args.push_back( make_pair( sizeof(cl_mem), (void *)&dst.data)); + cl_uint _stepBytes = dst.step; + args.push_back( make_pair( sizeof(cl_uint), (void *)&_stepBytes)); + int _offsetXBytes = dst.offset % dst.step; + int _offsetX = _offsetXBytes / dst.elemSize(); + CV_Assert((int)(_offsetX * dst.elemSize()) == _offsetXBytes); + int _offsetY = dst.offset / dst.step; + int _endX = (_offsetX + dst.cols); + int _endY = (_offsetY + dst.rows); + cl_int _rect[4] = {_offsetX, _offsetY, _endX, _endY}; + args.push_back( make_pair( sizeof(cl_int)*4, (void *)&_rect[0])); + + bool useDouble = src.depth() == CV_64F; + + float borderValue[4] = {0, 0, 0, 0}; // DON'T move into 'if' body + double borderValueDouble[4] = {0, 0, 0, 0}; // DON'T move into 'if' body + if ((borderType & ~BORDER_ISOLATED) == BORDER_CONSTANT) + { + if (useDouble) + args.push_back( make_pair( sizeof(double) * src.oclchannels(), (void *)&borderValue[0])); + else + args.push_back( make_pair( sizeof(float) * src.oclchannels(), (void *)&borderValueDouble[0])); + } + + double alphaDouble = alpha; // DON'T move into 'if' body + if (useDouble) + args.push_back( make_pair( sizeof(double), (void *)&alphaDouble)); + else + args.push_back( make_pair( sizeof(float), (void *)&alpha)); + + const char* btype = NULL; + + switch (borderType & ~BORDER_ISOLATED) + { + case BORDER_CONSTANT: + btype = "BORDER_CONSTANT"; break; - case 1: - sprintf(btype, "BORDER_REPLICATE"); + case BORDER_REPLICATE: + btype = "BORDER_REPLICATE"; break; - case 2: - sprintf(btype, "BORDER_REFLECT"); + case BORDER_REFLECT: + btype = "BORDER_REFLECT"; break; - case 3: + case BORDER_WRAP: CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!"); return; - case 4: - sprintf(btype, "BORDER_REFLECT_101"); + case BORDER_REFLECT101: + btype = "BORDER_REFLECT_101"; break; } - char build_options[150]; - sprintf(build_options, "-D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s", anchor.x, anchor.y, ksize.width, ksize.height, btype); + int requiredTop = anchor.y; + int requiredLeft = BLOCK_SIZE; // not this: anchor.x; + int requiredBottom = ksize.height - 1 - anchor.y; + int requiredRight = BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x; + int h = isIsolatedBorder ? src.rows : src.wholerows; + int w = isIsolatedBorder ? src.cols : src.wholecols; + bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight; - size_t blockSizeX = 256, blockSizeY = 1; - size_t gSize = blockSizeX - (ksize.width - 1); - size_t threads = (dst.offset % dst.step % 4 + dst.cols + 3) / 4; - size_t globalSizeX = threads % gSize == 0 ? threads / gSize * blockSizeX : (threads / gSize + 1) * blockSizeX; - size_t globalSizeY = ((dst.rows + 1) / 2) % blockSizeY == 0 ? ((dst.rows + 1) / 2) : (((dst.rows + 1) / 2) / blockSizeY + 1) * blockSizeY; + CV_Assert(w >= ksize.width && h >= ksize.height); // TODO Other cases are not tested well - size_t globalThreads[3] = { globalSizeX, globalSizeY, 1 }; - size_t localThreads[3] = { blockSizeX, blockSizeY, 1 }; + char build_options[1024]; + sprintf(build_options, "-D LOCAL_SIZE=%d -D BLOCK_SIZE_Y=%d -D DATA_DEPTH=%d -D DATA_CHAN=%d -D USE_DOUBLE=%d -D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d -D %s -D %s -D %s", + (int)BLOCK_SIZE, (int)BLOCK_SIZE_Y, + src.depth(), src.oclchannels(), useDouble ? 1 : 0, + anchor.x, anchor.y, ksize.width, ksize.height, + btype, + extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION", + isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED"); - vector > args; - args.push_back(make_pair(sizeof(cl_mem), &src.data)); - args.push_back(make_pair(sizeof(cl_mem), &dst.data)); - args.push_back(make_pair(sizeof(cl_float), (void *)&alpha)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.offset)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.step)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.offset)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.rows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.cols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.step)); - - openCLExecuteKernel(clCxt, &filtering_boxFilter, kernelName, globalThreads, localThreads, args, -1, -1, build_options); + size_t gt[3] = {divUp(dst.cols, BLOCK_SIZE - (ksize.width - 1)) * BLOCK_SIZE, divUp(dst.rows, BLOCK_SIZE_Y), 1}, lt[3] = {BLOCK_SIZE, 1, 1}; + openCLExecuteKernel(src.clCxt, &filtering_boxFilter, "boxFilter", gt, lt, args, -1, -1, build_options); } -static void GPUFilterBox_8u_C4R(const oclMat &src, oclMat &dst, - Size &ksize, const Point anchor, const int borderType) -{ - //Normalize the result by default - float alpha = ksize.height * ksize.width; - - CV_Assert(src.clCxt == dst.clCxt); - CV_Assert((src.cols == dst.cols) && - (src.rows == dst.rows)); - Context *clCxt = src.clCxt; - - string kernelName = "boxFilter_C4_D0"; - - char btype[30]; - - switch (borderType) - { - case 0: - sprintf(btype, "BORDER_CONSTANT"); - break; - case 1: - sprintf(btype, "BORDER_REPLICATE"); - break; - case 2: - sprintf(btype, "BORDER_REFLECT"); - break; - case 3: - CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!"); - return; - case 4: - sprintf(btype, "BORDER_REFLECT_101"); - break; - } - - char build_options[150]; - sprintf(build_options, "-D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s", anchor.x, anchor.y, ksize.width, ksize.height, btype); - - size_t blockSizeX = 256, blockSizeY = 1; - size_t gSize = blockSizeX - ksize.width / 2 * 2; - size_t globalSizeX = (src.cols) % gSize == 0 ? src.cols / gSize * blockSizeX : (src.cols / gSize + 1) * blockSizeX; - size_t rows_per_thread = 2; - size_t globalSizeY = ((src.rows + rows_per_thread - 1) / rows_per_thread) % blockSizeY == 0 ? ((src.rows + rows_per_thread - 1) / rows_per_thread) : (((src.rows + rows_per_thread - 1) / rows_per_thread) / blockSizeY + 1) * blockSizeY; - - size_t globalThreads[3] = { globalSizeX, globalSizeY, 1}; - size_t localThreads[3] = { blockSizeX, blockSizeY, 1}; - - vector > args; - args.push_back(make_pair(sizeof(cl_mem), &src.data)); - args.push_back(make_pair(sizeof(cl_mem), &dst.data)); - args.push_back(make_pair(sizeof(cl_float), (void *)&alpha)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.offset)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.step)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.offset)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.rows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.cols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.step)); - - openCLExecuteKernel(clCxt, &filtering_boxFilter, kernelName, globalThreads, localThreads, args, -1, -1, build_options); -} - -static void GPUFilterBox_32F_C1R(const oclMat &src, oclMat &dst, - Size &ksize, const Point anchor, const int borderType) -{ - //Normalize the result by default - float alpha = ksize.height * ksize.width; - - CV_Assert(src.clCxt == dst.clCxt); - CV_Assert((src.cols == dst.cols) && - (src.rows == dst.rows)); - Context *clCxt = src.clCxt; - - string kernelName = "boxFilter_C1_D5"; - - char btype[30]; - - switch (borderType) - { - case 0: - sprintf(btype, "BORDER_CONSTANT"); - break; - case 1: - sprintf(btype, "BORDER_REPLICATE"); - break; - case 2: - sprintf(btype, "BORDER_REFLECT"); - break; - case 3: - CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!"); - return; - case 4: - sprintf(btype, "BORDER_REFLECT_101"); - break; - } - - char build_options[150]; - sprintf(build_options, "-D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s", anchor.x, anchor.y, ksize.width, ksize.height, btype); - - size_t blockSizeX = 256, blockSizeY = 1; - size_t gSize = blockSizeX - ksize.width / 2 * 2; - size_t globalSizeX = (src.cols) % gSize == 0 ? src.cols / gSize * blockSizeX : (src.cols / gSize + 1) * blockSizeX; - size_t rows_per_thread = 2; - size_t globalSizeY = ((src.rows + rows_per_thread - 1) / rows_per_thread) % blockSizeY == 0 ? ((src.rows + rows_per_thread - 1) / rows_per_thread) : (((src.rows + rows_per_thread - 1) / rows_per_thread) / blockSizeY + 1) * blockSizeY; - - - size_t globalThreads[3] = { globalSizeX, globalSizeY, 1}; - size_t localThreads[3] = { blockSizeX, blockSizeY, 1}; - - vector > args; - args.push_back(make_pair(sizeof(cl_mem), &src.data)); - args.push_back(make_pair(sizeof(cl_mem), &dst.data)); - args.push_back(make_pair(sizeof(cl_float), (void *)&alpha)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.offset)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.step)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.offset)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.rows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.cols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.step)); - - openCLExecuteKernel(clCxt, &filtering_boxFilter, kernelName, globalThreads, localThreads, args, -1, -1, build_options); -} - -static void GPUFilterBox_32F_C4R(const oclMat &src, oclMat &dst, - Size &ksize, const Point anchor, const int borderType) -{ - //Normalize the result by default - float alpha = ksize.height * ksize.width; - - CV_Assert(src.clCxt == dst.clCxt); - CV_Assert((src.cols == dst.cols) && - (src.rows == dst.rows)); - Context *clCxt = src.clCxt; - - string kernelName = "boxFilter_C4_D5"; - - char btype[30]; - - switch (borderType) - { - case 0: - sprintf(btype, "BORDER_CONSTANT"); - break; - case 1: - sprintf(btype, "BORDER_REPLICATE"); - break; - case 2: - sprintf(btype, "BORDER_REFLECT"); - break; - case 3: - CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!"); - return; - case 4: - sprintf(btype, "BORDER_REFLECT_101"); - break; - } - - char build_options[150]; - sprintf(build_options, "-D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s", anchor.x, anchor.y, ksize.width, ksize.height, btype); - - size_t blockSizeX = 256, blockSizeY = 1; - size_t gSize = blockSizeX - ksize.width / 2 * 2; - size_t globalSizeX = (src.cols) % gSize == 0 ? src.cols / gSize * blockSizeX : (src.cols / gSize + 1) * blockSizeX; - size_t rows_per_thread = 2; - size_t globalSizeY = ((src.rows + rows_per_thread - 1) / rows_per_thread) % blockSizeY == 0 ? ((src.rows + rows_per_thread - 1) / rows_per_thread) : (((src.rows + rows_per_thread - 1) / rows_per_thread) / blockSizeY + 1) * blockSizeY; - - - size_t globalThreads[3] = { globalSizeX, globalSizeY, 1}; - size_t localThreads[3] = { blockSizeX, blockSizeY, 1}; - - vector > args; - args.push_back(make_pair(sizeof(cl_mem), &src.data)); - args.push_back(make_pair(sizeof(cl_mem), &dst.data)); - args.push_back(make_pair(sizeof(cl_float), (void *)&alpha)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.offset)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholerows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.wholecols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&src.step)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.offset)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.rows)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.cols)); - args.push_back(make_pair(sizeof(cl_int), (void *)&dst.step)); - - openCLExecuteKernel(clCxt, &filtering_boxFilter, kernelName, globalThreads, localThreads, args, -1, -1, build_options); -} - - -Ptr cv::ocl::getBoxFilter_GPU(int srcType, int dstType, +Ptr cv::ocl::getBoxFilter_GPU(int /*srcType*/, int /*dstType*/, const Size &ksize, Point anchor, int borderType) { - static const FilterBox_t FilterBox_callers[2][5] = {{0, GPUFilterBox_8u_C1R, 0, GPUFilterBox_8u_C4R, GPUFilterBox_8u_C4R}, - {0, GPUFilterBox_32F_C1R, 0, GPUFilterBox_32F_C4R, GPUFilterBox_32F_C4R} - }; - //Remove this check if more data types need to be supported. - CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC3 || srcType == CV_8UC4 || srcType == CV_32FC1 || - srcType == CV_32FC3 || srcType == CV_32FC4) && dstType == srcType); - normalizeAnchor(anchor, ksize); return Ptr(new GPUBoxFilter(ksize, anchor, - borderType, FilterBox_callers[(CV_MAT_DEPTH(srcType) == CV_32F)][CV_MAT_CN(srcType)])); + borderType, GPUFilterBox)); } Ptr cv::ocl::createBoxFilter_GPU(int srcType, int dstType, @@ -1372,8 +1268,11 @@ void cv::ocl::Scharr(const oclMat &src, oclMat &dst, int ddepth, int dx, int dy, sepFilter2D(src, dst, ddepth, kx, ky, Point(-1, -1), delta, bordertype); } -void cv::ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize, double scale) +void cv::ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize, double scale, + double delta, int borderType) { + CV_Assert(delta == 0); + if (!src.clCxt->supportsFeature(FEATURE_CL_DOUBLE) && src.type() == CV_64F) { CV_Error(CV_OpenCLDoubleNotSupported, "Selected device doesn't support double"); @@ -1382,17 +1281,17 @@ void cv::ocl::Laplacian(const oclMat &src, oclMat &dst, int ddepth, int ksize, d CV_Assert(ksize == 1 || ksize == 3); - int K[2][9] = + double K[2][9] = { {0, 1, 0, 1, -4, 1, 0, 1, 0}, {2, 0, 2, 0, -8, 0, 2, 0, 2} }; - Mat kernel(3, 3, CV_32S, (void *)K[ksize == 3]); + Mat kernel(3, 3, CV_64F, (void *)K[ksize == 3 ? 1 : 0]); if (scale != 1) kernel *= scale; - filter2D(src, dst, ddepth, kernel, Point(-1, -1)); + filter2D(src, dst, ddepth, kernel, Point(-1, -1), 0, borderType); } //////////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/modules/ocl/src/opencl/filtering_boxFilter.cl b/modules/ocl/src/opencl/filtering_boxFilter.cl index 030c13cc5..7f7fd018d 100644 --- a/modules/ocl/src/opencl/filtering_boxFilter.cl +++ b/modules/ocl/src/opencl/filtering_boxFilter.cl @@ -10,13 +10,9 @@ // License Agreement // For Open Source Computer Vision Library // -// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // -// @Authors -// Zhang Ying, zhangying913@gmail.com -// // Redistribution and use in source and binary forms, with or without modification, // are permitted provided that the following conditions are met: // @@ -79,400 +75,298 @@ #define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr)) #endif -#define THREADS 256 -#define ELEM(i, l_edge, r_edge, elem1, elem2) (i) >= (l_edge) && (i) < (r_edge) ? (elem1) : (elem2) - -inline void update_dst_C1_D0(__global uchar *dst, __local uint* temp, - int dst_rows, int dst_cols, - int dst_startX, int dst_x_off, - float alpha) -{ - if(get_local_id(0) < anX || get_local_id(0) >= (THREADS-ksX+anX+1)) - { - return; - } - - uint4 tmp_sum = 0; - int posX = dst_startX - dst_x_off + (get_local_id(0)-anX)*4; - int posY = (get_group_id(1) << 1); - - for(int i=-anX; i<=anX; i++) - { - tmp_sum += vload4(get_local_id(0), temp+i); - } - - if(posY < dst_rows && posX < dst_cols) - { - tmp_sum /= (uint4) alpha; - if(posX >= 0 && posX < dst_cols) - *(dst) = tmp_sum.x; - if(posX+1 >= 0 && posX+1 < dst_cols) - *(dst + 1) = tmp_sum.y; - if(posX+2 >= 0 && posX+2 < dst_cols) - *(dst + 2) = tmp_sum.z; - if(posX+3 >= 0 && posX+3 < dst_cols) - *(dst + 3) = tmp_sum.w; - } -} - - -inline void update_dst_C4_D0(__global uchar4 *dst, __local uint4* temp, - int dst_rows, int dst_cols, - int dst_startX, int dst_x_off, - float alpha) -{ - if(get_local_id(0) >= (THREADS-ksX+1)) - { - return; - } - - int posX = dst_startX - dst_x_off + get_local_id(0); - int posY = (get_group_id(1) << 1); - - uint4 temp_sum = 0; - for(int i=-anX; i<=anX; i++) - { - temp_sum += temp[get_local_id(0) + anX + i]; - } - - if(posX >= 0 && posX < dst_cols && posY >= 0 && posY < dst_rows) - *dst = convert_uchar4(convert_float4(temp_sum)/alpha); -} - -/////////////////////////////////////////////////////////////////////////////////////////////////// -/////////////////////////////////////////8uC1//////////////////////////////////////////////////////// -//////////////////////////////////////////////////////////////////////////////////////////////////// -__kernel void boxFilter_C1_D0(__global const uchar * restrict src, __global uchar *dst, float alpha, - int src_offset, int src_whole_rows, int src_whole_cols, int src_step, - int dst_offset, int dst_rows, int dst_cols, int dst_step - ) -{ - - int col = get_local_id(0); - const int gX = get_group_id(0); - const int gY = get_group_id(1); - int src_x_off = src_offset % src_step; - int src_y_off = src_offset / src_step; - int dst_x_off = dst_offset % dst_step; - int dst_y_off = dst_offset / dst_step; - - int head_off = dst_x_off%4; - int startX = ((gX * (THREADS-ksX+1)-anX) * 4) - head_off + src_x_off; - int startY = (gY << 1) - anY + src_y_off; - int dst_startX = (gX * (THREADS-ksX+1) * 4) - head_off + dst_x_off; - int dst_startY = (gY << 1) + dst_y_off; - - uint4 data[ksY+1]; - __local uint4 temp[2][THREADS]; - +#ifdef EXTRA_EXTRAPOLATION // border > src image size #ifdef BORDER_CONSTANT +// None +#elif defined BORDER_REPLICATE +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + x = max(min(x, maxX - 1), minX); \ + y = max(min(y, maxY - 1), minY); \ + } +#elif defined BORDER_WRAP +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + if (x < minX) \ + x -= ((x - maxX + 1) / maxX) * maxX; \ + if (x >= maxX) \ + x %= maxX; \ + if (y < minY) \ + y -= ((y - maxY + 1) / maxY) * maxY; \ + if (y >= maxY) \ + y %= maxY; \ + } +#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) +#define EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, delta) \ + { \ + if (maxX - minX == 1) \ + x = minX; \ + else \ + do \ + { \ + if (x < minX) \ + x = -(x - minX) - 1 + delta; \ + else \ + x = maxX - 1 - (x - maxX) - delta; \ + } \ + while (x >= maxX || x < minX); \ + \ + if (maxY - minY == 1) \ + y = minY; \ + else \ + do \ + { \ + if (y < minY) \ + y = -(y - minY) - 1 + delta; \ + else \ + y = maxY - 1 - (y - maxY) - delta; \ + } \ + while (y >= maxY || y < minY); \ + } +#ifdef BORDER_REFLECT +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 0) +#elif defined(BORDER_REFLECT_101) +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 1) +#endif +#else +#error No extrapolation method +#endif +#else +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + int _row = y - minY, _col = x - minX; \ + _row = ADDR_H(_row, 0, maxY - minY); \ + _row = ADDR_B(_row, maxY - minY, _row); \ + y = _row + minY; \ + \ + _col = ADDR_L(_col, 0, maxX - minX); \ + _col = ADDR_R(_col, maxX - minX, _col); \ + x = _col + minX; \ + } +#endif - for(int i=0; i < ksY+1; i++) +#if USE_DOUBLE +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#define FPTYPE double +#define CONVERT_TO_FPTYPE CAT(convert_double, VEC_SIZE) +#else +#define FPTYPE float +#define CONVERT_TO_FPTYPE CAT(convert_float, VEC_SIZE) +#endif + +#if DATA_DEPTH == 0 +#define BASE_TYPE uchar +#elif DATA_DEPTH == 1 +#define BASE_TYPE char +#elif DATA_DEPTH == 2 +#define BASE_TYPE ushort +#elif DATA_DEPTH == 3 +#define BASE_TYPE short +#elif DATA_DEPTH == 4 +#define BASE_TYPE int +#elif DATA_DEPTH == 5 +#define BASE_TYPE float +#elif DATA_DEPTH == 6 +#define BASE_TYPE double +#else +#error data_depth +#endif + +#define __CAT(x, y) x##y +#define CAT(x, y) __CAT(x, y) + +#define uchar1 uchar +#define char1 char +#define ushort1 ushort +#define short1 short +#define int1 int +#define float1 float +#define double1 double + +#define convert_uchar1_sat_rte convert_uchar_sat_rte +#define convert_char1_sat_rte convert_char_sat_rte +#define convert_ushort1_sat_rte convert_ushort_sat_rte +#define convert_short1_sat_rte convert_short_sat_rte +#define convert_int1_sat_rte convert_int_sat_rte +#define convert_float1 +#define convert_double1 + +#if DATA_DEPTH == 5 || DATA_DEPTH == 6 +#define CONVERT_TO_TYPE CAT(CAT(convert_, BASE_TYPE), VEC_SIZE) +#else +#define CONVERT_TO_TYPE CAT(CAT(CAT(convert_, BASE_TYPE), VEC_SIZE), _sat_rte) +#endif + +#define VEC_SIZE DATA_CHAN + +#define VEC_TYPE CAT(BASE_TYPE, VEC_SIZE) +#define TYPE VEC_TYPE + +#define SCALAR_TYPE CAT(FPTYPE, VEC_SIZE) + +#define INTERMEDIATE_TYPE CAT(FPTYPE, VEC_SIZE) + +struct RectCoords +{ + int x1, y1, x2, y2; +}; + +//#define DEBUG +#ifdef DEBUG +#define DEBUG_ONLY(x) x +#define ASSERT(condition) do { if (!(condition)) { printf("BUG in boxFilter kernel (global=%d,%d): " #condition "\n", get_global_id(0), get_global_id(1)); } } while (0) +#else +#define DEBUG_ONLY(x) +#define ASSERT(condition) +#endif + + +inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global TYPE *src, const unsigned int srcStepBytes, const struct RectCoords srcCoords +#ifdef BORDER_CONSTANT + , SCALAR_TYPE borderValue +#endif + ) +{ +#ifdef BORDER_ISOLATED + if(pos.x >= srcCoords.x1 && pos.y >= srcCoords.y1 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) +#else + if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) +#endif { - if(startY+i >=0 && startY+i < src_whole_rows && startX+col*4 >=0 && startX+col*4+3 %d, %d\n", pos.x, pos.y, selected_col, selected_row); + + pos = (int2)(selected_col, selected_row); + if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) { - data[i].x = *(src+(startY+i)*src_step + startX + col * 4); - data[i].y = *(src+(startY+i)*src_step + startX + col * 4 + 1); - data[i].z = *(src+(startY+i)*src_step + startX + col * 4 + 2); - data[i].w = *(src+(startY+i)*src_step + startX + col * 4 + 3); + __global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes); + return CONVERT_TO_FPTYPE(*ptr); } else { - data[i]=0; - int con = startY+i >=0 && startY+i < src_whole_rows && startX+col*4 >=0 && startX+col*4=0 && startY+i < src_whole_rows && startX+col*4+1 >=0 && startX+col*4+1=0 && startY+i < src_whole_rows && startX+col*4+2 >=0 && startX+col*4+2=0 && startY+i < src_whole_rows && startX+col*4+3 >=0 && startX+col*4+3src_whole_cols-1) - | (startY+i<0) | (startY+i>src_whole_rows-1); - if(not_all_in_range) - { - int selected_row; - int4 selected_col; - selected_row = ADDR_H(startY+i, 0, src_whole_rows); - selected_row = ADDR_B(startY+i, src_whole_rows, selected_row); - - selected_col.x = ADDR_L(startX+col*4, 0, src_whole_cols); - selected_col.x = ADDR_R(startX+col*4, src_whole_cols, selected_col.x); - - selected_col.y = ADDR_L(startX+col*4+1, 0, src_whole_cols); - selected_col.y = ADDR_R(startX+col*4+1, src_whole_cols, selected_col.y); - - selected_col.z = ADDR_L(startX+col*4+2, 0, src_whole_cols); - selected_col.z = ADDR_R(startX+col*4+2, src_whole_cols, selected_col.z); - - selected_col.w = ADDR_L(startX+col*4+3, 0, src_whole_cols); - selected_col.w = ADDR_R(startX+col*4+3, src_whole_cols, selected_col.w); - - data[i].x = *(src + selected_row * src_step + selected_col.x); - data[i].y = *(src + selected_row * src_step + selected_col.y); - data[i].z = *(src + selected_row * src_step + selected_col.z); - data[i].w = *(src + selected_row * src_step + selected_col.w); - } - else - { - data[i] = convert_uint4(vload4(col,(__global uchar*)(src+(startY+i)*src_step + startX))); - } - } #endif - uint4 tmp_sum = 0; - for(int i=1; i < ksY; i++) - { - tmp_sum += (data[i]); } - - int index = dst_startY * dst_step + dst_startX + (col-anX)*4; - - temp[0][col] = tmp_sum + (data[0]); - temp[1][col] = tmp_sum + (data[ksY]); - barrier(CLK_LOCAL_MEM_FENCE); - update_dst_C1_D0(dst+index, (__local uint *)(temp[0]), - dst_rows, dst_cols, dst_startX, dst_x_off, alpha); - update_dst_C1_D0(dst+index+dst_step, (__local uint *)(temp[1]), - dst_rows, dst_cols, dst_startX, dst_x_off, alpha); - } -/////////////////////////////////////////////////////////////////////////////////////////////////// -/////////////////////////////////////////8uC4//////////////////////////////////////////////////////// -//////////////////////////////////////////////////////////////////////////////////////////////////// -__kernel void boxFilter_C4_D0(__global const uchar4 * restrict src, __global uchar4 *dst, float alpha, - int src_offset, int src_whole_rows, int src_whole_cols, int src_step, - int dst_offset, int dst_rows, int dst_cols, int dst_step - ) -{ - int col = get_local_id(0); - const int gX = get_group_id(0); - const int gY = get_group_id(1); - - int src_x_off = (src_offset % src_step) >> 2; - int src_y_off = src_offset / src_step; - int dst_x_off = (dst_offset % dst_step) >> 2; - int dst_y_off = dst_offset / dst_step; - - int startX = gX * (THREADS-ksX+1) - anX + src_x_off; - int startY = (gY << 1) - anY + src_y_off; - int dst_startX = gX * (THREADS-ksX+1) + dst_x_off; - int dst_startY = (gY << 1) + dst_y_off; - - uint4 data[ksY+1]; - __local uint4 temp[2][THREADS]; +// INPUT PARAMETER: BLOCK_SIZE_Y (via defines) +__kernel +__attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1))) +void boxFilter(__global TYPE *src, const unsigned int srcStepBytes, const int4 srcRC, + __global TYPE *dst, const unsigned int dstStepBytes, const int4 dstRC, #ifdef BORDER_CONSTANT - bool con; - for(int i=0; i < ksY+1; i++) - { - con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows; - int cur_col = clamp(startX + col, 0, src_whole_cols); - - data[i].x = con ? src[(startY+i)*(src_step>>2) + cur_col].x : 0; - data[i].y = con ? src[(startY+i)*(src_step>>2) + cur_col].y : 0; - data[i].z = con ? src[(startY+i)*(src_step>>2) + cur_col].z : 0; - data[i].w = con ? src[(startY+i)*(src_step>>2) + cur_col].w : 0; - } -#else - for(int i=0; i < ksY+1; i++) - { - int selected_row; - int selected_col; - selected_row = ADDR_H(startY+i, 0, src_whole_rows); - selected_row = ADDR_B(startY+i, src_whole_rows, selected_row); - - selected_col = ADDR_L(startX+col, 0, src_whole_cols); - selected_col = ADDR_R(startX+col, src_whole_cols, selected_col); - - - data[i] = convert_uint4(src[selected_row * (src_step>>2) + selected_col]); - } - + SCALAR_TYPE borderValue, #endif - uint4 tmp_sum = 0; - for(int i=1; i < ksY; i++) - { - tmp_sum += (data[i]); - } - - int index = dst_startY * (dst_step>>2)+ dst_startX + col; - - temp[0][col] = tmp_sum + (data[0]); - temp[1][col] = tmp_sum + (data[ksY]); - barrier(CLK_LOCAL_MEM_FENCE); - update_dst_C4_D0(dst+index, (__local uint4 *)(temp[0]), - dst_rows, dst_cols, dst_startX, dst_x_off, alpha); - update_dst_C4_D0(dst+index+(dst_step>>2), (__local uint4 *)(temp[1]), - dst_rows, dst_cols, dst_startX, dst_x_off, alpha); - -} - -/////////////////////////////////////////////////////////////////////////////////////////////////// -/////////////////////////////////////////32fC1//////////////////////////////////////////////////////// -//////////////////////////////////////////////////////////////////////////////////////////////////// -__kernel void boxFilter_C1_D5(__global const float *restrict src, __global float *dst, float alpha, - int src_offset, int src_whole_rows, int src_whole_cols, int src_step, - int dst_offset, int dst_rows, int dst_cols, int dst_step - ) + FPTYPE alpha + ) { - int col = get_local_id(0); - const int gX = get_group_id(0); - const int gY = get_group_id(1); + const struct RectCoords srcCoords = {srcRC.s0, srcRC.s1, srcRC.s2, srcRC.s3}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY + const struct RectCoords dstCoords = {dstRC.s0, dstRC.s1, dstRC.s2, dstRC.s3}; - int src_x_off = (src_offset % src_step) >> 2; - int src_y_off = src_offset / src_step; - int dst_x_off = (dst_offset % dst_step) >> 2; - int dst_y_off = dst_offset / dst_step; + const int x = get_local_id(0) + (LOCAL_SIZE - (KERNEL_SIZE_X - 1)) * get_group_id(0) - ANCHOR_X; + const int y = get_global_id(1) * BLOCK_SIZE_Y; - int startX = gX * (THREADS-ksX+1) - anX + src_x_off; - int startY = (gY << 1) - anY + src_y_off; - int dst_startX = gX * (THREADS-ksX+1) + dst_x_off; - int dst_startY = (gY << 1) + dst_y_off; - float data[ksY+1]; - __local float temp[2][THREADS]; + const int local_id = get_local_id(0); + + INTERMEDIATE_TYPE data[KERNEL_SIZE_Y]; + __local INTERMEDIATE_TYPE sumOfCols[LOCAL_SIZE]; + + int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y); + for(int sy = 0; sy < KERNEL_SIZE_Y; sy++, srcPos.y++) + { + data[sy] = readSrcPixel(srcPos, src, srcStepBytes, srcCoords #ifdef BORDER_CONSTANT - bool con; - float ss; - for(int i=0; i < ksY+1; i++) - { - con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows; - - int cur_col = clamp(startX + col, 0, src_whole_cols); - ss = (startY+i)=0&&cur_col>=0&&cur_col>2) + cur_col]:(float)0; - - data[i] = con ? ss : 0.f; - } -#else - for(int i=0; i < ksY+1; i++) - { - int selected_row; - int selected_col; - selected_row = ADDR_H(startY+i, 0, src_whole_rows); - selected_row = ADDR_B(startY+i, src_whole_rows, selected_row); - - selected_col = ADDR_L(startX+col, 0, src_whole_cols); - selected_col = ADDR_R(startX+col, src_whole_cols, selected_col); - - data[i] = src[selected_row * (src_step>>2) + selected_col]; - } - + , borderValue #endif - float sum0 = 0.0, sum1 = 0.0, sum2 = 0.0; - for(int i=1; i < ksY; i++) - { - sum0 += (data[i]); + ); } - sum1 = sum0 + (data[0]); - sum2 = sum0 + (data[ksY]); - temp[0][col] = sum1; - temp[1][col] = sum2; - barrier(CLK_LOCAL_MEM_FENCE); - if(col < (THREADS-(ksX-1))) - { - col += anX; - int posX = dst_startX - dst_x_off + col - anX; - int posY = (gY << 1); - float tmp_sum[2]= {0.0, 0.0}; - for(int k=0; k<2; k++) - for(int i=-anX; i<=anX; i++) + INTERMEDIATE_TYPE tmp_sum = 0; + for(int sy = 0; sy < KERNEL_SIZE_Y; sy++) + { + tmp_sum += (data[sy]); + } + + sumOfCols[local_id] = tmp_sum; + barrier(CLK_LOCAL_MEM_FENCE); + + int2 pos = (int2)(dstCoords.x1 + x, dstCoords.y1 + y); + __global TYPE* dstPtr = (__global TYPE*)((__global char*)dst + pos.x * sizeof(TYPE) + pos.y * dstStepBytes); // Pointer can be out of bounds! + + int sy_index = 0; // current index in data[] array + int stepsY = min(dstCoords.y2 - pos.y, BLOCK_SIZE_Y); + ASSERT(stepsY > 0); + for (; ;) + { + ASSERT(pos.y < dstCoords.y2); + + if(local_id >= ANCHOR_X && local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X) && + pos.x >= dstCoords.x1 && pos.x < dstCoords.x2) + { + ASSERT(pos.y >= dstCoords.y1 && pos.y < dstCoords.y2); + + INTERMEDIATE_TYPE total_sum = 0; +#pragma unroll + for (int sx = 0; sx < KERNEL_SIZE_X; sx++) { - tmp_sum[k] += temp[k][col+i]; + total_sum += sumOfCols[local_id + sx - ANCHOR_X]; } - for(int i=0; i<2; i++) - { - if(posX >= 0 && posX < dst_cols && (posY+i) >= 0 && (posY+i) < dst_rows) - dst[(dst_startY+i) * (dst_step>>2)+ dst_startX + col - anX] = tmp_sum[i]/alpha; + *dstPtr = CONVERT_TO_TYPE(((INTERMEDIATE_TYPE)alpha) * total_sum); } - } -} - -/////////////////////////////////////////////////////////////////////////////////////////////////// -/////////////////////////////////////////32fC4//////////////////////////////////////////////////////// -//////////////////////////////////////////////////////////////////////////////////////////////////// -__kernel void boxFilter_C4_D5(__global const float4 *restrict src, __global float4 *dst, float alpha, - int src_offset, int src_whole_rows, int src_whole_cols, int src_step, - int dst_offset, int dst_rows, int dst_cols, int dst_step - ) -{ - int col = get_local_id(0); - const int gX = get_group_id(0); - const int gY = get_group_id(1); - - int src_x_off = (src_offset % src_step) >> 4; - int src_y_off = src_offset / src_step; - int dst_x_off = (dst_offset % dst_step) >> 4; - int dst_y_off = dst_offset / dst_step; - - int startX = gX * (THREADS-ksX+1) - anX + src_x_off; - int startY = (gY << 1) - anY + src_y_off; - int dst_startX = gX * (THREADS-ksX+1) + dst_x_off; - int dst_startY = (gY << 1) + dst_y_off; - float4 data[ksY+1]; - __local float4 temp[2][THREADS]; -#ifdef BORDER_CONSTANT - bool con; - float4 ss; - for(int i=0; i < ksY+1; i++) - { - con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows; - - int cur_col = clamp(startX + col, 0, src_whole_cols); - ss = (startY+i)=0&&cur_col>=0&&cur_col>4) + cur_col]:(float4)0; - - data[i] = con ? ss : (float4)(0.0,0.0,0.0,0.0); - } +#if BLOCK_SIZE_Y == 1 + break; #else - for(int i=0; i < ksY+1; i++) - { - int selected_row; - int selected_col; - selected_row = ADDR_H(startY+i, 0, src_whole_rows); - selected_row = ADDR_B(startY+i, src_whole_rows, selected_row); + if (--stepsY == 0) + break; - selected_col = ADDR_L(startX+col, 0, src_whole_cols); - selected_col = ADDR_R(startX+col, src_whole_cols, selected_col); + barrier(CLK_LOCAL_MEM_FENCE); - data[i] = src[selected_row * (src_step>>4) + selected_col]; - } + tmp_sum = sumOfCols[local_id]; // TODO FIX IT: workaround for BUG in OpenCL compiler + // only works with scalars: ASSERT(fabs(tmp_sum - sumOfCols[local_id]) < (INTERMEDIATE_TYPE)1e-6); + tmp_sum -= data[sy_index]; + data[sy_index] = readSrcPixel(srcPos, src, srcStepBytes, srcCoords +#ifdef BORDER_CONSTANT + , borderValue #endif - float4 sum0 = 0.0, sum1 = 0.0, sum2 = 0.0; - for(int i=1; i < ksY; i++) - { - sum0 += (data[i]); - } - sum1 = sum0 + (data[0]); - sum2 = sum0 + (data[ksY]); - temp[0][col] = sum1; - temp[1][col] = sum2; - barrier(CLK_LOCAL_MEM_FENCE); - if(col < (THREADS-(ksX-1))) - { - col += anX; - int posX = dst_startX - dst_x_off + col - anX; - int posY = (gY << 1); + ); + srcPos.y++; - float4 tmp_sum[2]= {(float4)(0.0,0.0,0.0,0.0), (float4)(0.0,0.0,0.0,0.0)}; - for(int k=0; k<2; k++) - for(int i=-anX; i<=anX; i++) - { - tmp_sum[k] += temp[k][col+i]; - } - for(int i=0; i<2; i++) - { - if(posX >= 0 && posX < dst_cols && (posY+i) >= 0 && (posY+i) < dst_rows) - dst[(dst_startY+i) * (dst_step>>4)+ dst_startX + col - anX] = tmp_sum[i]/alpha; - } + tmp_sum += data[sy_index]; + sumOfCols[local_id] = tmp_sum; + sy_index = (sy_index + 1 < KERNEL_SIZE_Y) ? sy_index + 1 : 0; + + barrier(CLK_LOCAL_MEM_FENCE); + + // next line + DEBUG_ONLY(pos.y++); + dstPtr = (__global TYPE*)((__global char*)dstPtr + dstStepBytes); // Pointer can be out of bounds! +#endif // BLOCK_SIZE_Y == 1 } } diff --git a/modules/ocl/src/opencl/filtering_filter2D.cl b/modules/ocl/src/opencl/filtering_filter2D.cl new file mode 100644 index 000000000..f96676689 --- /dev/null +++ b/modules/ocl/src/opencl/filtering_filter2D.cl @@ -0,0 +1,370 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// 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. +// +//M*/ + +#ifdef BORDER_REPLICATE +//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh +#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i)) +#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr)) +#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) :(i)) +#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr)) +#endif + +#ifdef BORDER_REFLECT +//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb +#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i)) +#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)) +#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i)-1 : (i)) +#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr)) +#endif + +#ifdef BORDER_REFLECT_101 +//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba +#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i)) +#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)) +#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i)) +#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr)) +#endif + +//blur function does not support BORDER_WRAP +#ifdef BORDER_WRAP +//BORDER_WRAP: cdefgh|abcdefgh|abcdefg +#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i)) +#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr)) +#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (i)+(b_edge) : (i)) +#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr)) +#endif + +#ifdef EXTRA_EXTRAPOLATION // border > src image size +#ifdef BORDER_CONSTANT +// None +#elif defined BORDER_REPLICATE +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + x = max(min(x, maxX - 1), minX); \ + y = max(min(y, maxY - 1), minY); \ + } +#elif defined BORDER_WRAP +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + if (x < minX) \ + x -= ((x - maxX + 1) / maxX) * maxX; \ + if (x >= maxX) \ + x %= maxX; \ + if (y < minY) \ + y -= ((y - maxY + 1) / maxY) * maxY; \ + if (y >= maxY) \ + y %= maxY; \ + } +#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) +#define EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, delta) \ + { \ + if (maxX - minX == 1) \ + x = minX; \ + else \ + do \ + { \ + if (x < minX) \ + x = -(x - minX) - 1 + delta; \ + else \ + x = maxX - 1 - (x - maxX) - delta; \ + } \ + while (x >= maxX || x < minX); \ + \ + if (maxY - minY == 1) \ + y = minY; \ + else \ + do \ + { \ + if (y < minY) \ + y = -(y - minY) - 1 + delta; \ + else \ + y = maxY - 1 - (y - maxY) - delta; \ + } \ + while (y >= maxY || y < minY); \ + } +#ifdef BORDER_REFLECT +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 0) +#elif defined(BORDER_REFLECT_101) +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 1) +#endif +#else +#error No extrapolation method +#endif +#else +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + int _row = y - minY, _col = x - minX; \ + _row = ADDR_H(_row, 0, maxY - minY); \ + _row = ADDR_B(_row, maxY - minY, _row); \ + y = _row + minY; \ + \ + _col = ADDR_L(_col, 0, maxX - minX); \ + _col = ADDR_R(_col, maxX - minX, _col); \ + x = _col + minX; \ + } +#endif + +#if USE_DOUBLE +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#define FPTYPE double +#define CONVERT_TO_FPTYPE CAT(convert_double, VEC_SIZE) +#else +#define FPTYPE float +#define CONVERT_TO_FPTYPE CAT(convert_float, VEC_SIZE) +#endif + +#if DATA_DEPTH == 0 +#define BASE_TYPE uchar +#elif DATA_DEPTH == 1 +#define BASE_TYPE char +#elif DATA_DEPTH == 2 +#define BASE_TYPE ushort +#elif DATA_DEPTH == 3 +#define BASE_TYPE short +#elif DATA_DEPTH == 4 +#define BASE_TYPE int +#elif DATA_DEPTH == 5 +#define BASE_TYPE float +#elif DATA_DEPTH == 6 +#define BASE_TYPE double +#else +#error data_depth +#endif + +#define __CAT(x, y) x##y +#define CAT(x, y) __CAT(x, y) + +#define uchar1 uchar +#define char1 char +#define ushort1 ushort +#define short1 short +#define int1 int +#define float1 float +#define double1 double + +#define convert_uchar1_sat_rte convert_uchar_sat_rte +#define convert_char1_sat_rte convert_char_sat_rte +#define convert_ushort1_sat_rte convert_ushort_sat_rte +#define convert_short1_sat_rte convert_short_sat_rte +#define convert_int1_sat_rte convert_int_sat_rte +#define convert_float1 +#define convert_double1 + +#if DATA_DEPTH == 5 || DATA_DEPTH == 6 +#define CONVERT_TO_TYPE CAT(CAT(convert_, BASE_TYPE), VEC_SIZE) +#else +#define CONVERT_TO_TYPE CAT(CAT(CAT(convert_, BASE_TYPE), VEC_SIZE), _sat_rte) +#endif + +#define VEC_SIZE DATA_CHAN + +#define VEC_TYPE CAT(BASE_TYPE, VEC_SIZE) +#define TYPE VEC_TYPE + +#define SCALAR_TYPE CAT(FPTYPE, VEC_SIZE) + +#define INTERMEDIATE_TYPE CAT(FPTYPE, VEC_SIZE) + +struct RectCoords +{ + int x1, y1, x2, y2; +}; + +//#define DEBUG +#ifdef DEBUG +#define DEBUG_ONLY(x) x +#define ASSERT(condition) do { if (!(condition)) { printf("BUG in boxFilter kernel (global=%d,%d): " #condition "\n", get_global_id(0), get_global_id(1)); } } while (0) +#else +#define DEBUG_ONLY(x) (void)0 +#define ASSERT(condition) (void)0 +#endif + + +inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global TYPE *src, const unsigned int srcStepBytes, const struct RectCoords srcCoords +#ifdef BORDER_CONSTANT + , SCALAR_TYPE borderValue +#endif + ) +{ +#ifdef BORDER_ISOLATED + if(pos.x >= srcCoords.x1 && pos.y >= srcCoords.y1 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) +#else + if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) +#endif + { + __global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes); + return CONVERT_TO_FPTYPE(*ptr); + } + else + { +#ifdef BORDER_CONSTANT + return borderValue; +#else + int selected_col = pos.x; + int selected_row = pos.y; + + EXTRAPOLATE(selected_col, selected_row, +#ifdef BORDER_ISOLATED + srcCoords.x1, srcCoords.y1, +#else + 0, 0, +#endif + srcCoords.x2, srcCoords.y2 + ); + + // debug border mapping + //printf("pos=%d,%d --> %d, %d\n", pos.x, pos.y, selected_col, selected_row); + + pos = (int2)(selected_col, selected_row); + 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); + return CONVERT_TO_FPTYPE(*ptr); + } + else + { + // for debug only + DEBUG_ONLY(printf("BUG in boxFilter kernel\n")); + return (FPTYPE)(0.0f); + } +#endif + } +} + +// INPUT PARAMETER: BLOCK_SIZE_Y (via defines) + +__kernel +__attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1))) +void filter2D(__global TYPE *src, const unsigned int srcStepBytes, const int4 srcRC, + __global TYPE *dst, const unsigned int dstStepBytes, const int4 dstRC, +#ifdef BORDER_CONSTANT + SCALAR_TYPE borderValue, +#endif + __constant FPTYPE* kernelData // transposed: [KERNEL_SIZE_X][KERNEL_SIZE_Y2_ALIGNED] + ) +{ + const struct RectCoords srcCoords = {srcRC.s0, srcRC.s1, srcRC.s2, srcRC.s3}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY + struct RectCoords dstCoords = {dstRC.s0, dstRC.s1, dstRC.s2, dstRC.s3}; + + 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; + const int y = get_global_id(1) * BLOCK_SIZE_Y; + + INTERMEDIATE_TYPE data[KERNEL_SIZE_Y]; + __local INTERMEDIATE_TYPE sumOfCols[LOCAL_SIZE]; + + int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y); + + int2 pos = (int2)(dstCoords.x1 + x, dstCoords.y1 + y); + __global TYPE* dstPtr = (__global TYPE*)((__global char*)dst + pos.x * sizeof(TYPE) + pos.y * dstStepBytes); // 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); + +#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); + for (; + pos.y < dstCoords.y2; + pos.y++, + dstPtr = (__global TYPE*)((__global char*)dstPtr + dstStepBytes)) +#endif + { + ASSERT(pos.y < dstCoords.y2); + + for ( +#if BLOCK_SIZE_Y > 1 + int sy = readAllpixels ? 0 : -1; sy < (readAllpixels ? KERNEL_SIZE_Y : 0); +#else + int sy = 0, sy_index = 0; sy < KERNEL_SIZE_Y; +#endif + sy++, srcPos.y++) + { + data[sy + sy_index] = readSrcPixel(srcPos, src, srcStepBytes, srcCoords +#ifdef BORDER_CONSTANT + , borderValue +#endif + ); + } + + INTERMEDIATE_TYPE total_sum = 0; + for (int sx = 0; sx < KERNEL_SIZE_X; sx++) + { + { + __constant FPTYPE* k = &kernelData[KERNEL_SIZE_Y2_ALIGNED * sx +#if BLOCK_SIZE_Y > 1 + + KERNEL_SIZE_Y - sy_index +#endif + ]; + INTERMEDIATE_TYPE tmp_sum = 0; + for (int sy = 0; sy < KERNEL_SIZE_Y; sy++) + { + tmp_sum += data[sy] * k[sy]; + } + + sumOfCols[local_id] = tmp_sum; + barrier(CLK_LOCAL_MEM_FENCE); + } + + int id = local_id + sx - ANCHOR_X; + if (id >= 0 && id < LOCAL_SIZE) + total_sum += sumOfCols[id]; + + barrier(CLK_LOCAL_MEM_FENCE); + } + + if (writeResult) + { + ASSERT(pos.y >= dstCoords.y1 && pos.y < dstCoords.y2); + *dstPtr = CONVERT_TO_TYPE(total_sum); + } + +#if BLOCK_SIZE_Y > 1 + readAllpixels = false; +#if BLOCK_SIZE_Y > KERNEL_SIZE_Y + sy_index = (sy_index + 1 <= KERNEL_SIZE_Y) ? sy_index + 1 : 1; +#else + sy_index++; +#endif +#endif // BLOCK_SIZE_Y == 1 + } +} diff --git a/modules/ocl/src/opencl/filtering_laplacian.cl b/modules/ocl/src/opencl/filtering_laplacian.cl deleted file mode 100644 index ea22967df..000000000 --- a/modules/ocl/src/opencl/filtering_laplacian.cl +++ /dev/null @@ -1,381 +0,0 @@ -/*M/////////////////////////////////////////////////////////////////////////////////////// -// -// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. -// -// By downloading, copying, installing or using the software you agree to this license. -// If you do not agree to this license, do not download, install, -// copy or use the software. -// -// -// License Agreement -// For Open Source Computer Vision Library -// -// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. -// Third party copyrights are property of their respective owners. -// -// @Authors -// Pang Erping, erping@multicorewareinc.com -// Jia Haipeng, jiahaipeng95@gmail.com -// Peng Xiao, pengxiao@outlook.com -// -// Redistribution and use in source and binary forms, with or without modification, -// are permitted provided that the following conditions are met: -// -// * Redistribution's of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// -// * Redistribution's in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// * The name of the copyright holders may not be used to endorse or promote products -// derived from this software without specific prior written permission. -// -// This software is provided by the copyright holders and contributors as is and -// any express or implied warranties, including, but not limited to, the implied -// warranties of merchantability and fitness for a particular purpose are disclaimed. -// In no event shall the Intel Corporation or contributors be liable for any direct, -// indirect, incidental, special, exemplary, or consequential damages -// (including, but not limited to, procurement of substitute goods or services; -// loss of use, data, or profits; or business interruption) however caused -// and on any theory of liability, whether in contract, strict liability, -// 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. -// -//M*/ - -/////////////////////////////////////////////////////////////////////////////////////////////////// -/////////////////////////////////Macro for border type//////////////////////////////////////////// -///////////////////////////////////////////////////////////////////////////////////////////////// -#ifdef BORDER_REPLICATE - -//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh -#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i)) -#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr)) -#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) : (i)) -#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr)) -#endif - -#ifdef BORDER_REFLECT -#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? ((l_edge)<<1)-(i)-1 : (i)) -#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)) -#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? ((t_edge)<<1)-(i)-1 : (i)) -#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr)) -#endif - -#ifdef BORDER_REFLECT_101 -//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba -#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? ((l_edge)<<1)-(i) : (i)) -#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)) -#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? ((t_edge)<<1)-(i) : (i)) -#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr)) -#endif - -#ifdef IMG_C_1_0 -#define T_IMG uchar -#define T_IMGx4 uchar4 -#define T_IMG_C1 uchar -#define CONVERT_TYPE convert_uchar_sat -#define CONVERT_TYPEx4 convert_uchar4_sat -#endif -#ifdef IMG_C_4_0 -#define T_IMG uchar4 -#define T_IMGx4 uchar16 -#define T_IMG_C1 uchar -#define CONVERT_TYPE convert_uchar4_sat -#define CONVERT_TYPEx4 convert_uchar16_sat -#endif -#ifdef IMG_C_1_5 -#define T_IMG float -#define T_IMGx4 float4 -#define T_IMG_C1 float -#define CONVERT_TYPE convert_float -#define CONVERT_TYPEx4 convert_float4 -#endif -#ifdef IMG_C_4_5 -#define T_IMG float4 -#define T_IMGx4 float16 -#define T_IMG_C1 float -#define CONVERT_TYPE convert_float4 -#define CONVERT_TYPEx4 convert_float16 -#endif - -#ifndef CN -#define CN 1 -#endif - -#if CN == 1 -#define T_SUM float -#define T_SUMx4 float4 -#define CONVERT_TYPE_SUM convert_float -#define CONVERT_TYPE_SUMx4 convert_float4 -#define SUM_ZERO (0.0f) -#define SUM_ZEROx4 (0.0f, 0.0f, 0.0f, 0.0f) -#define VLOAD4 vload4 -#define SX x -#define SY y -#define SZ z -#define SW w -#elif CN == 4 -#define T_SUM float4 -#define T_SUMx4 float16 -#define CONVERT_TYPE_SUM convert_float4 -#define CONVERT_TYPE_SUMx4 convert_float16 -#define SUM_ZERO (0.0f, 0.0f, 0.0f, 0.0f) -#define SUM_ZEROx4 (0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f) -#define VLOAD4 vload16 -#define SX s0123 -#define SY s4567 -#define SZ s89ab -#define SW scdef -#endif - -#ifndef FILTER_SIZE -#define FILTER_SIZE 3 -#endif - -#define LOCAL_GROUP_SIZE 16 - -#define LOCAL_WIDTH ((FILTER_SIZE/2)*2 + LOCAL_GROUP_SIZE) -#define LOCAL_HEIGHT ((FILTER_SIZE/2)*2 + LOCAL_GROUP_SIZE) - -#define FILTER_RADIUS (FILTER_SIZE >> 1) - -__kernel void filter2D( - __global T_IMG *src, - __global T_IMG *dst, - int src_step, - int dst_step, - __constant float *mat_kernel, - __local T_IMG *local_data, - int wholerows, - int wholecols, - int src_offset_x, - int src_offset_y, - int dst_offset_x, - int dst_offset_y, - int cols, - int rows, - int operate_cols -) -{ - int groupStartCol = get_group_id(0) * get_local_size(0); - int groupStartRow = get_group_id(1) * get_local_size(1); - - int localCol = get_local_id(0); - int localRow = get_local_id(1); - int globalCol = groupStartCol + localCol; - int globalRow = groupStartRow + localRow; - const int src_offset = mad24(src_offset_y, src_step, src_offset_x); - const int dst_offset = mad24(dst_offset_y, dst_step, dst_offset_x); - -#ifdef BORDER_CONSTANT - for(int i = localRow; i < LOCAL_HEIGHT; i += get_local_size(1)) - { - int curRow = groupStartRow + i; - for(int j = localCol; j < LOCAL_WIDTH; j += get_local_size(0)) - { - int curCol = groupStartCol + j; - if(curRow < FILTER_RADIUS - src_offset_y || (curRow - FILTER_RADIUS) >= wholerows - src_offset_y|| - curCol < FILTER_RADIUS - src_offset_x || (curCol - FILTER_RADIUS) >= wholecols - src_offset_x) - { - local_data[(i) * LOCAL_WIDTH + j] = 0; - } - else - { - local_data[(i) * LOCAL_WIDTH + j] = src[(curRow - FILTER_RADIUS) * src_step + curCol - FILTER_RADIUS + src_offset]; - } - } - } -#else - for(int i = localRow; i < LOCAL_HEIGHT; i += get_local_size(1)) - { - int curRow = groupStartRow + i; - - curRow = ADDR_H(curRow, FILTER_RADIUS - src_offset_y, wholerows - src_offset_y); - - curRow = ADDR_B(curRow - FILTER_RADIUS, wholerows - src_offset_y, curRow - FILTER_RADIUS); - - for(int j = localCol; j < LOCAL_WIDTH; j += get_local_size(0)) - { - int curCol = groupStartCol + j; - curCol = ADDR_L(curCol, FILTER_RADIUS - src_offset_x, wholecols - src_offset_x); - curCol = ADDR_R(curCol - FILTER_RADIUS, wholecols - src_offset_x, curCol - FILTER_RADIUS); - if(curRow < wholerows && curCol < wholecols) - { - local_data[(i) * LOCAL_WIDTH + j] = src[(curRow) * src_step + curCol + src_offset]; - } - } - } -#endif - - barrier(CLK_LOCAL_MEM_FENCE); - if(globalRow < rows && globalCol < cols) - { - T_SUM sum = (T_SUM)(SUM_ZERO); - int filterIdx = 0; - for(int i = 0; i < FILTER_SIZE; i++) - { - int offset = (i + localRow) * LOCAL_WIDTH; - - for(int j = 0; j < FILTER_SIZE; j++) - { - sum += CONVERT_TYPE_SUM(local_data[offset + j + localCol]) * mat_kernel[filterIdx++]; - } - } - dst[(globalRow)*dst_step + (globalCol) + dst_offset] = CONVERT_TYPE(sum); - } -} - -/// following is specific for 3x3 kernels - -////////////////////////////////////////////////////////////////////////////////////////////////////// -/////////////////////////////Macro for define elements number per thread///////////////////////////// -//////////////////////////////////////////////////////////////////////////////////////////////////// - -#define ANX 1 -#define ANY 1 - -#define ROWS_PER_GROUP 4 -#define ROWS_PER_GROUP_BITS 2 -#define ROWS_FETCH (ROWS_PER_GROUP + ANY + ANY) //(ROWS_PER_GROUP + anY * 2) - -#define THREADS_PER_ROW 64 -#define THREADS_PER_ROW_BIT 6 - -#define ELEMENTS_PER_THREAD 4 -#define ELEMENTS_PER_THREAD_BIT 2 - -#define LOCAL_MEM_STEP 260 //divup((get_local_size(0) + anX * 2), 4) * 4 - -/////////////////////////////////////////////////////////////////////////////////////////////////// -/////////////////////////////////////////8uC1//////////////////////////////////////////////////////// -//////////////////////////////////////////////////////////////////////////////////////////////////// - -__kernel void filter2D_3x3( - __global T_IMG *src, - __global T_IMG *dst, - int src_step, - int dst_step, - __constant float *mat_kernel, - __local T_IMG *local_data, - int wholerows, - int wholecols, - int src_offset_x, - int src_offset_y, - int dst_offset_x, - int dst_offset_y, - int cols, - int rows, - int operate_cols -) -{ - int gX = get_global_id(0); - int gY = get_global_id(1); - - int lX = get_local_id(0); - - int groupX_size = get_local_size(0); - int groupX_id = get_group_id(0); - -#define dst_align (dst_offset_x & 3) - int cols_start_index_group = src_offset_x - dst_align + groupX_size * groupX_id - ANX; - int rows_start_index = src_offset_y + (gY << ROWS_PER_GROUP_BITS) - ANY; - - if((gY << 2) < rows) - { - for(int i = 0; i < ROWS_FETCH; ++i) - { - if((rows_start_index - src_offset_y) + i < rows + ANY) - { -#ifdef BORDER_CONSTANT - int selected_row = rows_start_index + i; - int selected_cols = cols_start_index_group + lX; - - T_IMG data = src[mad24(selected_row, src_step, selected_cols)]; - int con = selected_row >= 0 && selected_row < wholerows && selected_cols >= 0 && selected_cols < wholecols; - data = con ? data : (T_IMG)(0); - local_data[mad24(i, LOCAL_MEM_STEP, lX)] = data; - - if(lX < (ANX << 1)) - { - selected_cols = cols_start_index_group + lX + groupX_size; - - data = src[mad24(selected_row, src_step, selected_cols)]; - con = selected_row >= 0 && selected_row < wholerows && selected_cols >= 0 && selected_cols < wholecols; - data = con ? data : (T_IMG)(0); - local_data[mad24(i, LOCAL_MEM_STEP, lX) + groupX_size] = data; - } -#else - int selected_row = ADDR_H(rows_start_index + i, 0, wholerows); - selected_row = ADDR_B(rows_start_index + i, wholerows, selected_row); - - int selected_cols = ADDR_L(cols_start_index_group + lX, 0, wholecols); - selected_cols = ADDR_R(cols_start_index_group + lX, wholecols, selected_cols); - - T_IMG data = src[mad24(selected_row, src_step, selected_cols)]; - - local_data[mad24(i, LOCAL_MEM_STEP, lX)] = data; - - if(lX < (ANX << 1)) - { - selected_cols = cols_start_index_group + lX + groupX_size; - selected_cols = ADDR_R(selected_cols, wholecols, selected_cols); - - data = src[mad24(selected_row, src_step, selected_cols)]; - local_data[mad24(i, LOCAL_MEM_STEP, lX) + groupX_size] = data; - } -#endif - } - } - } - barrier(CLK_LOCAL_MEM_FENCE); - - int process_col = groupX_size * groupX_id + ((lX % THREADS_PER_ROW) << 2); - if(((gY << 2) < rows) && (process_col < operate_cols)) - { - int dst_cols_start = dst_offset_x; - int dst_cols_end = dst_offset_x + cols; - int dst_cols_index = (dst_offset_x + process_col) & 0xfffffffc; - - int dst_rows_end = dst_offset_y + rows; - int dst_rows_index = dst_offset_y + (gY << ROWS_PER_GROUP_BITS) + (lX >> THREADS_PER_ROW_BIT); - dst = dst + mad24(dst_rows_index, dst_step, dst_cols_index); - - T_IMGx4 dst_data = *(__global T_IMGx4 *)dst; - - T_SUMx4 sum = (T_SUMx4)SUM_ZEROx4; - T_IMGx4 data; - - for(int i = 0; i < FILTER_SIZE; i++) - { -#pragma unroll - for(int j = 0; j < FILTER_SIZE; j++) - { - if(dst_rows_index < dst_rows_end) - { - int local_row = (lX >> THREADS_PER_ROW_BIT) + i; - int local_cols = ((lX % THREADS_PER_ROW) << ELEMENTS_PER_THREAD_BIT) + j; - - data = VLOAD4(0, (__local T_IMG_C1 *)(local_data + local_row * LOCAL_MEM_STEP + local_cols)); - sum = sum + (mat_kernel[i * FILTER_SIZE + j] * CONVERT_TYPE_SUMx4(data)); - } - } - } - - if(dst_rows_index < dst_rows_end) - { - T_IMGx4 tmp_dst = CONVERT_TYPEx4(sum); - tmp_dst.SX = ((dst_cols_index + 0 >= dst_cols_start) && (dst_cols_index + 0 < dst_cols_end)) ? - tmp_dst.SX : dst_data.SX; - tmp_dst.SY = ((dst_cols_index + 1 >= dst_cols_start) && (dst_cols_index + 1 < dst_cols_end)) ? - tmp_dst.SY : dst_data.SY; - tmp_dst.SZ = ((dst_cols_index + 2 >= dst_cols_start) && (dst_cols_index + 2 < dst_cols_end)) ? - tmp_dst.SZ : dst_data.SZ; - tmp_dst.SW = ((dst_cols_index + 3 >= dst_cols_start) && (dst_cols_index + 3 < dst_cols_end)) ? - tmp_dst.SW : dst_data.SW; - *(__global T_IMGx4 *)dst = tmp_dst; - } - } -} diff --git a/modules/ocl/test/test_filters.cpp b/modules/ocl/test/test_filters.cpp index 86ff834d3..a8583b28a 100644 --- a/modules/ocl/test/test_filters.cpp +++ b/modules/ocl/test/test_filters.cpp @@ -59,10 +59,15 @@ using namespace cv; PARAM_TEST_CASE(FilterTestBase, MatType, int, // kernel size Size, // dx, dy - int, // border type, or iteration + int, // border type + double, // optional parameter bool) // roi or not { + bool isFP; + int type, borderType, ksize; + Size size; + double param; bool useRoi; Mat src, dst_whole, src_roi, dst_roi; @@ -72,31 +77,53 @@ PARAM_TEST_CASE(FilterTestBase, MatType, { type = GET_PARAM(0); ksize = GET_PARAM(1); + size = GET_PARAM(2); borderType = GET_PARAM(3); - useRoi = GET_PARAM(4); + param = GET_PARAM(4); + useRoi = GET_PARAM(5); + + isFP = (CV_MAT_DEPTH(type) == CV_32F || CV_MAT_DEPTH(type) == CV_64F); } - void random_roi() + void random_roi(int minSize = 1) { - Size roiSize = randomSize(1, MAX_VALUE); + if (minSize == 0) + minSize = ksize; + Size roiSize = randomSize(minSize, MAX_VALUE); Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); - randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256); + randomSubMat(src, src_roi, roiSize, srcBorder, type, isFP ? 0 : 5, isFP ? 1 : 256); Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); - randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, 5, 16); + randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, isFP ? 0.20 : 60, isFP ? 0.25 : 70); generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder); generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder); } - void Near(double threshold = 0.0) + void Near() + { + if (isFP) + Near(1e-6, true); + else + Near(1, false); + } + + void Near(double threshold, bool relative) { Mat roi, whole; gdst_whole.download(whole); gdst_roi.download(roi); - EXPECT_MAT_NEAR(dst_whole, whole, threshold); - EXPECT_MAT_NEAR(dst_roi, roi, threshold); + if (relative) + { + EXPECT_MAT_NEAR_RELATIVE(dst_whole, whole, threshold); + EXPECT_MAT_NEAR_RELATIVE(dst_roi, roi, threshold); + } + else + { + EXPECT_MAT_NEAR(dst_whole, whole, threshold); + EXPECT_MAT_NEAR(dst_roi, roi, threshold); + } } }; @@ -111,12 +138,12 @@ OCL_TEST_P(Blur, Mat) for (int j = 0; j < LOOP_TIMES; j++) { - random_roi(); + random_roi(0); // TODO NOTE: min value for size is kernel size (temporary bypass border issues in CPU implementation) blur(src_roi, dst_roi, kernelSize, Point(-1, -1), borderType); ocl::blur(gsrc_roi, gdst_roi, kernelSize, Point(-1, -1), borderType); // TODO anchor - Near(1.0); + Near(); } } @@ -127,64 +154,51 @@ typedef FilterTestBase LaplacianTest; OCL_TEST_P(LaplacianTest, Accuracy) { + double scale = param; + for (int j = 0; j < LOOP_TIMES; j++) { random_roi(); - // border type is used as a scale factor for the Laplacian kernel - double scale = static_cast(borderType); + Laplacian(src_roi, dst_roi, -1, ksize, scale, 0, borderType); + ocl::Laplacian(gsrc_roi, gdst_roi, -1, ksize, scale, 0, borderType); - Laplacian(src_roi, dst_roi, -1, ksize, scale); - ocl::Laplacian(gsrc_roi, gdst_roi, -1, ksize, scale); - - Near(1e-5); + Near(); } } ///////////////////////////////////////////////////////////////////////////////////////////////// // erode & dilate -struct ErodeDilate : - public FilterTestBase -{ - int iterations; - - virtual void SetUp() - { - type = GET_PARAM(0); - ksize = GET_PARAM(1); - iterations = GET_PARAM(3); - useRoi = GET_PARAM(4); - } -}; - -typedef ErodeDilate Erode; +typedef FilterTestBase Erode; OCL_TEST_P(Erode, Mat) { // erode or dilate kernel Size kernelSize(ksize, ksize); Mat kernel; + int iterations = (int)param; for (int j = 0; j < LOOP_TIMES; j++) { - kernel = randomMat(kernelSize, CV_8UC1, 0, 3); - random_roi(); - cv::erode(src_roi, dst_roi, kernel, Point(-1, -1), iterations); - ocl::erode(gsrc_roi, gdst_roi, kernel, Point(-1, -1), iterations); // TODO iterations, borderType + kernel = randomMat(kernelSize, CV_8UC1, 0, 3); - Near(1e-5); + cv::erode(src_roi, dst_roi, kernel, Point(-1, -1), iterations);//, borderType); + ocl::erode(gsrc_roi, gdst_roi, kernel, Point(-1, -1), iterations);//, borderType); + + Near(); } } -typedef ErodeDilate Dilate; +typedef FilterTestBase Dilate; OCL_TEST_P(Dilate, Mat) { // erode or dilate kernel Mat kernel; + int iterations = (int)param; for (int j = 0; j < LOOP_TIMES; j++) { @@ -195,79 +209,56 @@ OCL_TEST_P(Dilate, Mat) cv::dilate(src_roi, dst_roi, kernel, Point(-1, -1), iterations); ocl::dilate(gsrc_roi, gdst_roi, kernel, Point(-1, -1), iterations); // TODO iterations, borderType - Near(1e-5); + Near(); } } ///////////////////////////////////////////////////////////////////////////////////////////////// // Sobel -struct SobelTest : - public FilterTestBase -{ - int dx, dy; - - virtual void SetUp() - { - type = GET_PARAM(0); - ksize = GET_PARAM(1); - borderType = GET_PARAM(3); - useRoi = GET_PARAM(4); - - Size d = GET_PARAM(2); - dx = d.width, dy = d.height; - } -}; +typedef FilterTestBase SobelTest; OCL_TEST_P(SobelTest, Mat) { + int dx = size.width, dy = size.height; + double scale = param; + for (int j = 0; j < LOOP_TIMES; j++) { random_roi(); - Sobel(src_roi, dst_roi, -1, dx, dy, ksize, /* scale */ 0.00001, /* delta */0, borderType); - ocl::Sobel(gsrc_roi, gdst_roi, -1, dx, dy, ksize, /* scale */ 0.00001, /* delta */ 0, borderType); + Sobel(src_roi, dst_roi, -1, dx, dy, ksize, scale, /* delta */0, borderType); + ocl::Sobel(gsrc_roi, gdst_roi, -1, dx, dy, ksize, scale, /* delta */0, borderType); - Near(1); + Near(); } } ///////////////////////////////////////////////////////////////////////////////////////////////// // Scharr -typedef SobelTest ScharrTest; +typedef FilterTestBase ScharrTest; OCL_TEST_P(ScharrTest, Mat) { + int dx = size.width, dy = size.height; + double scale = param; + for (int j = 0; j < LOOP_TIMES; j++) { random_roi(); - Scharr(src_roi, dst_roi, -1, dx, dy, /* scale */ 1, /* delta */ 0, borderType); - ocl::Scharr(gsrc_roi, gdst_roi, -1, dx, dy, /* scale */ 1, /* delta */ 0, borderType); + Scharr(src_roi, dst_roi, -1, dx, dy, scale, /* delta */ 0, borderType); + ocl::Scharr(gsrc_roi, gdst_roi, -1, dx, dy, scale, /* delta */ 0, borderType); - Near(1); + Near(); } } ///////////////////////////////////////////////////////////////////////////////////////////////// // GaussianBlur -struct GaussianBlurTest : - public FilterTestBase -{ - double sigma1, sigma2; - - virtual void SetUp() - { - type = GET_PARAM(0); - ksize = GET_PARAM(1); - borderType = GET_PARAM(3); - - sigma1 = rng.uniform(0.1, 1.0); - sigma2 = rng.uniform(0.1, 1.0); - } -}; +typedef FilterTestBase GaussianBlurTest; OCL_TEST_P(GaussianBlurTest, Mat) { @@ -275,10 +266,13 @@ OCL_TEST_P(GaussianBlurTest, Mat) { random_roi(); + double sigma1 = rng.uniform(0.1, 1.0); + double sigma2 = rng.uniform(0.1, 1.0); + GaussianBlur(src_roi, dst_roi, Size(ksize, ksize), sigma1, sigma2, borderType); ocl::GaussianBlur(gsrc_roi, gdst_roi, Size(ksize, ksize), sigma1, sigma2, borderType); - Near(1); + Near(); } } @@ -289,19 +283,24 @@ typedef FilterTestBase Filter2D; OCL_TEST_P(Filter2D, Mat) { - const Size kernelSize(ksize, ksize); - Mat kernel; - for (int j = 0; j < LOOP_TIMES; j++) { - kernel = randomMat(kernelSize, CV_32FC1, 0.0, 1.0); - random_roi(); - cv::filter2D(src_roi, dst_roi, -1, kernel, Point(-1, -1), 0.0, borderType); // TODO anchor - ocl::filter2D(gsrc_roi, gdst_roi, -1, kernel, Point(-1, -1), borderType); + Point anchor(-1, -1); + if (size.width >= 0) + anchor.x = size.width % ksize; + if (size.height >= 0) + anchor.y = size.height % ksize; - Near(1); + const Size kernelSize(ksize, ksize); + Mat kernel = randomMat(kernelSize, CV_32FC1, 0, 1.0); + kernel *= 1.0 / (double)(ksize * ksize); + + cv::filter2D(src_roi, dst_roi, -1, kernel, anchor, 0.0, borderType); + ocl::filter2D(gsrc_roi, gdst_roi, -1, kernel, anchor, 0.0, borderType); + + Near(); } } @@ -322,7 +321,7 @@ OCL_TEST_P(Bilateral, Mat) cv::bilateralFilter(src_roi, dst_roi, ksize, sigmacolor, sigmaspace, borderType); ocl::bilateralFilter(gsrc_roi, gdst_roi, ksize, sigmacolor, sigmaspace, borderType); - Near(1); + Near(); } } @@ -342,7 +341,7 @@ OCL_TEST_P(AdaptiveBilateral, Mat) adaptiveBilateralFilter(src_roi, dst_roi, kernelSize, 5, Point(-1, -1), borderType); // TODO anchor ocl::adaptiveBilateralFilter(gsrc_roi, gdst_roi, kernelSize, 5, Point(-1, -1), borderType); - Near(1); + Near(); } } @@ -366,80 +365,97 @@ OCL_TEST_P(MedianFilter, Mat) ////////////////////////////////////////////////////////////////////////////////////////////////////////////// +#define FILTER_BORDER_SET_NO_ISOLATED \ + Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_WRAP, (int)BORDER_REFLECT_101/*, \ + (int)BORDER_CONSTANT|BORDER_ISOLATED, (int)BORDER_REPLICATE|BORDER_ISOLATED, \ + (int)BORDER_REFLECT|BORDER_ISOLATED, (int)BORDER_WRAP|BORDER_ISOLATED, \ + (int)BORDER_REFLECT_101|BORDER_ISOLATED*/) // WRAP and ISOLATED are not supported by cv:: version + +#define FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED \ + Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, /*(int)BORDER_WRAP,*/ (int)BORDER_REFLECT_101/*, \ + (int)BORDER_CONSTANT|BORDER_ISOLATED, (int)BORDER_REPLICATE|BORDER_ISOLATED, \ + (int)BORDER_REFLECT|BORDER_ISOLATED, (int)BORDER_WRAP|BORDER_ISOLATED, \ + (int)BORDER_REFLECT_101|BORDER_ISOLATED*/) // WRAP and ISOLATED are not supported by cv:: version + + INSTANTIATE_TEST_CASE_P(Filter, Blur, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4), Values(3, 5, 7), Values(Size(0, 0)), // not used - Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_REFLECT_101), + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(0.0), // not used Bool())); INSTANTIATE_TEST_CASE_P(Filter, LaplacianTest, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), Values(1, 3), Values(Size(0, 0)), // not used - Values(1, 2), // value is used as scale factor for kernel + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(1.0, 0.2, 3.0), // scalar Bool())); INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), Values(3, 5, 7), Values(Size(0, 0)), // not used - testing::Range(1, 4), + Values(0), // not used + Values(1.0, 2.0, 3.0), Bool())); INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), Values(3, 5, 7), Values(Size(0, 0)), // not used - testing::Range(1, 4), + Values(0), // not used + Values(1.0, 2.0, 3.0), Bool())); INSTANTIATE_TEST_CASE_P(Filter, SobelTest, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), Values(3, 5), - Values(Size(1, 0), Size(1, 1), Size(2, 0), Size(2, 1)), - Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101, - (int)BORDER_REPLICATE, (int)BORDER_REFLECT), + Values(Size(1, 0), Size(1, 1), Size(2, 0), Size(2, 1)), // dx, dy + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(0.0), // not used Bool())); INSTANTIATE_TEST_CASE_P(Filter, ScharrTest, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), - Values(0), // not used - Values(Size(0, 1), Size(1, 0)), - Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101, - (int)BORDER_REPLICATE, (int)BORDER_REFLECT), + Values(1), + Values(Size(0, 1), Size(1, 0)), // dx, dy + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(1.0, 0.2), // scalar Bool())); INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4), Values(3, 5), Values(Size(0, 0)), // not used - Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101, - (int)BORDER_REPLICATE, (int)BORDER_REFLECT), + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(0.0), // not used Bool())); INSTANTIATE_TEST_CASE_P(Filter, Filter2D, testing::Combine( Values(CV_8UC1, CV_32FC1, CV_32FC4), - Values(3, 15, 25), - Values(Size(0, 0)), // not used - Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101, - (int)BORDER_REPLICATE, (int)BORDER_REFLECT), + Values(3, 15), // TODO 25: CPU implementation has some issues + Values(Size(-1, -1), Size(0, 0), Size(2, 1)), // anchor + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(0.0), // not used Bool())); INSTANTIATE_TEST_CASE_P(Filter, Bilateral, Combine( Values(CV_8UC1, CV_8UC3), Values(5, 9), Values(Size(0, 0)), // not used - Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, - (int)BORDER_REFLECT, (int)BORDER_WRAP, (int)BORDER_REFLECT_101), + FILTER_BORDER_SET_NO_ISOLATED, + Values(0.0), // not used Bool())); INSTANTIATE_TEST_CASE_P(Filter, AdaptiveBilateral, Combine( Values(CV_8UC1, CV_8UC3), Values(5, 9), Values(Size(0, 0)), // not used - Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, - (int)BORDER_REFLECT, (int)BORDER_REFLECT_101), + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(0.0), // not used Bool())); INSTANTIATE_TEST_CASE_P(Filter, MedianFilter, Combine( @@ -447,6 +463,7 @@ INSTANTIATE_TEST_CASE_P(Filter, MedianFilter, Combine( Values(3, 5), Values(Size(0, 0)), // not used Values(0), // not used + Values(0.0), // not used Bool())); #endif // HAVE_OPENCL diff --git a/modules/ocl/test/utility.hpp b/modules/ocl/test/utility.hpp index 47da0134d..1970572fb 100644 --- a/modules/ocl/test/utility.hpp +++ b/modules/ocl/test/utility.hpp @@ -72,6 +72,13 @@ double checkNorm(const cv::Mat &m); double checkNorm(const cv::Mat &m1, const cv::Mat &m2); double checkSimilarity(const cv::Mat &m1, const cv::Mat &m2); +inline double checkNormRelative(const Mat &m1, const Mat &m2) +{ + return cv::norm(m1, m2, cv::NORM_INF) / + std::max((double)std::numeric_limits::epsilon(), + (double)std::max(cv::norm(m1, cv::NORM_INF), norm(m2, cv::NORM_INF))); +} + #define EXPECT_MAT_NORM(mat, eps) \ { \ EXPECT_LE(checkNorm(cv::Mat(mat)), eps) \ @@ -84,6 +91,13 @@ double checkSimilarity(const cv::Mat &m1, const cv::Mat &m2); EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps); \ } +#define EXPECT_MAT_NEAR_RELATIVE(mat1, mat2, eps) \ +{ \ + ASSERT_EQ(mat1.type(), mat2.type()); \ + ASSERT_EQ(mat1.size(), mat2.size()); \ + EXPECT_LE(checkNormRelative(cv::Mat(mat1), cv::Mat(mat2)), eps); \ +} + #define EXPECT_MAT_SIMILAR(mat1, mat2, eps) \ { \ ASSERT_EQ(mat1.type(), mat2.type()); \