Merge pull request #2046 from ilya-lavrenov:tapi_inrange

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<uchar> _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);
 

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;
+            }
+    }
+}

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