diff --git a/modules/core/src/arithm.cpp b/modules/core/src/arithm.cpp index 449303cc31..b58eda1aa9 100644 --- a/modules/core/src/arithm.cpp +++ b/modules/core/src/arithm.cpp @@ -2877,11 +2877,121 @@ static InRangeFunc getInRangeFunc(int depth) return inRangeTab[depth]; } +static bool ocl_inRange( InputArray _src, InputArray _lowerb, + InputArray _upperb, OutputArray _dst ) +{ + int skind = _src.kind(), lkind = _lowerb.kind(), ukind = _upperb.kind(); + Size ssize = _src.size(), lsize = _lowerb.size(), usize = _upperb.size(); + int stype = _src.type(), ltype = _lowerb.type(), utype = _upperb.type(); + int sdepth = CV_MAT_DEPTH(stype), ldepth = CV_MAT_DEPTH(ltype), udepth = CV_MAT_DEPTH(utype); + int cn = CV_MAT_CN(stype); + bool lbScalar = false, ubScalar = false; + + if( (lkind == _InputArray::MATX && skind != _InputArray::MATX) || + ssize != lsize || stype != ltype ) + { + if( !checkScalar(_lowerb, stype, lkind, skind) ) + CV_Error( CV_StsUnmatchedSizes, + "The lower bounary is neither an array of the same size and same type as src, nor a scalar"); + lbScalar = true; + } + + if( (ukind == _InputArray::MATX && skind != _InputArray::MATX) || + ssize != usize || stype != utype ) + { + if( !checkScalar(_upperb, stype, ukind, skind) ) + CV_Error( CV_StsUnmatchedSizes, + "The upper bounary is neither an array of the same size and same type as src, nor a scalar"); + ubScalar = true; + } + + if (lbScalar != ubScalar) + return false; + + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0, + haveScalar = lbScalar && ubScalar; + + if ( (!doubleSupport && sdepth == CV_64F) || + (!haveScalar && (sdepth != ldepth || sdepth != udepth)) ) + return false; + + ocl::Kernel ker("inrange", ocl::core::inrange_oclsrc, + format("%s-D cn=%d -D T=%s%s", haveScalar ? "-D HAVE_SCALAR " : "", + cn, ocl::typeToStr(sdepth), doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (ker.empty()) + return false; + + _dst.create(ssize, CV_8UC1); + UMat src = _src.getUMat(), dst = _dst.getUMat(), lscalaru, uscalaru; + Mat lscalar, uscalar; + + if (lbScalar && ubScalar) + { + lscalar = _lowerb.getMat(); + uscalar = _upperb.getMat(); + + size_t esz = src.elemSize(); + size_t blocksize = 36; + + AutoBuffer _buf(blocksize*(((int)lbScalar + (int)ubScalar)*esz + cn) + 2*cn*sizeof(int) + 128); + uchar *buf = alignPtr(_buf + blocksize*cn, 16); + + if( ldepth != sdepth && sdepth < CV_32S ) + { + int* ilbuf = (int*)alignPtr(buf + blocksize*esz, 16); + int* iubuf = ilbuf + cn; + + BinaryFunc sccvtfunc = getConvertFunc(ldepth, CV_32S); + sccvtfunc(lscalar.data, 0, 0, 0, (uchar*)ilbuf, 0, Size(cn, 1), 0); + sccvtfunc(uscalar.data, 0, 0, 0, (uchar*)iubuf, 0, Size(cn, 1), 0); + int minval = cvRound(getMinVal(sdepth)), maxval = cvRound(getMaxVal(sdepth)); + + for( int k = 0; k < cn; k++ ) + { + if( ilbuf[k] > iubuf[k] || ilbuf[k] > maxval || iubuf[k] < minval ) + ilbuf[k] = minval+1, iubuf[k] = minval; + } + lscalar = Mat(cn, 1, CV_32S, ilbuf); + uscalar = Mat(cn, 1, CV_32S, iubuf); + } + + lscalar.convertTo(lscalar, stype); + uscalar.convertTo(uscalar, stype); + } + else + { + lscalaru = _lowerb.getUMat(); + uscalaru = _upperb.getUMat(); + } + + ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src), + dstarg = ocl::KernelArg::WriteOnly(dst); + + if (haveScalar) + { + lscalar.copyTo(lscalaru); + uscalar.copyTo(uscalaru); + + ker.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(lscalaru), + ocl::KernelArg::PtrReadOnly(uscalaru)); + } + else + ker.args(srcarg, dstarg, ocl::KernelArg::ReadOnlyNoSize(lscalaru), + ocl::KernelArg::ReadOnlyNoSize(uscalaru)); + + size_t globalsize[2] = { ssize.width, ssize.height }; + return ker.run(2, globalsize, NULL, false); +} + } void cv::inRange(InputArray _src, InputArray _lowerb, InputArray _upperb, OutputArray _dst) { + if (ocl::useOpenCL() && _src.dims() <= 2 && _lowerb.dims() <= 2 && + _upperb.dims() <= 2 && _dst.isUMat() && ocl_inRange(_src, _lowerb, _upperb, _dst)) + return; + int skind = _src.kind(), lkind = _lowerb.kind(), ukind = _upperb.kind(); Mat src = _src.getMat(), lb = _lowerb.getMat(), ub = _upperb.getMat(); @@ -2905,14 +3015,14 @@ void cv::inRange(InputArray _src, InputArray _lowerb, ubScalar = true; } - CV_Assert( ((int)lbScalar ^ (int)ubScalar) == 0 ); + CV_Assert(lbScalar == ubScalar); int cn = src.channels(), depth = src.depth(); size_t esz = src.elemSize(); size_t blocksize0 = (size_t)(BLOCK_SIZE + esz-1)/esz; - _dst.create(src.dims, src.size, CV_8U); + _dst.create(src.dims, src.size, CV_8UC1); Mat dst = _dst.getMat(); InRangeFunc func = getInRangeFunc(depth); diff --git a/modules/core/src/opencl/inrange.cl b/modules/core/src/opencl/inrange.cl new file mode 100644 index 0000000000..7549cf3949 --- /dev/null +++ b/modules/core/src/opencl/inrange.cl @@ -0,0 +1,89 @@ +/*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) 2013, OpenCV Foundation, 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 copyright holders 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 inrange(__global const uchar * src1ptr, int src1_step, int src1_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, +#ifdef HAVE_SCALAR + __global const T * src2, __global const T * src3 +#else + __global const uchar * src2ptr, int src2_step, int src2_offset, + __global const uchar * src3ptr, int src3_step, int src3_offset +#endif + ) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < dst_cols && y < dst_rows) + { + int src1_index = mad24(y, src1_step, x*(int)sizeof(T)*cn + src1_offset); + int dst_index = mad24(y, dst_step, x + dst_offset); + __global const T * src1 = (__global const T *)(src1ptr + src1_index); + __global uchar * dst = dstptr + dst_index; + +#ifndef HAVE_SCALAR + int src2_index = mad24(y, src2_step, x*(int)sizeof(T)*cn + src2_offset); + int src3_index = mad24(y, src3_step, x*(int)sizeof(T)*cn + src3_offset); + __global const T * src2 = (__global const T *)(src2ptr + src2_index); + __global const T * src3 = (__global const T *)(src3ptr + src3_index); +#endif + + dst[0] = 255; + + #pragma unroll + for (int c = 0; c < cn; ++c) + if ( src2[c] > src1[c] || src3[c] < src1[c] ) + { + dst[0] = 0; + break; + } + } +} diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp index 58edceccd2..7bc0b5ac0e 100644 --- a/modules/core/test/ocl/test_arithm.cpp +++ b/modules/core/test/ocl/test_arithm.cpp @@ -1241,6 +1241,89 @@ OCL_TEST_P(Normalize, Mat) } } +//////////////////////////////////////// InRange /////////////////////////////////////////////// + +PARAM_TEST_CASE(InRange, MatDepth, Channels, bool /*Scalar or not*/, bool /*Roi*/) +{ + int depth; + int cn; + bool scalars, use_roi; + cv::Scalar val1, val2; + + TEST_DECLARE_INPUT_PARAMETER(src1) + TEST_DECLARE_INPUT_PARAMETER(src2) + TEST_DECLARE_INPUT_PARAMETER(src3) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + depth = GET_PARAM(0); + cn = GET_PARAM(1); + scalars = GET_PARAM(2); + use_roi = GET_PARAM(3); + } + + virtual void generateTestData() + { + const int type = CV_MAKE_TYPE(depth, cn); + + Size roiSize = randomSize(1, MAX_VALUE); + Border src1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src1, src1_roi, roiSize, src1Border, type, -40, 40); + + Border src2Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src2, src2_roi, roiSize, src2Border, type, -40, 40); + + Border src3Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src3, src3_roi, roiSize, src3Border, type, -40, 40); + + Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, CV_8UC1, 5, 16); + + val1 = cv::Scalar(rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0), + rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0)); + val2 = cv::Scalar(rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0), + rng.uniform(-100.0, 100.0), rng.uniform(-100.0, 100.0)); + + UMAT_UPLOAD_INPUT_PARAMETER(src1) + UMAT_UPLOAD_INPUT_PARAMETER(src2) + UMAT_UPLOAD_INPUT_PARAMETER(src3) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + } + + void Near() + { + OCL_EXPECT_MATS_NEAR(dst, 0) + } +}; + +OCL_TEST_P(InRange, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::inRange(src1_roi, src2_roi, src3_roi, dst_roi)); + OCL_ON(cv::inRange(usrc1_roi, usrc2_roi, usrc3_roi, udst_roi)); + + Near(); + } +} + +OCL_TEST_P(InRange, Scalar) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::inRange(src1_roi, val1, val2, dst_roi)); + OCL_ON(cv::inRange(usrc1_roi, val1, val2, udst_roi)); + + Near(); + } +} + + //////////////////////////////////////// Instantiation ///////////////////////////////////////// OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool())); @@ -1276,7 +1359,7 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx_Mask, Combine(OCL_ALL_DEPTHS, ::te OCL_INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); OCL_INSTANTIATE_TEST_CASE_P(Arithm, Sqrt, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); OCL_INSTANTIATE_TEST_CASE_P(Arithm, Normalize, Combine(OCL_ALL_DEPTHS, Values(Channels(1)), Bool())); - +OCL_INSTANTIATE_TEST_CASE_P(Arithm, InRange, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool())); } } // namespace cvtest::ocl