mirror of
https://github.com/opencv/opencv.git
synced 2024-11-29 13:47:32 +08:00
added cv::inRange to T-API
This commit is contained in:
Ilya Lavrenov
parent
0966e5ffa1
commit
4c23059209
@ -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);
|
||||
|
||||
|
||||
89
modules/core/src/opencl/inrange.cl
Normal file
89
modules/core/src/opencl/inrange.cl
Normal file
@ -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;
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -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
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user