Merge pull request #2033 from ilya-lavrenov:tapi_normalize

This commit is contained in:
Andrey Pavlenko 2013-12-23 19:19:00 +04:00 committed by OpenCV Buildbot
commit 0966e5ffa1
5 changed files with 199 additions and 16 deletions

View File

@ -1357,43 +1357,65 @@ void cv::LUT( InputArray _src, InputArray _lut, OutputArray _dst )
func(ptrs[0], lut.data, ptrs[1], len, cn, lutcn);
}
namespace cv {
static bool ocl_normalize( InputArray _src, OutputArray _dst, InputArray _mask, int rtype,
double scale, double shift )
{
UMat src = _src.getUMat(), dst = _dst.getUMat();
if( _mask.empty() )
src.convertTo( dst, rtype, scale, shift );
else
{
UMat temp;
src.convertTo( temp, rtype, scale, shift );
temp.copyTo( dst, _mask );
}
return true;
}
}
void cv::normalize( InputArray _src, OutputArray _dst, double a, double b,
int norm_type, int rtype, InputArray _mask )
{
Mat src = _src.getMat(), mask = _mask.getMat();
double scale = 1, shift = 0;
if( norm_type == CV_MINMAX )
{
double smin = 0, smax = 0;
double dmin = MIN( a, b ), dmax = MAX( a, b );
minMaxLoc( _src, &smin, &smax, 0, 0, mask );
minMaxLoc( _src, &smin, &smax, 0, 0, _mask );
scale = (dmax - dmin)*(smax - smin > DBL_EPSILON ? 1./(smax - smin) : 0);
shift = dmin - smin*scale;
}
else if( norm_type == CV_L2 || norm_type == CV_L1 || norm_type == CV_C )
{
scale = norm( src, norm_type, mask );
scale = norm( _src, norm_type, _mask );
scale = scale > DBL_EPSILON ? a/scale : 0.;
shift = 0;
}
else
CV_Error( CV_StsBadArg, "Unknown/unsupported norm type" );
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
if( rtype < 0 )
rtype = _dst.fixedType() ? _dst.depth() : src.depth();
rtype = _dst.fixedType() ? _dst.depth() : depth;
_dst.createSameSize(_src, CV_MAKETYPE(rtype, cn));
_dst.create(src.dims, src.size, CV_MAKETYPE(rtype, src.channels()));
Mat dst = _dst.getMat();
if (ocl::useOpenCL() && _dst.isUMat() &&
ocl_normalize(_src, _dst, _mask, rtype, scale, shift))
return;
if( !mask.data )
Mat src = _src.getMat(), dst = _dst.getMat();
if( _mask.empty() )
src.convertTo( dst, rtype, scale, shift );
else
{
Mat temp;
src.convertTo( temp, rtype, scale, shift );
temp.copyTo( dst, mask );
temp.copyTo( dst, _mask );
}
}

View File

@ -41,6 +41,52 @@
//
//M*/
#ifdef COPY_TO_MASK
#define DEFINE_DATA \
int src_index = mad24(y, src_step, x*(int)sizeof(T)*scn + src_offset); \
int dst_index = mad24(y, dst_step, x*(int)sizeof(T)*scn + dst_offset); \
\
__global const T * src = (__global const T *)(srcptr + src_index); \
__global T * dst = (__global T *)(dstptr + dst_index)
__kernel void copyToMask(__global const uchar * srcptr, int src_step, int src_offset,
__global const uchar * maskptr, int mask_step, int mask_offset,
__global uchar * dstptr, int dst_step, int dst_offset,
int dst_rows, int dst_cols)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < dst_cols && y < dst_rows)
{
int mask_index = mad24(y, mask_step, x * mcn + mask_offset);
__global const uchar * mask = (__global const uchar *)(maskptr + mask_index);
#if mcn == 1
if (mask[0])
{
DEFINE_DATA;
#pragma unroll
for (int c = 0; c < scn; ++c)
dst[c] = src[c];
}
#elif scn == mcn
DEFINE_DATA;
#pragma unroll
for (int c = 0; c < scn; ++c)
if (mask[c])
dst[c] = src[c];
#else
#error "(mcn == 1 || mcn == scn) should be true"
#endif
}
}
#else
__kernel void setMask(__global const uchar* mask, int maskstep, int maskoffset,
__global uchar* dstptr, int dststep, int dstoffset,
int rows, int cols, dstT value )
@ -71,3 +117,5 @@ __kernel void set(__global uchar* dstptr, int dststep, int dstoffset,
*(__global dstT*)(dstptr + dst_index) = value;
}
}
#endif

View File

@ -661,6 +661,45 @@ void UMat::copyTo(OutputArray _dst) const
}
}
void UMat::copyTo(OutputArray _dst, InputArray _mask) const
{
if( _mask.empty() )
{
copyTo(_dst);
return;
}
int cn = channels(), mtype = _mask.type(), mdepth = CV_MAT_DEPTH(mtype), mcn = CV_MAT_CN(mtype);
CV_Assert( mdepth == CV_8U && (mcn == 1 || mcn == cn) );
if (ocl::useOpenCL() && _dst.isUMat() && dims <= 2)
{
UMatData * prevu = _dst.getUMat().u;
_dst.create( dims, size, type() );
UMat dst = _dst.getUMat();
if( prevu != dst.u ) // do not leave dst uninitialized
dst = Scalar(0);
ocl::Kernel k("copyToMask", ocl::core::copyset_oclsrc,
format("-D COPY_TO_MASK -D T=%s -D scn=%d -D mcn=%d",
ocl::memopTypeToStr(depth()), cn, mcn));
if (!k.empty())
{
k.args(ocl::KernelArg::ReadOnlyNoSize(*this), ocl::KernelArg::ReadOnlyNoSize(_mask.getUMat()),
ocl::KernelArg::WriteOnly(dst));
size_t globalsize[2] = { cols, rows };
if (k.run(2, globalsize, NULL, false))
return;
}
}
Mat src = getMat(ACCESS_READ);
src.copyTo(_dst, _mask);
}
void UMat::convertTo(OutputArray _dst, int _type, double alpha, double beta) const
{
bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON;

View File

@ -1219,6 +1219,28 @@ OCL_TEST_P(Sqrt, Mat)
}
}
//////////////////////////////// Normalize ////////////////////////////////////////////////
typedef ArithmTestBase Normalize;
OCL_TEST_P(Normalize, Mat)
{
static int modes[] = { CV_MINMAX, CV_L2, CV_L1, CV_C };
for (int j = 0; j < test_loop_times; j++)
{
generateTestData();
for (int i = 0, size = sizeof(modes) / sizeof(modes[0]); i < size; ++i)
{
OCL_OFF(cv::normalize(src1_roi, dst1_roi, 10, 110, modes[i], src1_roi.type(), mask_roi));
OCL_ON(cv::normalize(usrc1_roi, udst1_roi, 10, 110, modes[i], src1_roi.type(), umask_roi));
Near(1);
}
}
}
//////////////////////////////////////// Instantiation /////////////////////////////////////////
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
@ -1253,6 +1275,8 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx, Combine(OCL_ALL_DEPTHS, OCL_ALL_C
OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx_Mask, Combine(OCL_ALL_DEPTHS, ::testing::Values(Channels(1)), Bool()));
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()));
} } // namespace cvtest::ocl

View File

@ -54,7 +54,7 @@ namespace ocl {
////////////////////////////////converto/////////////////////////////////////////////////
PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool)
PARAM_TEST_CASE(ConvertTo, MatDepth, MatDepth, Channels, bool)
{
int src_depth, cn, dstType;
bool use_roi;
@ -85,8 +85,6 @@ PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool)
}
};
typedef MatrixTestBase ConvertTo;
OCL_TEST_P(ConvertTo, Accuracy)
{
for (int j = 0; j < test_loop_times; j++)
@ -103,7 +101,51 @@ OCL_TEST_P(ConvertTo, Accuracy)
}
}
typedef MatrixTestBase CopyTo;
//////////////////////////////// CopyTo /////////////////////////////////////////////////
PARAM_TEST_CASE(CopyTo, MatDepth, Channels, bool, bool)
{
int depth, cn;
bool use_roi, use_mask;
TEST_DECLARE_INPUT_PARAMETER(src)
TEST_DECLARE_INPUT_PARAMETER(mask)
TEST_DECLARE_OUTPUT_PARAMETER(dst)
virtual void SetUp()
{
depth = GET_PARAM(0);
cn = GET_PARAM(1);
use_roi = GET_PARAM(2);
use_mask = GET_PARAM(3);
}
void generateTestData()
{
const int type = CV_MAKE_TYPE(depth, cn);
Size roiSize = randomSize(1, MAX_VALUE);
Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
if (use_mask)
{
Border maskBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
int mask_cn = randomDouble(0.0, 2.0) > 1.0 ? cn : 1;
randomSubMat(mask, mask_roi, roiSize, maskBorder, CV_8UC(mask_cn), 0, 2);
cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
}
Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(dst, dst_roi, roiSize, dstBorder, type, 5, 16);
UMAT_UPLOAD_INPUT_PARAMETER(src)
if (use_mask)
UMAT_UPLOAD_INPUT_PARAMETER(mask)
UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
}
};
OCL_TEST_P(CopyTo, Accuracy)
{
@ -111,8 +153,16 @@ OCL_TEST_P(CopyTo, Accuracy)
{
generateTestData();
OCL_OFF(src_roi.copyTo(dst_roi));
OCL_ON(usrc_roi.copyTo(udst_roi));
if (use_mask)
{
OCL_OFF(src_roi.copyTo(dst_roi, mask_roi));
OCL_ON(usrc_roi.copyTo(udst_roi, umask_roi));
}
else
{
OCL_OFF(src_roi.copyTo(dst_roi));
OCL_ON(usrc_roi.copyTo(udst_roi));
}
OCL_EXPECT_MATS_NEAR(dst, 0);
}
@ -122,7 +172,7 @@ OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, ConvertTo, Combine(
OCL_ALL_DEPTHS, OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, CopyTo, Combine(
OCL_ALL_DEPTHS, Values((MatDepth)0), OCL_ALL_CHANNELS, Bool()));
OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool()));
} } // namespace cvtest::ocl