Enable OpenCL version of norm and convertScaleAbs or 32F data

Fix error in minmaxloc.cl
Change test for convertScaleAbs
Fix minMaxIdx for _src2 align
Change epsilon on the tests
pull/3088/head
vbystricky 11 years ago
parent 12c69ad490
commit 1fe403f461
  1. 8
      modules/core/src/convert.cpp
  2. 4
      modules/core/src/opencl/minmaxloc.cl
  3. 8
      modules/core/src/stat.cpp
  4. 58
      modules/core/test/ocl/test_arithm.cpp

@ -1760,13 +1760,12 @@ static bool ocl_convertScaleAbs( InputArray _src, OutputArray _dst, double alpha
kercn = ocl::predictOptimalVectorWidth(_src, _dst), rowsPerWI = d.isIntel() ? 4 : 1;
bool doubleSupport = d.doubleFPConfig() > 0;
if (depth == CV_32F || depth == CV_64F)
if (!doubleSupport && depth == CV_64F)
return false;
char cvt[2][50];
int wdepth = std::max(depth, CV_32F);
ocl::Kernel k("KF", ocl::core::arithm_oclsrc,
format("-D OP_CONVERT_SCALE_ABS -D UNARY_OP -D dstT=%s -D srcT1=%s"
String build_opt = format("-D OP_CONVERT_SCALE_ABS -D UNARY_OP -D dstT=%s -D srcT1=%s"
" -D workT=%s -D wdepth=%d -D convertToWT1=%s -D convertToDT=%s"
" -D workT1=%s -D rowsPerWI=%d%s",
ocl::typeToStr(CV_8UC(kercn)),
@ -1775,7 +1774,8 @@ static bool ocl_convertScaleAbs( InputArray _src, OutputArray _dst, double alpha
ocl::convertTypeStr(depth, wdepth, kercn, cvt[0]),
ocl::convertTypeStr(wdepth, CV_8U, kercn, cvt[1]),
ocl::typeToStr(wdepth), rowsPerWI,
doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
doubleSupport ? " -D DOUBLE_SUPPORT" : "");
ocl::Kernel k("KF", ocl::core::arithm_oclsrc, build_opt);
if (k.empty())
return false;

@ -98,7 +98,7 @@
#ifdef OP_CALC2
#define CALC_MAX2(p) \
maxval2 = MAX(maxval2, temp.p);
maxval2 = MAX(maxval2, temp2.p);
#else
#define CALC_MAX2(p)
#endif
@ -196,7 +196,7 @@ __kernel void minmaxloc(__global const uchar * srcptr, int src_step, int src_off
#ifdef HAVE_SRC2
#ifdef HAVE_SRC2_CONT
src2_index = mul24(id, srcTSIZE);
src2_index = id * srcTSIZE; //mul24(id, srcTSIZE);
#else
src2_index = mad24(id / cols, src2_step, mul24(id % cols, srcTSIZE));
#endif

@ -1444,7 +1444,7 @@ static bool ocl_minMaxIdx( InputArray _src, double* minVal, double* maxVal, int*
bool doubleSupport = dev.doubleFPConfig() > 0, haveMask = !_mask.empty(),
haveSrc2 = _src2.kind() != _InputArray::NONE;
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type),
kercn = haveMask ? cn : std::min(4, ocl::predictOptimalVectorWidth(_src));
kercn = haveMask ? cn : std::min(4, ocl::predictOptimalVectorWidth(_src, _src2));
CV_Assert( (cn == 1 && (!haveMask || _mask.type() == CV_8U)) ||
(cn >= 1 && !minLoc && !maxLoc) );
@ -2190,9 +2190,6 @@ static bool ocl_norm( InputArray _src, int normType, InputArray _mask, double &
(!doubleSupport && depth == CV_64F))
return false;
if( depth == CV_32F && (!_mask.empty() || normType == NORM_INF) )
return false;
UMat src = _src.getUMat();
if (normType == NORM_INF)
@ -2548,9 +2545,6 @@ static bool ocl_norm( InputArray _src1, InputArray _src2, int normType, InputArr
normType &= ~NORM_RELATIVE;
bool normsum = normType == NORM_L1 || normType == NORM_L2 || normType == NORM_L2SQR;
if ( !normsum || !_mask.empty() )
return false;
if (normsum)
{
if (!ocl_sum(_src1, sc1, normType == NORM_L2 || normType == NORM_L2SQR ?

@ -341,7 +341,7 @@ OCL_TEST_P(Mul, Mat_Scalar_Scale)
OCL_OFF(cv::multiply(src1_roi, val, dst1_roi, val[0]));
OCL_ON(cv::multiply(usrc1_roi, val, udst1_roi, val[0]));
Near(udst1_roi.depth() >= CV_32F ? 2e-2 : 1);
Near(udst1_roi.depth() >= CV_32F ? 1e-2 : 1);
}
}
@ -397,7 +397,7 @@ OCL_TEST_P(Div, Mat_Scale)
OCL_OFF(cv::divide(src1_roi, src2_roi, dst1_roi, val[0]));
OCL_ON(cv::divide(usrc1_roi, usrc2_roi, udst1_roi, val[0]));
Near(udst1_roi.depth() >= CV_32F ? 2e-2 : 1);
Near(udst1_roi.depth() >= CV_32F ? 4e-3 : 1);
}
}
@ -1303,7 +1303,7 @@ OCL_TEST_P(Norm, NORM_INF_2args)
OCL_OFF(const double cpuRes = cv::norm(src1_roi, src2_roi, type));
OCL_ON(const double gpuRes = cv::norm(usrc1_roi, usrc2_roi, type));
EXPECT_NEAR(cpuRes, gpuRes, 0.2);
EXPECT_PRED3(relativeError, cpuRes, gpuRes, 2e-2);
}
}
@ -1323,7 +1323,7 @@ OCL_TEST_P(Norm, NORM_INF_2args_mask)
OCL_OFF(const double cpuRes = cv::norm(src1_roi, src2_roi, type, mask_roi));
OCL_ON(const double gpuRes = cv::norm(usrc1_roi, usrc2_roi, type, umask_roi));
EXPECT_NEAR(cpuRes, gpuRes, 0.1);
EXPECT_PRED3(relativeError, cpuRes, gpuRes, 2e-2);
}
}
@ -1547,7 +1547,49 @@ OCL_TEST_P(InRange, Scalar)
//////////////////////////////// ConvertScaleAbs ////////////////////////////////////////////////
typedef ArithmTestBase ConvertScaleAbs;
PARAM_TEST_CASE(ConvertScaleAbs, MatDepth, Channels, bool)
{
int depth;
int cn;
bool use_roi;
cv::Scalar val;
TEST_DECLARE_INPUT_PARAMETER(src);
TEST_DECLARE_OUTPUT_PARAMETER(dst);
virtual void SetUp()
{
depth = GET_PARAM(0);
cn = GET_PARAM(1);
use_roi = GET_PARAM(2);
}
virtual void generateTestData()
{
const int stype = CV_MAKE_TYPE(depth, cn);
const int dtype = CV_MAKE_TYPE(CV_8U, cn);
Size roiSize = randomSize(1, MAX_VALUE);
Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src, src_roi, roiSize, srcBorder, stype, 2, 11); // FIXIT: Test with minV, maxV
Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(dst, dst_roi, roiSize, dstBorder, dtype, 5, 16);
val = 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(src);
UMAT_UPLOAD_OUTPUT_PARAMETER(dst);
}
void Near(double threshold = 0.)
{
OCL_EXPECT_MATS_NEAR(dst, threshold);
}
};
OCL_TEST_P(ConvertScaleAbs, Mat)
{
@ -1555,10 +1597,10 @@ OCL_TEST_P(ConvertScaleAbs, Mat)
{
generateTestData();
OCL_OFF(cv::convertScaleAbs(src1_roi, dst1_roi, val[0], val[1]));
OCL_ON(cv::convertScaleAbs(usrc1_roi, udst1_roi, val[0], val[1]));
OCL_OFF(cv::convertScaleAbs(src_roi, dst_roi, val[0], val[1]));
OCL_ON(cv::convertScaleAbs(usrc_roi, udst_roi, val[0], val[1]));
Near(depth <= CV_32S ? 1 : 1e-6);
Near(1);
}
}

Loading…
Cancel
Save