diff --git a/modules/core/src/copy.cpp b/modules/core/src/copy.cpp index ee34ef451..fa6fa69f1 100644 --- a/modules/core/src/copy.cpp +++ b/modules/core/src/copy.cpp @@ -47,6 +47,7 @@ // */ #include "precomp.hpp" +#include "opencl_kernels.hpp" namespace cv { @@ -701,12 +702,73 @@ void copyMakeConstBorder_8u( const uchar* src, size_t srcstep, cv::Size srcroi, } +namespace cv { + +static bool ocl_copyMakeBorder( InputArray _src, OutputArray _dst, int top, int bottom, + int left, int right, int borderType, const Scalar& value ) +{ + int type = _src.type(), cn = CV_MAT_CN(type); + bool isolated = (borderType & BORDER_ISOLATED) != 0; + borderType &= ~cv::BORDER_ISOLATED; + + if ( !(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE || borderType == BORDER_REFLECT || + borderType == BORDER_WRAP || borderType == BORDER_REFLECT_101) || + cn == 3 || cn > 4) + return false; + + const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" }; + ocl::Kernel k("copyMakeBorder", ocl::core::copymakeborder_oclsrc, + format("-D T=%s -D %s", ocl::memopTypeToStr(type), borderMap[borderType])); + if (k.empty()) + return false; + + UMat src = _src.getUMat(); + if( src.isSubmatrix() && !isolated ) + { + Size wholeSize; + Point ofs; + src.locateROI(wholeSize, ofs); + int dtop = std::min(ofs.y, top); + int dbottom = std::min(wholeSize.height - src.rows - ofs.y, bottom); + int dleft = std::min(ofs.x, left); + int dright = std::min(wholeSize.width - src.cols - ofs.x, right); + src.adjustROI(dtop, dbottom, dleft, dright); + top -= dtop; + left -= dleft; + bottom -= dbottom; + right -= dright; + } + + _dst.create(src.rows + top + bottom, src.cols + left + right, type); + UMat dst = _dst.getUMat(); + + if (top == 0 && left == 0 && bottom == 0 && right == 0) + { + if(src.u != dst.u || src.step != dst.step) + src.copyTo(dst); + return true; + } + + k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst), + top, left, ocl::KernelArg::Constant(Mat(1, 1, type, value))); + + size_t globalsize[2] = { dst.cols, dst.rows }; + return k.run(2, globalsize, NULL, false); +} + +} + void cv::copyMakeBorder( InputArray _src, OutputArray _dst, int top, int bottom, int left, int right, int borderType, const Scalar& value ) { - Mat src = _src.getMat(); CV_Assert( top >= 0 && bottom >= 0 && left >= 0 && right >= 0 ); + if (ocl::useOpenCL() && _dst.isUMat() && _src.dims() <= 2 && + ocl_copyMakeBorder(_src, _dst, top, bottom, left, right, borderType, value)) + return; + + Mat src = _src.getMat(); + if( src.isSubmatrix() && (borderType & BORDER_ISOLATED) == 0 ) { Size wholeSize; diff --git a/modules/core/src/opencl/copymakeborder.cl b/modules/core/src/opencl/copymakeborder.cl new file mode 100644 index 000000000..bb264421f --- /dev/null +++ b/modules/core/src/opencl/copymakeborder.cl @@ -0,0 +1,132 @@ +// 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 +// Niko Li, newlife20080214@gmail.com +// Zero Lin zero.lin@amd.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. +// +// + +#ifdef DOUBLE_SUPPORT +#ifdef cl_amd_fp64 +#pragma OPENCL EXTENSION cl_amd_fp64:enable +#elif defined (cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#endif +#endif + +#ifdef BORDER_CONSTANT +#define EXTRAPOLATE(x, y, v) v = scalar; +#elif defined BORDER_REPLICATE +#define EXTRAPOLATE(x, y, v) \ + { \ + x = max(min(x, src_cols - 1), 0); \ + y = max(min(y, src_rows - 1), 0); \ + v = *(__global const T *)(srcptr + mad24(y, src_step, x * (int)sizeof(T) + src_offset)); \ + } +#elif defined BORDER_WRAP +#define EXTRAPOLATE(x, y, v) \ + { \ + if (x < 0) \ + x -= ((x - src_cols + 1) / src_cols) * src_cols; \ + if (x >= src_cols) \ + x %= src_cols; \ + \ + if (y < 0) \ + y -= ((y - src_rows + 1) / src_rows) * src_rows; \ + if( y >= src_rows ) \ + y %= src_rows; \ + v = *(__global const T *)(srcptr + mad24(y, src_step, x * (int)sizeof(T) + src_offset)); \ + } +#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) +#ifdef BORDER_REFLECT +#define DELTA int delta = 0 +#else +#define DELTA int delta = 1 +#endif +#define EXTRAPOLATE(x, y, v) \ + { \ + DELTA; \ + if (src_cols == 1) \ + x = 0; \ + else \ + do \ + { \ + if( x < 0 ) \ + x = -x - 1 + delta; \ + else \ + x = src_cols - 1 - (x - src_cols) - delta; \ + } \ + while (x >= src_cols || x < 0); \ + \ + if (src_rows == 1) \ + y = 0; \ + else \ + do \ + { \ + if( y < 0 ) \ + y = -y - 1 + delta; \ + else \ + y = src_rows - 1 - (y - src_rows) - delta; \ + } \ + while (y >= src_rows || y < 0); \ + v = *(__global const T *)(srcptr + mad24(y, src_step, x * (int)sizeof(T) + src_offset)); \ + } +#else +#error No extrapolation method +#endif + +#define NEED_EXTRAPOLATION(gx, gy) (gx >= src_cols || gy >= src_rows || gx < 0 || gy < 0) + +__kernel void copyMakeBorder(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols, + __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, + int top, int left, T scalar) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < dst_cols && y < dst_rows) + { + int src_x = x - left; + int src_y = y - top; + + int dst_index = mad24(y, dst_step, x * (int)sizeof(T) + dst_offset); + __global T * dst = (__global T *)(dstptr + dst_index); + + if (NEED_EXTRAPOLATION(src_x, src_y)) + EXTRAPOLATE(src_x, src_y, dst[0]) + else + { + int src_index = mad24(src_y, src_step, src_x * (int)sizeof(T) + src_offset); + __global const T * src = (__global const T *)(srcptr + src_index); + dst[0] = src[0]; + } + } +} diff --git a/modules/core/src/umatrix.cpp b/modules/core/src/umatrix.cpp index c35d4ccfa..e6d9c9e38 100644 --- a/modules/core/src/umatrix.cpp +++ b/modules/core/src/umatrix.cpp @@ -580,6 +580,7 @@ Mat UMat::getMat(int accessFlags) const Mat hdr(dims, size.p, type(), u->data + offset, step.p); hdr.flags = flags; hdr.u = u; + hdr.flags = flags; hdr.datastart = u->data; hdr.data = hdr.datastart + offset; hdr.datalimit = hdr.dataend = u->data + u->size; diff --git a/modules/imgproc/src/opencl/threshold.cl b/modules/imgproc/src/opencl/threshold.cl new file mode 100644 index 000000000..f5b6fbb1d --- /dev/null +++ b/modules/imgproc/src/opencl/threshold.cl @@ -0,0 +1,81 @@ +/*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 +// 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: +// +// * 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 DOUBLE_SUPPORT +#ifdef cl_amd_fp64 +#pragma OPENCL EXTENSION cl_amd_fp64:enable +#elif defined (cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#endif +#endif + +__kernel void threshold(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols, + T thresh, T max_val) +{ + int gx = get_global_id(0); + int gy = get_global_id(1); + + if (gx < cols && gy < rows) + { + int src_index = mad24(gy, src_step, src_offset + gx * (int)sizeof(T)); + int dst_index = mad24(gy, dst_step, dst_offset + gx * (int)sizeof(T)); + + T sdata = *(__global const T *)(srcptr + src_index); + __global T * dst = (__global T *)(dstptr + dst_index); + +#ifdef THRESH_BINARY + dst[0] = sdata > thresh ? max_val : (T)(0); +#elif defined THRESH_BINARY_INV + dst[0] = sdata > thresh ? (T)(0) : max_val; +#elif defined THRESH_TRUNC + dst[0] = sdata > thresh ? thresh : sdata; +#elif defined THRESH_TOZERO + dst[0] = sdata > thresh ? sdata : (T)(0); +#elif defined THRESH_TOZERO_INV + dst[0] = sdata > thresh ? (T)(0) : sdata; +#endif + } +} diff --git a/modules/imgproc/src/thresh.cpp b/modules/imgproc/src/thresh.cpp index 6e80b3b2a..ce853a783 100644 --- a/modules/imgproc/src/thresh.cpp +++ b/modules/imgproc/src/thresh.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" namespace cv { @@ -705,10 +706,47 @@ private: int thresholdType; }; +static bool ocl_threshold( InputArray _src, OutputArray _dst, double & thresh, double maxval, int thresh_type ) +{ + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), ktype = CV_MAKE_TYPE(depth, 1); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if ( !(thresh_type == THRESH_BINARY || thresh_type == THRESH_BINARY_INV || thresh_type == THRESH_TRUNC || + thresh_type == THRESH_TOZERO || thresh_type == THRESH_TOZERO_INV) || + (!doubleSupport && depth == CV_64F)) + return false; + + const char * const thresholdMap[] = { "THRESH_BINARY", "THRESH_BINARY_INV", "THRESH_TRUNC", + "THRESH_TOZERO", "THRESH_TOZERO_INV" }; + ocl::Kernel k("threshold", ocl::imgproc::threshold_oclsrc, + format("-D %s -D T=%s%s", thresholdMap[thresh_type], + ocl::typeToStr(ktype), doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (k.empty()) + return false; + + UMat src = _src.getUMat(); + _dst.create(src.size(), type); + UMat dst = _dst.getUMat(); + + if (depth <= CV_32S) + thresh = cvFloor(thresh); + + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst, cn), + ocl::KernelArg::Constant(Mat(1, 1, ktype, thresh)), + ocl::KernelArg::Constant(Mat(1, 1, ktype, maxval))); + + size_t globalsize[2] = { dst.cols * cn, dst.rows }; + return k.run(2, globalsize, NULL, false); +} + } double cv::threshold( InputArray _src, OutputArray _dst, double thresh, double maxval, int type ) { + if (ocl::useOpenCL() && _src.dims() <= 2 && _dst.isUMat() && + ocl_threshold(_src, _dst, thresh, maxval, type)) + return thresh; + Mat src = _src.getMat(); bool use_otsu = (type & THRESH_OTSU) != 0; type &= THRESH_MASK; diff --git a/modules/imgproc/test/ocl/test_imgproc.cpp b/modules/imgproc/test/ocl/test_imgproc.cpp new file mode 100644 index 000000000..ae833406f --- /dev/null +++ b/modules/imgproc/test/ocl/test_imgproc.cpp @@ -0,0 +1,455 @@ +/*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. +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Niko Li, newlife20080214@gmail.com +// Jia Haipeng, jiahaipeng95@gmail.com +// Shengen Yan, yanshengen@gmail.com +// Jiang Liyuan, lyuan001.good@163.com +// Rock Li, Rock.Li@amd.com +// Wu Zailong, bullet@yeah.net +// Xu Pang, pangxu010@163.com +// Sen Liu, swjtuls1987@126.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*/ + +#include "test_precomp.hpp" +#include "cvconfig.h" +#include "opencv2/ts/ocl_test.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +/////////////////////////////////////////////////////////////////////////////// + +PARAM_TEST_CASE(ImgprocTestBase, MatType, + int, // blockSize + int, // border type + bool) // roi or not +{ + int type, borderType, blockSize; + bool useRoi; + + TEST_DECLARE_INPUT_PARAMETER(src) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + type = GET_PARAM(0); + blockSize = GET_PARAM(1); + borderType = GET_PARAM(2); + useRoi = GET_PARAM(3); + } + + virtual void random_roi() + { + Size roiSize = randomSize(1, MAX_VALUE); + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, type, 5, 16); + + UMAT_UPLOAD_INPUT_PARAMETER(src) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + } + + void Near(double threshold = 0.0, bool relative = false) + { + if (relative) + { + EXPECT_MAT_NEAR_RELATIVE(dst, udst, threshold); + EXPECT_MAT_NEAR_RELATIVE(dst_roi, udst_roi, threshold); + } + else + { + EXPECT_MAT_NEAR(dst, udst, threshold); + EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold); + } + } +}; + +////////////////////////////////copyMakeBorder//////////////////////////////////////////// + +PARAM_TEST_CASE(CopyMakeBorder, MatDepth, // depth + Channels, // channels + bool, // isolated or not + BorderType, // border type + bool) // roi or not +{ + int type, borderType; + bool useRoi; + + TestUtils::Border border; + Scalar val; + + TEST_DECLARE_INPUT_PARAMETER(src) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1)); + borderType = GET_PARAM(3); + + if (GET_PARAM(2)) + borderType |= BORDER_ISOLATED; + + useRoi = GET_PARAM(4); + } + + void random_roi() + { + border = randomBorder(0, MAX_VALUE << 2); + val = randomScalar(-MAX_VALUE, MAX_VALUE); + + Size roiSize = randomSize(1, MAX_VALUE); + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + dstBorder.top += border.top; + dstBorder.lef += border.lef; + dstBorder.rig += border.rig; + dstBorder.bot += border.bot; + + randomSubMat(dst, dst_roi, roiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE); + + UMAT_UPLOAD_INPUT_PARAMETER(src) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + } + + void Near(double threshold = 0.0) + { + EXPECT_MAT_NEAR(dst, udst, threshold); + EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold); + } +}; + +OCL_TEST_P(CopyMakeBorder, Mat) +{ + for (int i = 0; i < test_loop_times; ++i) + { + random_roi(); + + OCL_OFF(cv::copyMakeBorder(src_roi, dst_roi, border.top, border.bot, border.lef, border.rig, borderType, val)); + OCL_ON(cv::copyMakeBorder(usrc_roi, udst_roi, border.top, border.bot, border.lef, border.rig, borderType, val)); + + Near(); + } +} + +////////////////////////////////equalizeHist////////////////////////////////////////////// + +typedef ImgprocTestBase EqualizeHist; + +OCL_TEST_P(EqualizeHist, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + OCL_OFF(cv::equalizeHist(src_roi, dst_roi)); + OCL_ON(cv::equalizeHist(usrc_roi, udst_roi)); + + Near(1.1); + } +} + +////////////////////////////////cornerMinEigenVal////////////////////////////////////////// + +struct CornerTestBase : + public ImgprocTestBase +{ + virtual void random_roi() + { + Mat image = readImageType("gpu/stereobm/aloe-L.png", type); + ASSERT_FALSE(image.empty()); + + bool isFP = CV_MAT_DEPTH(type) >= CV_32F; + float val = 255.0f; + if (isFP) + { + image.convertTo(image, -1, 1.0 / 255); + val /= 255.0f; + } + + Size roiSize = image.size(); + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + + Size wholeSize = Size(roiSize.width + srcBorder.lef + srcBorder.rig, roiSize.height + srcBorder.top + srcBorder.bot); + src = randomMat(wholeSize, type, -val, val, false); + src_roi = src(Rect(srcBorder.lef, srcBorder.top, roiSize.width, roiSize.height)); + image.copyTo(src_roi); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, CV_32FC1, 5, 16); + + UMAT_UPLOAD_INPUT_PARAMETER(src) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + } +}; + +typedef CornerTestBase CornerMinEigenVal; + +OCL_TEST_P(CornerMinEigenVal, DISABLED_Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + int apertureSize = 3; + + OCL_OFF(cv::cornerMinEigenVal(src_roi, dst_roi, blockSize, apertureSize, borderType)); + OCL_ON(cv::cornerMinEigenVal(usrc_roi, udst_roi, blockSize, apertureSize, borderType)); + + Near(1e-5, true); + } +} + +////////////////////////////////cornerHarris////////////////////////////////////////// + +typedef CornerTestBase CornerHarris; + +OCL_TEST_P(CornerHarris, DISABLED_Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + int apertureSize = 3; + double k = randomDouble(0.01, 0.9); + + OCL_OFF(cv::cornerHarris(src_roi, dst_roi, blockSize, apertureSize, k, borderType)); + OCL_ON(cv::cornerHarris(usrc_roi, udst_roi, blockSize, apertureSize, k, borderType)); + + Near(1e-5, true); + } +} + +//////////////////////////////////integral///////////////////////////////////////////////// + +struct Integral : + public ImgprocTestBase +{ + int sdepth; + + virtual void SetUp() + { + type = GET_PARAM(0); + blockSize = GET_PARAM(1); + sdepth = GET_PARAM(2); + useRoi = GET_PARAM(3); + } +}; + +OCL_TEST_P(Integral, Mat1) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + OCL_OFF(cv::integral(src_roi, dst_roi, sdepth)); + OCL_ON(cv::integral(usrc_roi, udst_roi, sdepth)); + + Near(); + } +} + +OCL_TEST_P(Integral, Mat2) +{ + Mat dst1; + UMat udst1; + + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + OCL_OFF(cv::integral(src_roi, dst_roi, dst1, sdepth)); + OCL_ON(cv::integral(usrc_roi, udst_roi, udst1, sdepth)); + + Near(); + if (cv::ocl::Device::getDefault().doubleFPConfig() > 0) + EXPECT_MAT_NEAR(dst1, udst1, 0.); + } +} + +/////////////////////////////////////////////////////////////////////////////////////////////////// +//// threshold + +struct Threshold : + public ImgprocTestBase +{ + int thresholdType; + + virtual void SetUp() + { + type = GET_PARAM(0); + blockSize = GET_PARAM(1); + thresholdType = GET_PARAM(2); + useRoi = GET_PARAM(3); + } +}; + +OCL_TEST_P(Threshold, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + double maxVal = randomDouble(20.0, 127.0); + double thresh = randomDouble(0.0, maxVal); + + OCL_OFF(cv::threshold(src_roi, dst_roi, thresh, maxVal, thresholdType)); + OCL_ON(cv::threshold(usrc_roi, udst_roi, thresh, maxVal, thresholdType)); + + Near(1); + } +} + + +///////////////////////////////////////////////////////////////////////////////////////////////////////// +//// CLAHE + +PARAM_TEST_CASE(CLAHETest, Size, double, bool) +{ + Size gridSize; + double clipLimit; + bool useRoi; + + TEST_DECLARE_INPUT_PARAMETER(src) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + gridSize = GET_PARAM(0); + clipLimit = GET_PARAM(1); + useRoi = GET_PARAM(2); + } + + void random_roi() + { + Size roiSize = randomSize(std::max(gridSize.height, gridSize.width), MAX_VALUE); + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, CV_8UC1, 5, 256); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, CV_8UC1, 5, 16); + + UMAT_UPLOAD_INPUT_PARAMETER(src) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + } + + void Near(double threshold = 0.0) + { + EXPECT_MAT_NEAR(dst, udst, threshold); + EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold); + } +}; + +OCL_TEST_P(CLAHETest, Accuracy) +{ + for (int i = 0; i < test_loop_times; ++i) + { + random_roi(); + + Ptr clahe = cv::createCLAHE(clipLimit, gridSize); + + OCL_OFF(clahe->apply(src_roi, dst_roi)); + OCL_ON(clahe->apply(usrc_roi, udst_roi)); + + Near(1.0); + } +} + +///////////////////////////////////////////////////////////////////////////////////// + +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, EqualizeHist, Combine( + Values((MatType)CV_8UC1), + Values(0), // not used + Values(0), // not used + Bool())); + +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, CornerMinEigenVal, Combine( + Values((MatType)CV_8UC1, (MatType)CV_32FC1), + Values(3, 5), + Values((int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_REFLECT101), + Bool())); + +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, CornerHarris, Combine( + Values((MatType)CV_8UC1, CV_32FC1), + Values(3, 5), + Values( (int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_REFLECT_101), + Bool())); + +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, Integral, Combine( + Values((MatType)CV_8UC1), // TODO does not work with CV_32F, CV_64F + Values(0), // not used + Values((MatType)CV_32SC1, (MatType)CV_32FC1), + Bool())); + +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, Threshold, Combine( + Values(CV_8UC1, CV_8UC2, CV_8UC3, CV_8UC4, + CV_16SC1, CV_16SC2, CV_16SC3, CV_16SC4, + CV_32FC1, CV_32FC2, CV_32FC3, CV_32FC4), + Values(0), + Values(ThreshOp(THRESH_BINARY), + ThreshOp(THRESH_BINARY_INV), ThreshOp(THRESH_TRUNC), + ThreshOp(THRESH_TOZERO), ThreshOp(THRESH_TOZERO_INV)), + Bool())); + +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, CLAHETest, Combine( + Values(Size(4, 4), Size(32, 8), Size(8, 64)), + Values(0.0, 10.0, 62.0, 300.0), + Bool())); + +OCL_INSTANTIATE_TEST_CASE_P(ImgprocTestBase, CopyMakeBorder, Combine( + testing::Values((MatDepth)CV_8U, (MatDepth)CV_16S, (MatDepth)CV_32S, (MatDepth)CV_32F), + testing::Values(Channels(1), Channels(3), (Channels)4), + Bool(), // border isolated or not + Values((BorderType)BORDER_CONSTANT, (BorderType)BORDER_REPLICATE, (BorderType)BORDER_REFLECT, + (BorderType)BORDER_WRAP, (BorderType)BORDER_REFLECT_101), + Bool())); + +} } // namespace cvtest::ocl + +#endif // HAVE_OPENCL diff --git a/modules/imgproc/test/ocl/test_warp.cpp b/modules/imgproc/test/ocl/test_warp.cpp index 9c22e17f7..c05c33590 100644 --- a/modules/imgproc/test/ocl/test_warp.cpp +++ b/modules/imgproc/test/ocl/test_warp.cpp @@ -224,7 +224,7 @@ OCL_TEST_P(Resize, Mat) ///////////////////////////////////////////////////////////////////////////////////////////////// // remap -PARAM_TEST_CASE(Remap, MatDepth, Channels, std::pair, Border, bool) +PARAM_TEST_CASE(Remap, MatDepth, Channels, std::pair, BorderType, bool) { int srcType, map1Type, map2Type; int borderType; @@ -349,11 +349,11 @@ OCL_INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_LINEAR, Combine( Values(std::pair((MatType)CV_32FC1, (MatType)CV_32FC1), std::pair((MatType)CV_16SC2, (MatType)CV_16UC1), std::pair((MatType)CV_32FC2, noType)), - Values((Border)BORDER_CONSTANT, - (Border)BORDER_REPLICATE, - (Border)BORDER_WRAP, - (Border)BORDER_REFLECT, - (Border)BORDER_REFLECT_101), + Values((BorderType)BORDER_CONSTANT, + (BorderType)BORDER_REPLICATE, + (BorderType)BORDER_WRAP, + (BorderType)BORDER_REFLECT, + (BorderType)BORDER_REFLECT_101), Bool())); OCL_INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_NEAREST, Combine( @@ -363,11 +363,11 @@ OCL_INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_NEAREST, Combine( std::pair((MatType)CV_32FC2, noType), std::pair((MatType)CV_16SC2, (MatType)CV_16UC1), std::pair((MatType)CV_16SC2, noType)), - Values((Border)BORDER_CONSTANT, - (Border)BORDER_REPLICATE, - (Border)BORDER_WRAP, - (Border)BORDER_REFLECT, - (Border)BORDER_REFLECT_101), + Values((BorderType)BORDER_CONSTANT, + (BorderType)BORDER_REPLICATE, + (BorderType)BORDER_WRAP, + (BorderType)BORDER_REFLECT, + (BorderType)BORDER_REFLECT_101), Bool())); } } // namespace cvtest::ocl diff --git a/modules/ts/include/opencv2/ts/ocl_test.hpp b/modules/ts/include/opencv2/ts/ocl_test.hpp index 75517c93f..1b66799ed 100644 --- a/modules/ts/include/opencv2/ts/ocl_test.hpp +++ b/modules/ts/include/opencv2/ts/ocl_test.hpp @@ -264,7 +264,7 @@ struct CV_EXPORTS TestUtils #define UMAT_UPLOAD_INPUT_PARAMETER(name) \ { \ name.copyTo(u ## name); \ - Size wholeSize; Point ofs; name ## _roi.locateROI(wholeSize, ofs); \ + Size _wholeSize; Point ofs; name ## _roi.locateROI(_wholeSize, ofs); \ u ## name ## _roi = u ## name(Rect(ofs.x, ofs.y, name ## _roi.size().width, name ## _roi.size().height)); \ } #define UMAT_UPLOAD_OUTPUT_PARAMETER(name) UMAT_UPLOAD_INPUT_PARAMETER(name) @@ -306,7 +306,8 @@ IMPLEMENT_PARAM_CLASS(Channels, int) #define OCL_ALL_CHANNELS Values(1, 2, 3, 4) CV_ENUM(Interpolation, INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_AREA) -CV_ENUM(Border, BORDER_CONSTANT, BORDER_REPLICATE, BORDER_WRAP, BORDER_REFLECT, BORDER_REFLECT_101) +CV_ENUM(ThreshOp, THRESH_BINARY, THRESH_BINARY_INV, THRESH_TRUNC, THRESH_TOZERO, THRESH_TOZERO_INV) +CV_ENUM(BorderType, BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101) #define OCL_INSTANTIATE_TEST_CASE_P(prefix, test_case_name, generator) \ INSTANTIATE_TEST_CASE_P(OCL_ ## prefix, test_case_name, generator)