mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2024-11-29 13:47:32 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Merge pull request #1866 from ilya-lavrenov:ocl_fix2

Browse Source
This commit is contained in:
Andrey Pavlenko 2013-11-26 23:27:48 +04:00 committed by OpenCV Buildbot
commit 96a228fd7e
6 changed files with 109 additions and 88 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

25
modules/core/include/opencv2/core/mat.hpp
View File

@ -2340,31 +2340,6 @@ CV_EXPORTS MatExpr max(double s, const Mat& a);
CV_EXPORTS MatExpr abs(const Mat& m);
CV_EXPORTS MatExpr abs(const MatExpr& e);
namespace traits {
template <typename T>
struct GetMatForRead
{
};
template <>
struct GetMatForRead<Mat>
{
static const Mat get(const Mat& m) { return m; }
};
template <>
struct GetMatForRead<UMat>
{
static const Mat get(const UMat& m) { return m.getMat(ACCESS_READ); }
};
} // namespace traits
template <typename T>
const Mat getMatForRead(const T& mat)
{
return traits::GetMatForRead<T>::get(mat);
}
} // cv
#include "opencv2/core/mat.inl.hpp"

60
modules/core/src/arithm.cpp
View File

@ -1283,47 +1283,29 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
int type1 = _src1.type(), depth1 = CV_MAT_DEPTH(type1), cn = CV_MAT_CN(type1);
bool haveMask = !_mask.empty();
if( (haveMask || haveScalar) && cn > 4 )
if( ((haveMask || haveScalar) && cn > 4) || cn == 3) // TODO need fix for 3 channels
return false;
int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), wdepth = CV_MAT_DEPTH(wtype);
int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), wdepth = std::max(CV_32S, CV_MAT_DEPTH(wtype));
wtype = CV_MAKETYPE(wdepth, cn);
int type2 = haveScalar ? _src2.type() : wtype, depth2 = CV_MAT_DEPTH(type2);
int type2 = haveScalar ? wtype : _src2.type(), depth2 = CV_MAT_DEPTH(type2);
int kercn = haveMask || haveScalar ? cn : 1;
UMat src1 = _src1.getUMat(), src2;
UMat dst = _dst.getUMat(), mask = _mask.getUMat();
char opts[1024];
int kercn = haveMask || haveScalar ? cn : 1;
if( (depth1 == depth2 || haveScalar) && ddepth == depth1 && wdepth == depth1 )
{
const char* oclopstr = oclop2str[oclop];
if( wdepth <= CV_16S )
{
oclopstr = oclop == OCL_OP_ADD ? "OCL_OP_ADD_SAT" :
oclop == OCL_OP_SUB ? "OCL_OP_SUB_SAT" :
oclop == OCL_OP_RSUB ? "OCL_OP_RSUB_SAT" : oclopstr;
}
sprintf(opts, "-D %s%s -D %s -D dstT=%s",
(haveMask ? "MASK_" : ""), (haveScalar ? "UNARY_OP" : "BINARY_OP"),
oclop2str[oclop], ocl::typeToStr(CV_MAKETYPE(ddepth, kercn)));
}
else
{
char cvtstr[3][32];
sprintf(opts, "-D %s%s -D %s -D srcT1=%s -D srcT2=%s "
"-D dstT=%s -D workT=%s -D convertToWT1=%s "
"-D convertToWT2=%s -D convertToDT=%s",
(haveMask ? "MASK_" : ""), (haveScalar ? "UNARY_OP" : "BINARY_OP"),
oclop2str[oclop], ocl::typeToStr(CV_MAKETYPE(depth1, kercn)),
ocl::typeToStr(CV_MAKETYPE(depth2, kercn)),
ocl::typeToStr(CV_MAKETYPE(ddepth, kercn)),
ocl::typeToStr(CV_MAKETYPE(wdepth, kercn)),
ocl::convertTypeStr(depth1, wdepth, kercn, cvtstr[0]),
ocl::convertTypeStr(depth2, wdepth, kercn, cvtstr[1]),
ocl::convertTypeStr(wdepth, ddepth, kercn, cvtstr[2]));
}
char cvtstr[3][32], opts[1024];
sprintf(opts, "-D %s%s -D %s -D srcT1=%s -D srcT2=%s "
"-D dstT=%s -D workT=%s -D convertToWT1=%s "
"-D convertToWT2=%s -D convertToDT=%s",
(haveMask ? "MASK_" : ""), (haveScalar ? "UNARY_OP" : "BINARY_OP"),
oclop2str[oclop], ocl::typeToStr(CV_MAKETYPE(depth1, kercn)),
ocl::typeToStr(CV_MAKETYPE(depth2, kercn)),
ocl::typeToStr(CV_MAKETYPE(ddepth, kercn)),
ocl::typeToStr(CV_MAKETYPE(wdepth, kercn)),
ocl::convertTypeStr(depth1, wdepth, kercn, cvtstr[0]),
ocl::convertTypeStr(depth2, wdepth, kercn, cvtstr[1]),
ocl::convertTypeStr(wdepth, ddepth, kercn, cvtstr[2]));
const uchar* usrdata_p = (const uchar*)usrdata;
const double* usrdata_d = (const double*)usrdata;
@ -1336,7 +1318,6 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
usrdata_f[i] = (float)usrdata_d[i];
usrdata_p = (const uchar*)usrdata_f;
}
size_t usrdata_esz = CV_ELEM_SIZE(wdepth);
ocl::Kernel k("KF", ocl::core::arithm_oclsrc, opts);
if( k.empty() )
@ -1368,6 +1349,7 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
}
else
{
size_t usrdata_esz = CV_ELEM_SIZE(wdepth);
src2 = _src2.getUMat();
ocl::KernelArg src2arg = ocl::KernelArg::ReadOnlyNoSize(src2, cscale);
@ -1392,8 +1374,8 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
}
}
size_t globalsize[] = { src1.cols*(cn/kercn), src1.rows };
return k.run(2, globalsize, 0, false);
size_t globalsize[] = { src1.cols*cscale, src1.rows };
return k.run(2, globalsize, NULL, false);
}
@ -1410,8 +1392,7 @@ static void arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
int wtype, dims1 = psrc1->dims(), dims2 = psrc2->dims();
Size sz1 = dims1 <= 2 ? psrc1->size() : Size();
Size sz2 = dims2 <= 2 ? psrc2->size() : Size();
bool use_opencl = (kind1 == _InputArray::UMAT || kind2 == _InputArray::UMAT || _dst.kind() == _OutputArray::UMAT) &&
ocl::useOpenCL() && dims1 <= 2 && dims2 <= 2;
bool use_opencl = _dst.kind() == _OutputArray::UMAT && ocl::useOpenCL() && dims1 <= 2 && dims2 <= 2;
bool src1Scalar = checkScalar(*psrc1, type2, kind1, kind2);
bool src2Scalar = checkScalar(*psrc2, type1, kind2, kind1);
@ -1426,6 +1407,7 @@ static void arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst,
(!usrdata ? type1 : std::max(depth1, CV_32F)),
usrdata, oclop, false))
return;
Mat src1 = psrc1->getMat(), src2 = psrc2->getMat(), dst = _dst.getMat();
Size sz = getContinuousSize(src1, src2, dst, src1.channels());
tab[depth1](src1.data, src1.step, src2.data, src2.step, dst.data, dst.step, sz, usrdata);

1
modules/core/src/ocl.cpp
View File

@ -2347,7 +2347,6 @@ struct Program::Impl
void** deviceList = deviceListBuf;
for( i = 0; i < n; i++ )
deviceList[i] = ctx.device(i).ptr();
printf("Building the OpenCL program ...\n");
retval = clBuildProgram(handle, n,
(const cl_device_id*)deviceList,
buildflags.c_str(), 0, 0);

12
modules/core/src/opencl/arithm.cl
View File

@ -89,21 +89,12 @@
#define EXTRA_PARAMS
#if defined OP_ADD_SAT
#define PROCESS_ELEM dstelem = add_sat(srcelem1, srcelem2)
#elif defined OP_ADD
#if defined OP_ADD
#define PROCESS_ELEM dstelem = convertToDT(srcelem1 + srcelem2)
#elif defined OP_SUB_SAT
#define PROCESS_ELEM dstelem = sub_sat(srcelem1, srcelem2)
#elif defined OP_SUB
#define PROCESS_ELEM dstelem = convertToDT(srcelem1 - srcelem2)
#elif defined OP_RSUB_SAT
#define PROCESS_ELEM dstelem = sub_sat(srcelem2, srcelem1)
#elif defined OP_RSUB
#define PROCESS_ELEM dstelem = convertToDT(srcelem2 - srcelem1)
@ -226,7 +217,6 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1,
int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset);
PROCESS_ELEM;
//printf("(x=%d, y=%d). %d, %d, %d\n", x, y, (int)srcelem1, (int)srcelem2, (int)dstelem);
}
}

56
modules/core/test/ocl/test_arithm.cpp
View File

@ -149,7 +149,7 @@ OCL_TEST_P(Add, Scalar)
generateTestData();
OCL_OFF(cv::add(src1_roi, val, dst1_roi));
OCL_ON(cv::add(usrc1_roi, val, udst1_roi));
OCL_ON(cv::add(val, usrc1_roi, udst1_roi));
Near(1e-5);
}
}
@ -166,12 +166,62 @@ OCL_TEST_P(Add, Scalar_Mask)
}
}
//////////////////////////////////////// Subtract //////////////////////////////////////////////
typedef ArithmTestBase Subtract;
OCL_TEST_P(Subtract, Mat)
{
for (int j = 0; j < test_loop_times; j++)
{
generateTestData();
OCL_OFF(cv::subtract(src1_roi, src2_roi, dst1_roi));
OCL_ON(cv::subtract(usrc1_roi, usrc2_roi, udst1_roi));
Near(0);
}
}
OCL_TEST_P(Subtract, Mat_Mask)
{
for (int j = 0; j < test_loop_times; j++)
{
generateTestData();
OCL_OFF(cv::subtract(src1_roi, src2_roi, dst1_roi, mask_roi));
OCL_ON(cv::subtract(usrc1_roi, usrc2_roi, udst1_roi, umask_roi));
Near(0);
}
}
OCL_TEST_P(Subtract, Scalar)
{
for (int j = 0; j < test_loop_times; j++)
{
generateTestData();
OCL_OFF(cv::subtract(val, src1_roi, dst1_roi));
OCL_ON(cv::subtract(val, usrc1_roi, udst1_roi));
Near(1e-5);
}
}
OCL_TEST_P(Subtract, Scalar_Mask)
{
for (int j = 0; j < test_loop_times; j++)
{
generateTestData();
OCL_OFF(cv::subtract(src1_roi, val, dst1_roi, mask_roi));
OCL_ON(cv::subtract(usrc1_roi, val, udst1_roi, umask_roi));
Near(1e-5);
}
}
//////////////////////////////////////// Instantiation /////////////////////////////////////////
// TODO FIXIT Invalid "add" implementation
//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(OCL_ALL_DEPTHS, ::testing::Values(1, 2, 4), Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Arithm, Subtract, Combine(OCL_ALL_DEPTHS, ::testing::Values(1, 2, 4), Bool()));
} } // namespace cvtest::ocl

43
modules/ts/include/opencv2/ts/ocl_test.hpp
View File

@ -60,6 +60,31 @@ namespace ocl {
using namespace cv;
using namespace testing;
namespace traits {
template <typename T>
struct GetMatForRead
{
};
template <>
struct GetMatForRead<Mat>
{
static const Mat get(const Mat& m) { return m; }
};
template <>
struct GetMatForRead<UMat>
{
static const Mat get(const UMat& m) { return m.getMat(ACCESS_READ); }
};
} // namespace traits
template <typename T>
const Mat getMatForRead(const T& mat)
{
return traits::GetMatForRead<T>::get(mat);
}
extern int test_loop_times;
#define MAX_VALUE 357
@ -203,32 +228,32 @@ struct TestUtils
template <typename T1>
static double checkNorm(const T1& m)
{
return checkNorm(cv::getMatForRead(m));
return checkNorm(getMatForRead(m));
}
template <typename T1, typename T2>
static double checkNorm(const T1& m1, const T2& m2)
{
return checkNorm(cv::getMatForRead(m1), cv::getMatForRead(m2));
return checkNorm(getMatForRead(m1), getMatForRead(m2));
}
template <typename T1, typename T2>
static double checkSimilarity(const T1& m1, const T2& m2)
{
return checkSimilarity(cv::getMatForRead(m1), cv::getMatForRead(m2));
return checkSimilarity(getMatForRead(m1), getMatForRead(m2));
}
template <typename T1, typename T2>
static inline double checkNormRelative(const T1& m1, const T2& m2)
{
const Mat _m1 = cv::getMatForRead(m1);
const Mat _m2 = cv::getMatForRead(m2);
const Mat _m1 = getMatForRead(m1);
const Mat _m2 = getMatForRead(m2);
return checkNormRelative(_m1, _m2);
}
template <typename T1, typename T2, typename T3>
static void showDiff(const T1& src, const T2& gold, const T3& actual, double eps, bool alwaysShow = false)
{
const Mat _src = cv::getMatForRead(src);
const Mat _gold = cv::getMatForRead(gold);
const Mat _actual = cv::getMatForRead(actual);
const Mat _src = getMatForRead(src);
const Mat _gold = getMatForRead(gold);
const Mat _actual = getMatForRead(actual);
showDiff(_src, _gold, _actual, eps, alwaysShow);
}
};
@ -277,7 +302,7 @@ IMPLEMENT_PARAM_CLASS(Channels, int)
#define OCL_OFF(fn) cv::ocl::setUseOpenCL(false); fn
#define OCL_ON(fn) cv::ocl::setUseOpenCL(true); fn
#define OCL_ALL_DEPTHS Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F)
#define OCL_ALL_DEPTHS Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F)
#define OCL_ALL_CHANNELS Values(1, 2, 3, 4)
#define OCL_INSTANTIATE_TEST_CASE_P(prefix, test_case_name, generator) \