diff --git a/modules/core/src/arithm.cpp b/modules/core/src/arithm.cpp index c3ba8c6d64..bcd11d2566 100644 --- a/modules/core/src/arithm.cpp +++ b/modules/core/src/arithm.cpp @@ -929,11 +929,11 @@ static bool ocl_binary_op(InputArray _src1, InputArray _src2, OutputArray _dst, int srcdepth = CV_MAT_DEPTH(srctype); int cn = CV_MAT_CN(srctype); - if( oclop < 0 || ((haveMask || haveScalar) && cn > 4) ) - return false; + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; - UMat src1 = _src1.getUMat(), src2; - UMat dst = _dst.getUMat(), mask = _mask.getUMat(); + if( oclop < 0 || ((haveMask || haveScalar) && (cn > 4 || cn == 3)) || + (!doubleSupport && srcdepth == CV_64F)) + return false; char opts[1024]; int kercn = haveMask || haveScalar ? cn : 1; @@ -946,6 +946,9 @@ static bool ocl_binary_op(InputArray _src1, InputArray _src2, OutputArray _dst, if( k.empty() ) return false; + UMat src1 = _src1.getUMat(), src2; + UMat dst = _dst.getUMat(), mask = _mask.getUMat(); + int cscale = cn/kercn; ocl::KernelArg src1arg = ocl::KernelArg::ReadOnlyNoSize(src1, cscale); ocl::KernelArg dstarg = haveMask ? ocl::KernelArg::ReadWrite(dst, cscale) : @@ -1280,24 +1283,28 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst, void* usrdata, int oclop, bool haveScalar ) { + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; int type1 = _src1.type(), depth1 = CV_MAT_DEPTH(type1), cn = CV_MAT_CN(type1); bool haveMask = !_mask.empty(); - if( ((haveMask || haveScalar) && cn > 4) || cn == 3) // TODO need fix for 3 channels + if( ((haveMask || haveScalar) && (cn > 4 || cn == 3)) ) return false; int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype), wdepth = std::max(CV_32S, CV_MAT_DEPTH(wtype)); + if (!doubleSupport) + wdepth = std::min(wdepth, CV_32F); + wtype = CV_MAKETYPE(wdepth, cn); int type2 = haveScalar ? wtype : _src2.type(), depth2 = CV_MAT_DEPTH(type2); - int kercn = haveMask || haveScalar ? cn : 1; + if (!doubleSupport && (depth2 == CV_64F || depth1 == CV_64F)) + return false; - UMat src1 = _src1.getUMat(), src2; - UMat dst = _dst.getUMat(), mask = _mask.getUMat(); + int kercn = haveMask || haveScalar ? cn : 1; 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", + "-D convertToWT2=%s -D convertToDT=%s%s", (haveMask ? "MASK_" : ""), (haveScalar ? "UNARY_OP" : "BINARY_OP"), oclop2str[oclop], ocl::typeToStr(CV_MAKETYPE(depth1, kercn)), ocl::typeToStr(CV_MAKETYPE(depth2, kercn)), @@ -1305,7 +1312,8 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst, 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])); + ocl::convertTypeStr(wdepth, ddepth, kercn, cvtstr[2]), + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); const uchar* usrdata_p = (const uchar*)usrdata; const double* usrdata_d = (const double*)usrdata; @@ -1323,6 +1331,9 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst, if( k.empty() ) return false; + UMat src1 = _src1.getUMat(), src2; + UMat dst = _dst.getUMat(), mask = _mask.getUMat(); + int cscale = cn/kercn; ocl::KernelArg src1arg = ocl::KernelArg::ReadOnlyNoSize(src1, cscale); @@ -1337,9 +1348,7 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst, Mat src2sc = _src2.getMat(); if( !src2sc.empty() ) - { convertAndUnrollScalar(src2sc, wtype, (uchar*)buf, 1); - } ocl::KernelArg scalararg = ocl::KernelArg(0, 0, 0, buf, esz); if( !haveMask ) @@ -1369,12 +1378,10 @@ static bool ocl_arithm_op(InputArray _src1, InputArray _src2, OutputArray _dst, CV_Error(Error::StsNotImplemented, "unsupported number of extra parameters"); } else - { k.args(src1arg, src2arg, maskarg, dstarg); - } } - size_t globalsize[] = { src1.cols*cscale, src1.rows }; + size_t globalsize[] = { src1.cols * cscale, src1.rows }; return k.run(2, globalsize, NULL, false); } @@ -2075,7 +2082,7 @@ void cv::multiply(InputArray src1, InputArray src2, OutputArray dst, double scale, int dtype) { arithm_op(src1, src2, dst, noArray(), dtype, getMulTab(), - true, &scale, scale == 1. ? OCL_OP_MUL : OCL_OP_MUL_SCALE); + true, &scale, std::abs(scale - 1.0) < DBL_EPSILON ? OCL_OP_MUL : OCL_OP_MUL_SCALE); } void cv::divide(InputArray src1, InputArray src2, @@ -2581,6 +2588,42 @@ static double getMaxVal(int depth) return tab[depth]; } +static bool ocl_compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op) +{ + if ( !((_src1.isMat() || _src1.isUMat()) && (_src2.isMat() || _src2.isUMat())) ) + return false; + + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + int type = _src1.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), type2 = _src2.type(); + if (!doubleSupport && (depth == CV_64F || _src2.depth() == CV_64F)) + return false; + + const char * const operationMap[] = { "==", ">", ">=", "<", "<=", "!=" }; + ocl::Kernel k("KF", ocl::core::arithm_oclsrc, + format("-D BINARY_OP -D srcT1=%s -D workT=srcT1" + " -D OP_CMP -D CMP_OPERATOR=%s%s", + ocl::typeToStr(CV_MAKE_TYPE(depth, 1)), + operationMap[op], + doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (k.empty()) + return false; + + CV_Assert(type == type2); + UMat src1 = _src1.getUMat(), src2 = _src2.getUMat(); + Size size = src1.size(); + CV_Assert(size == src2.size()); + + _dst.create(size, CV_8UC(cn)); + UMat dst = _dst.getUMat(); + + k.args(ocl::KernelArg::ReadOnlyNoSize(src1), + ocl::KernelArg::ReadOnlyNoSize(src2), + ocl::KernelArg::WriteOnly(dst, cn)); + + size_t globalsize[2] = { dst.cols * cn, dst.rows }; + return k.run(2, globalsize, NULL, false); +} + } void cv::compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op) @@ -2588,6 +2631,10 @@ void cv::compare(InputArray _src1, InputArray _src2, OutputArray _dst, int op) CV_Assert( op == CMP_LT || op == CMP_LE || op == CMP_EQ || op == CMP_NE || op == CMP_GE || op == CMP_GT ); + if (ocl::useOpenCL() && _src1.dims() <= 2 && _src2.dims() <= 2 && _dst.isUMat() && + ocl_compare(_src1, _src2, _dst, op)) + return; + int kind1 = _src1.kind(), kind2 = _src2.kind(); Mat src1 = _src1.getMat(), src2 = _src2.getMat(); diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp index 29601bec23..79959435d3 100644 --- a/modules/core/src/mathfuncs.cpp +++ b/modules/core/src/mathfuncs.cpp @@ -497,10 +497,49 @@ void phase( InputArray src1, InputArray src2, OutputArray dst, bool angleInDegre } } +static bool ocl_cartToPolar( InputArray _src1, InputArray _src2, + OutputArray _dst1, OutputArray _dst2, bool angleInDegrees ) +{ + int type = _src1.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if ( !(_src1.dims() <= 2 && _src2.dims() <= 2 && + (depth == CV_32F || depth == CV_64F) && type == _src2.type()) || + (depth == CV_64F && !doubleSupport) ) + return false; + + ocl::Kernel k("KF", ocl::core::arithm_oclsrc, + format("-D BINARY_OP -D dstT=%s -D OP_CTP_%s%s", + ocl::typeToStr(CV_MAKE_TYPE(depth, 1)), + angleInDegrees ? "AD" : "AR", + doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (k.empty()) + return false; + + UMat src1 = _src1.getUMat(), src2 = _src2.getUMat(); + Size size = src1.size(); + CV_Assert( size == src2.size() ); + + _dst1.create(size, type); + _dst2.create(size, type); + UMat dst1 = _dst1.getUMat(), dst2 = _dst2.getUMat(); + + k.args(ocl::KernelArg::ReadOnlyNoSize(src1), + ocl::KernelArg::ReadOnlyNoSize(src2), + ocl::KernelArg::WriteOnly(dst1, cn), + ocl::KernelArg::WriteOnlyNoSize(dst2)); + + size_t globalsize[2] = { dst1.cols * cn, dst1.rows }; + return k.run(2, globalsize, NULL, false); +} void cartToPolar( InputArray src1, InputArray src2, OutputArray dst1, OutputArray dst2, bool angleInDegrees ) { + if (ocl::useOpenCL() && dst1.isUMat() && dst2.isUMat() && + ocl_cartToPolar(src1, src2, dst1, dst2, angleInDegrees)) + return; + Mat X = src1.getMat(), Y = src2.getMat(); int type = X.type(), depth = X.depth(), cn = X.channels(); CV_Assert( X.size == Y.size && type == Y.type() && (depth == CV_32F || depth == CV_64F)); @@ -644,12 +683,50 @@ static void SinCos_32f( const float *angle, float *sinval, float* cosval, } +static bool ocl_polarToCart( InputArray _mag, InputArray _angle, + OutputArray _dst1, OutputArray _dst2, bool angleInDegrees ) +{ + int type = _angle.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if ( !doubleSupport && depth == CV_64F ) + return false; + + ocl::Kernel k("KF", ocl::core::arithm_oclsrc, + format("-D dstT=%s -D BINARY_OP -D OP_PTC_%s%s", + ocl::typeToStr(CV_MAKE_TYPE(depth, 1)), + angleInDegrees ? "AD" : "AR", + doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (k.empty()) + return false; + + UMat mag = _mag.getUMat(), angle = _angle.getUMat(); + Size size = angle.size(); + CV_Assert(mag.size() == size); + + _dst1.create(size, type); + _dst2.create(size, type); + UMat dst1 = _dst1.getUMat(), dst2 = _dst2.getUMat(); + + k.args(ocl::KernelArg::ReadOnlyNoSize(mag), ocl::KernelArg::ReadOnlyNoSize(angle), + ocl::KernelArg::WriteOnly(dst1, cn), ocl::KernelArg::WriteOnlyNoSize(dst2)); + + size_t globalsize[2] = { dst1.cols * cn, dst1.rows }; + return k.run(2, globalsize, NULL, false); +} + void polarToCart( InputArray src1, InputArray src2, OutputArray dst1, OutputArray dst2, bool angleInDegrees ) { + int type = src2.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + CV_Assert((depth == CV_32F || depth == CV_64F) && (src1.empty() || src1.type() == type)); + + if (ocl::useOpenCL() && !src1.empty() && src2.dims() <= 2 && dst1.isUMat() && dst2.isUMat() && + ocl_polarToCart(src1, src2, dst1, dst2, angleInDegrees)) + return; + Mat Mag = src1.getMat(), Angle = src2.getMat(); - int type = Angle.type(), depth = Angle.depth(), cn = Angle.channels(); - CV_Assert( Mag.empty() || (Angle.size == Mag.size && type == Mag.type() && (depth == CV_32F || depth == CV_64F))); + CV_Assert( Mag.empty() || Angle.size == Mag.size); dst1.create( Angle.dims, Angle.size, type ); dst2.create( Angle.dims, Angle.size, type ); Mat X = dst1.getMat(), Y = dst2.getMat(); @@ -1955,9 +2032,42 @@ static IPowFunc ipowTab[] = (IPowFunc)iPow32s, (IPowFunc)iPow32f, (IPowFunc)iPow64f, 0 }; +static bool ocl_pow(InputArray _src, double power, OutputArray _dst) +{ + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if ( !(_src.dims() <= 2 && (depth == CV_32F || depth == CV_64F)) || + (depth == CV_64F && !doubleSupport) ) + return false; + + ocl::Kernel k("KF", ocl::core::arithm_oclsrc, + format("-D dstT=%s -D OP_POW -D UNARY_OP%s", ocl::typeToStr(CV_MAKE_TYPE(depth, 1)), + doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (k.empty()) + return false; + + UMat src = _src.getUMat(); + _dst.create(src.size(), type); + UMat dst = _dst.getUMat(); + + ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src), + dstarg = ocl::KernelArg::WriteOnly(dst, cn); + + if (depth == CV_32F) + k.args(srcarg, dstarg, (float)power); + else + k.args(srcarg, dstarg, power); + + size_t globalsize[2] = { dst.cols * cn, dst.rows }; + return k.run(2, globalsize, NULL, false); +} void pow( InputArray _src, double power, OutputArray _dst ) { + if (ocl::useOpenCL() && _dst.isUMat() && ocl_pow(_src, power, _dst)) + return; + Mat src = _src.getMat(); int type = src.type(), depth = src.depth(), cn = src.channels(); diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index 919452dc98..4e9be9807c 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" /****************************************************************************************\ * [scaled] Identity matrix initialization * @@ -2368,10 +2369,37 @@ void cv::vconcat(InputArray _src, OutputArray dst) } //////////////////////////////////////// set identity //////////////////////////////////////////// + +namespace cv { + +static bool ocl_setIdentity( InputOutputArray _m, const Scalar& s ) +{ + int type = _m.type(), cn = CV_MAT_CN(type); + if (cn == 3) + return false; + + ocl::Kernel k("setIdentity", ocl::core::set_identity_oclsrc, + format("-D T=%s", ocl::memopTypeToStr(type))); + if (k.empty()) + return false; + + UMat m = _m.getUMat(); + k.args(ocl::KernelArg::WriteOnly(m), ocl::KernelArg::Constant(Mat(1, 1, type, s))); + + size_t globalsize[2] = { m.cols, m.rows }; + return k.run(2, globalsize, NULL, false); +} + +} + void cv::setIdentity( InputOutputArray _m, const Scalar& s ) { + CV_Assert( _m.dims() <= 2 ); + + if (ocl::useOpenCL() && _m.isUMat() && ocl_setIdentity(_m, s)) + return; + Mat m = _m.getMat(); - CV_Assert( m.dims <= 2 ); int i, j, rows = m.rows, cols = m.cols, type = m.type(); if( type == CV_32FC1 ) @@ -2548,18 +2576,63 @@ static TransposeInplaceFunc transposeInplaceTab[] = 0, 0, 0, 0, 0, 0, 0, transposeI_32sC6, 0, 0, 0, 0, 0, 0, 0, transposeI_32sC8 }; +static inline int divUp(int a, int b) +{ + return (a + b - 1) / b; +} + +static bool ocl_transpose( InputArray _src, OutputArray _dst ) +{ + const int TILE_DIM = 32, BLOCK_ROWS = 8; + int type = _src.type(), cn = CV_MAT_CN(type); + + if (cn == 3) + return false; + + UMat src = _src.getUMat(); + _dst.create(src.cols, src.rows, type); + UMat dst = _dst.getUMat(); + + String kernelName("transpose"); + bool inplace = dst.u == src.u; + + if (inplace) + { + CV_Assert(dst.cols == dst.rows); + kernelName += "_inplace"; + } + + ocl::Kernel k(kernelName.c_str(), ocl::core::transpose_oclsrc, + format("-D T=%s -D TILE_DIM=%d -D BLOCK_ROWS=%d", + ocl::memopTypeToStr(type), TILE_DIM, BLOCK_ROWS)); + if (inplace) + k.args(ocl::KernelArg::ReadWriteNoSize(dst), dst.rows); + else + k.args(ocl::KernelArg::ReadOnly(src), + ocl::KernelArg::WriteOnlyNoSize(dst)); + + size_t localsize[3] = { TILE_DIM, BLOCK_ROWS, 1 }; + size_t globalsize[3] = { src.cols, inplace ? src.rows : divUp(src.rows, TILE_DIM) * BLOCK_ROWS, 1 }; + + return k.run(2, globalsize, localsize, false); +} + } void cv::transpose( InputArray _src, OutputArray _dst ) { + int type = _src.type(), esz = CV_ELEM_SIZE(type); + CV_Assert( _src.dims() <= 2 && esz <= 32 ); + + if (ocl::useOpenCL() && _dst.isUMat() && ocl_transpose(_src, _dst)) + return; + Mat src = _src.getMat(); if( src.empty() ) { _dst.release(); return; } - size_t esz = src.elemSize(); - CV_Assert( src.dims <= 2 && esz <= (size_t)32 ); _dst.create(src.cols, src.rows, src.type()); Mat dst = _dst.getMat(); @@ -2576,6 +2649,7 @@ void cv::transpose( InputArray _src, OutputArray _dst ) { TransposeInplaceFunc func = transposeInplaceTab[esz]; CV_Assert( func != 0 ); + CV_Assert( dst.cols == dst.rows ); func( dst.data, dst.step, dst.rows ); } else diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 64460efb0d..f733dd11fb 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -3145,7 +3145,7 @@ const char* memopTypeToStr(int t) "ushort", "ushort2", "ushort3", "ushort4", "int", "int2", "int3", "int4", "int", "int2", "int3", "int4", - "long", "long2", "long3", "long4", + "int2", "int4", "?", "int8", "?", "?", "?", "?" }; int cn = CV_MAT_CN(t); diff --git a/modules/core/src/opencl/arithm.cl b/modules/core/src/opencl/arithm.cl index a9c23645ac..b4cdb53f2c 100644 --- a/modules/core/src/opencl/arithm.cl +++ b/modules/core/src/opencl/arithm.cl @@ -57,19 +57,22 @@ -D workDepth= [-D cn=]" - for mixed-type operations */ -#if defined (DOUBLE_SUPPORT) +#ifdef DOUBLE_SUPPORT #ifdef cl_khr_fp64 #pragma OPENCL EXTENSION cl_khr_fp64:enable #elif defined (cl_amd_fp64) #pragma OPENCL EXTENSION cl_amd_fp64:enable #endif +#define CV_EPSILON DBL_EPSILON +#define CV_PI M_PI +#else +#define CV_EPSILON FLT_EPSILON +#define CV_PI M_PI_F #endif -#define CV_32S 4 -#define CV_32F 5 - #define dstelem *(__global dstT*)(dstptr + dst_index) -#define noconvert(x) x +#define dstelem2 *(__global dstT*)(dstptr2 + dst_index2) +#define noconvert #ifndef workT @@ -88,6 +91,7 @@ #endif #define EXTRA_PARAMS +#define EXTRA_INDEX #if defined OP_ADD #define PROCESS_ELEM dstelem = convertToDT(srcelem1 + srcelem2) @@ -99,7 +103,9 @@ #define PROCESS_ELEM dstelem = convertToDT(srcelem2 - srcelem1) #elif defined OP_ABSDIFF -#define PROCESS_ELEM dstelem = abs_diff(srcelem1, srcelem2) +#define PROCESS_ELEM \ + workT v = srcelem1 - srcelem2; \ + dstelem = convertToDT(v >= (workT)(0) ? v : -v); #elif defined OP_AND #define PROCESS_ELEM dstelem = srcelem1 & srcelem2 @@ -169,6 +175,9 @@ #elif defined OP_EXP #define PROCESS_ELEM dstelem = exp(srcelem1) +#elif defined OP_POW +#define PROCESS_ELEM dstelem = pow(srcelem1, srcelem2) + #elif defined OP_SQRT #define PROCESS_ELEM dstelem = sqrt(srcelem1) @@ -178,6 +187,10 @@ dstT v = (dstT)(srcelem1);\ dstelem = v > (dstT)(0) ? log(v) : log(-v) #elif defined OP_CMP +#define dstT uchar +#define srcT2 srcT1 +#define convertToWT1 +#define convertToWT2 #define PROCESS_ELEM dstelem = convert_uchar(srcelem1 CMP_OPERATOR srcelem2 ? 255 : 0) #elif defined OP_CONVERT @@ -188,15 +201,55 @@ dstelem = v > (dstT)(0) ? log(v) : log(-v) #define EXTRA_PARAMS , workT alpha, workT beta #define PROCESS_ELEM dstelem = convertToDT(srcelem1*alpha + beta) +#elif defined OP_CTP_AD || defined OP_CTP_AR +#ifdef OP_CTP_AD +#define TO_DEGREE cartToPolar *= (180 / CV_PI); +#elif defined OP_CTP_AR +#define TO_DEGREE +#endif +#define PROCESS_ELEM \ + dstT x = srcelem1, y = srcelem2; \ + dstT x2 = x * x, y2 = y * y; \ + dstT magnitude = sqrt(x2 + y2); \ + dstT tmp = y >= 0 ? 0 : CV_PI * 2; \ + tmp = x < 0 ? CV_PI : tmp; \ + dstT tmp1 = y >= 0 ? CV_PI * 0.5f : CV_PI * 1.5f; \ + dstT cartToPolar = y2 <= x2 ? x * y / (x2 + 0.28f * y2 + CV_EPSILON) + tmp : (tmp1 - x * y / (y2 + 0.28f * x2 + CV_EPSILON)); \ + TO_DEGREE \ + dstelem = magnitude; \ + dstelem2 = cartToPolar + +#elif defined OP_PTC_AD || defined OP_PTC_AR +#ifdef OP_PTC_AD +#define FROM_DEGREE \ + dstT ascale = CV_PI/180.0f; \ + dstT alpha = y * ascale +#else +#define FROM_DEGREE \ + dstT alpha = y +#endif +#define PROCESS_ELEM \ + dstT x = srcelem1, y = srcelem2; \ + FROM_DEGREE; \ + dstelem = cos(alpha) * x; \ + dstelem2 = sin(alpha) * x + #else #error "unknown op type" #endif +#if defined OP_CTP_AD || defined OP_CTP_AR || defined OP_PTC_AD || defined OP_PTC_AR + #undef EXTRA_PARAMS + #define EXTRA_PARAMS , __global uchar* dstptr2, int dststep2, int dstoffset2 + #undef EXTRA_INDEX + #define EXTRA_INDEX int dst_index2 = mad24(y, dststep2, x*(int)sizeof(dstT) + dstoffset2) +#endif + #if defined UNARY_OP || defined MASK_UNARY_OP #undef srcelem2 #if defined OP_AND || defined OP_OR || defined OP_XOR || defined OP_ADD || defined OP_SAT_ADD || \ defined OP_SUB || defined OP_SAT_SUB || defined OP_RSUB || defined OP_SAT_RSUB || \ - defined OP_ABSDIFF || defined OP_CMP || defined OP_MIN || defined OP_MAX + defined OP_ABSDIFF || defined OP_CMP || defined OP_MIN || defined OP_MAX || defined OP_POW #undef EXTRA_PARAMS #define EXTRA_PARAMS , workT srcelem2 #endif @@ -217,6 +270,7 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1, int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1); int src2_index = mad24(y, srcstep2, x*(int)sizeof(srcT2) + srcoffset2); int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); + EXTRA_INDEX; PROCESS_ELEM; } @@ -260,6 +314,7 @@ __kernel void KF(__global const uchar* srcptr1, int srcstep1, int srcoffset1, { int src1_index = mad24(y, srcstep1, x*(int)sizeof(srcT1) + srcoffset1); int dst_index = mad24(y, dststep, x*(int)sizeof(dstT) + dstoffset); + EXTRA_INDEX; PROCESS_ELEM; } diff --git a/modules/core/src/opencl/reduce.cl b/modules/core/src/opencl/reduce.cl new file mode 100644 index 0000000000..526cc51010 --- /dev/null +++ b/modules/core/src/opencl/reduce.cl @@ -0,0 +1,130 @@ +//////////////////////////////////////////////////////////////////////////////////////// +// +// 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. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Shengen Yan,yanshengen@gmail.com +// +// 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 Intel Corporation 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. +// + +#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 + +#define noconvert + +#if defined OP_SUM || defined OP_SUM_ABS || defined OP_SUM_SQR +#if OP_SUM +#define FUNC(a, b) a += b +#elif OP_SUM_ABS +#define FUNC(a, b) a += b >= (dstT)(0) ? b : -b +#elif OP_SUM_SQR +#define FUNC(a, b) a += b * b +#endif +#define DEFINE_ACCUMULATOR \ + dstT accumulator = (dstT)(0) +#define REDUCE_GLOBAL \ + dstT temp = convertToDT(src[0]); \ + FUNC(accumulator, temp) +#define REDUCE_LOCAL_1 \ + localmem[lid - WGS2_ALIGNED] += accumulator +#define REDUCE_LOCAL_2 \ + localmem[lid] += localmem[lid2] + +#elif defined OP_COUNT_NON_ZERO +#define dstT int +#define DEFINE_ACCUMULATOR \ + dstT accumulator = (dstT)(0); \ + srcT zero = (srcT)(0), one = (srcT)(1) +#define REDUCE_GLOBAL \ + accumulator += src[0] == zero ? zero : one +#define REDUCE_LOCAL_1 \ + localmem[lid - WGS2_ALIGNED] += accumulator +#define REDUCE_LOCAL_2 \ + localmem[lid] += localmem[lid2] + +#else +#error "No operation" + +#endif + +__kernel void reduce(__global const uchar * srcptr, int step, int offset, int cols, + int total, int groupnum, __global uchar * dstptr) +{ + int lid = get_local_id(0); + int gid = get_group_id(0); + int id = get_global_id(0); + + __local dstT localmem[WGS2_ALIGNED]; + DEFINE_ACCUMULATOR; + + for (int grain = groupnum * WGS; id < total; id += grain) + { + int src_index = mad24(id / cols, step, offset + (id % cols) * (int)sizeof(srcT)); + __global const srcT * src = (__global const srcT *)(srcptr + src_index); + REDUCE_GLOBAL; + } + + if (lid < WGS2_ALIGNED) + localmem[lid] = accumulator; + barrier(CLK_LOCAL_MEM_FENCE); + + if (lid >= WGS2_ALIGNED) + REDUCE_LOCAL_1; + barrier(CLK_LOCAL_MEM_FENCE); + + for (int lsize = WGS2_ALIGNED >> 1; lsize > 0; lsize >>= 1) + { + if (lid < lsize) + { + int lid2 = lsize + lid; + REDUCE_LOCAL_2; + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + if (lid == 0) + { + __global dstT * dst = (__global dstT *)(dstptr + (int)sizeof(dstT) * gid); + dst[0] = localmem[0]; + } +} diff --git a/modules/core/src/opencl/set_identity.cl b/modules/core/src/opencl/set_identity.cl new file mode 100644 index 0000000000..de8caaf85b --- /dev/null +++ b/modules/core/src/opencl/set_identity.cl @@ -0,0 +1,59 @@ +/*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, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Jin Ma jin@multicorewareinc.com +// +// 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 Intel Corporation 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*/ + +__kernel void setIdentity(__global uchar * srcptr, int src_step, int src_offset, int rows, int cols, + T scalar) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < cols && y < rows) + { + int src_index = mad24(y, src_step, src_offset + x * (int)sizeof(T)); + __global T * src = (__global T *)(srcptr + src_index); + + src[0] = x == y ? scalar : (T)(0); + } +} diff --git a/modules/core/src/opencl/transpose.cl b/modules/core/src/opencl/transpose.cl new file mode 100644 index 0000000000..da9608c0d2 --- /dev/null +++ b/modules/core/src/opencl/transpose.cl @@ -0,0 +1,124 @@ +/*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. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Jia Haipeng, jiahaipeng95@gmail.com +// +// 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 Intel Corporation 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*/ + +#define LDS_STEP TILE_DIM + +__kernel void transpose(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols, + __global uchar * dstptr, int dst_step, int dst_offset) +{ + int gp_x = get_group_id(0), gp_y = get_group_id(1); + int gs_x = get_num_groups(0), gs_y = get_num_groups(1); + + int groupId_x, groupId_y; + + if (src_rows == src_cols) + { + groupId_y = gp_x; + groupId_x = (gp_x + gp_y) % gs_x; + } + else + { + int bid = gp_x + gs_x * gp_y; + groupId_y = bid % gs_y; + groupId_x = ((bid / gs_y) + groupId_y) % gs_x; + } + + int lx = get_local_id(0); + int ly = get_local_id(1); + + int x = groupId_x * TILE_DIM + lx; + int y = groupId_y * TILE_DIM + ly; + + int x_index = groupId_y * TILE_DIM + lx; + int y_index = groupId_x * TILE_DIM + ly; + + __local T title[TILE_DIM * LDS_STEP]; + + if (x < src_cols && y < src_rows) + { + int index_src = mad24(y, src_step, x * (int)sizeof(T) + src_offset); + + for (int i = 0; i < TILE_DIM; i += BLOCK_ROWS) + if (y + i < src_rows) + { + __global const T * src = (__global const T *)(srcptr + index_src); + title[(ly + i) * LDS_STEP + lx] = src[0]; + index_src = mad24(BLOCK_ROWS, src_step, index_src); + } + } + barrier(CLK_LOCAL_MEM_FENCE); + + if (x_index < src_rows && y_index < src_cols) + { + int index_dst = mad24(y_index, dst_step, x_index * (int)sizeof(T) + dst_offset); + + for (int i = 0; i < TILE_DIM; i += BLOCK_ROWS) + if ((y_index + i) < src_cols) + { + __global T * dst = (__global T *)(dstptr + index_dst); + dst[0] = title[lx * LDS_STEP + ly + i]; + index_dst = mad24(BLOCK_ROWS, dst_step, index_dst); + } + } +} + +__kernel void transpose_inplace(__global uchar * srcptr, int src_step, int src_offset, int src_rows) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (y < src_rows && x < y) + { + int src_index = mad24(y, src_step, src_offset + x * (int)sizeof(T)); + int dst_index = mad24(x, src_step, src_offset + y * (int)sizeof(T)); + + __global T * src = (__global T *)(srcptr + src_index); + __global T * dst = (__global T *)(srcptr + dst_index); + + T tmp = dst[0]; + dst[0] = src[0]; + src[0] = tmp; + } +} diff --git a/modules/core/src/stat.cpp b/modules/core/src/stat.cpp index bb2e1f4932..b19be3b476 100644 --- a/modules/core/src/stat.cpp +++ b/modules/core/src/stat.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" #include #include @@ -448,10 +449,77 @@ static SumSqrFunc getSumSqrTab(int depth) return sumSqrTab[depth]; } +template Scalar ocl_part_sum(Mat m) +{ + CV_Assert(m.rows == 1); + + Scalar s = Scalar::all(0); + int cn = m.channels(); + const T * const ptr = m.ptr(0); + + for (int x = 0, w = m.cols * cn; x < w; ) + for (int c = 0; c < cn; ++c, ++x) + s[c] += ptr[x]; + + return s; +} + +enum { OCL_OP_SUM = 0, OCL_OP_SUM_ABS = 1, OCL_OP_SUM_SQR = 2 }; + +static bool ocl_sum( InputArray _src, Scalar & res, int sum_op ) +{ + CV_Assert(sum_op == OCL_OP_SUM || sum_op == OCL_OP_SUM_ABS || sum_op == OCL_OP_SUM_SQR); + + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if ( (!doubleSupport && depth == CV_64F) || cn > 4 || cn == 3 || _src.dims() > 2 ) + return false; + + int dbsize = ocl::Device::getDefault().maxComputeUnits(); + size_t wgs = ocl::Device::getDefault().maxWorkGroupSize(); + + int ddepth = std::max(CV_32S, depth), dtype = CV_MAKE_TYPE(ddepth, cn); + + int wgs2_aligned = 1; + while (wgs2_aligned < (int)wgs) + wgs2_aligned <<= 1; + wgs2_aligned >>= 1; + + static const char * const opMap[3] = { "OP_SUM", "OP_SUM_ABS", "OP_SUM_SQR" }; + char cvt[40]; + ocl::Kernel k("reduce", ocl::core::reduce_oclsrc, + format("-D srcT=%s -D dstT=%s -D convertToDT=%s -D %s -D WGS=%d -D WGS2_ALIGNED=%d%s", + ocl::typeToStr(type), ocl::typeToStr(dtype), ocl::convertTypeStr(depth, ddepth, cn, cvt), + opMap[sum_op], (int)wgs, wgs2_aligned, + doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (k.empty()) + return false; + + UMat src = _src.getUMat(), db(1, dbsize, dtype); + k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(), + dbsize, ocl::KernelArg::PtrWriteOnly(db)); + + size_t globalsize = dbsize * wgs; + if (k.run(1, &globalsize, &wgs, true)) + { + typedef Scalar (*part_sum)(Mat m); + part_sum funcs[3] = { ocl_part_sum, ocl_part_sum, ocl_part_sum }, + func = funcs[ddepth - CV_32S]; + res = func(db.getMat(ACCESS_READ)); + return true; + } + return false; +} + } cv::Scalar cv::sum( InputArray _src ) { + Scalar _res; + if (ocl::useOpenCL() && _src.isUMat() && ocl_sum(_src, _res, OCL_OP_SUM)) + return _res; + Mat src = _src.getMat(); int k, cn = src.channels(), depth = src.depth(); @@ -542,12 +610,55 @@ cv::Scalar cv::sum( InputArray _src ) return s; } +namespace cv { + +static bool ocl_countNonZero( InputArray _src, int & res ) +{ + int type = _src.type(), depth = CV_MAT_DEPTH(type); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if (depth == CV_64F && !doubleSupport) + return false; + + int dbsize = ocl::Device::getDefault().maxComputeUnits(); + size_t wgs = ocl::Device::getDefault().maxWorkGroupSize(); + + int wgs2_aligned = 1; + while (wgs2_aligned < (int)wgs) + wgs2_aligned <<= 1; + wgs2_aligned >>= 1; + + ocl::Kernel k("reduce", ocl::core::reduce_oclsrc, + format("-D srcT=%s -D OP_COUNT_NON_ZERO -D WGS=%d -D WGS2_ALIGNED=%d%s", + ocl::typeToStr(type), (int)wgs, + wgs2_aligned, doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + if (k.empty()) + return false; + + UMat src = _src.getUMat(), db(1, dbsize, CV_32SC1); + k.args(ocl::KernelArg::ReadOnlyNoSize(src), src.cols, (int)src.total(), + dbsize, ocl::KernelArg::PtrWriteOnly(db)); + + size_t globalsize = dbsize * wgs; + if (k.run(1, &globalsize, &wgs, true)) + return res = saturate_cast(cv::sum(db.getMat(ACCESS_READ))[0]), true; + return false; +} + +} + int cv::countNonZero( InputArray _src ) { + CV_Assert( _src.channels() == 1 ); + + int res = -1; + if (ocl::useOpenCL() && _src.isUMat() && ocl_countNonZero(_src, res)) + return res; + Mat src = _src.getMat(); CountNonZeroFunc func = getCountNonZeroTab(src.depth()); - CV_Assert( src.channels() == 1 && func != 0 ); + CV_Assert( func != 0 ); const Mat* arrays[] = {&src, 0}; uchar* ptrs[1]; @@ -693,9 +804,54 @@ cv::Scalar cv::mean( InputArray _src, InputArray _mask ) return s*(nz0 ? 1./nz0 : 0); } +namespace cv { + +static bool ocl_meanStdDev( InputArray _src, OutputArray _mean, OutputArray _sdv ) +{ + Scalar mean, stddev; + if (!ocl_sum(_src, mean, OCL_OP_SUM)) + return false; + if (!ocl_sum(_src, stddev, OCL_OP_SUM_SQR)) + return false; + + double total = 1.0 / _src.total(); + int k, j, cn = _src.channels(); + for (int i = 0; i < cn; ++i) + { + mean[i] *= total; + stddev[i] = std::sqrt(std::max(stddev[i] * total - mean[i] * mean[i] , 0.)); + } + + for( j = 0; j < 2; j++ ) + { + const double * const sptr = j == 0 ? &mean[0] : &stddev[0]; + _OutputArray _dst = j == 0 ? _mean : _sdv; + if( !_dst.needed() ) + continue; + + if( !_dst.fixedSize() ) + _dst.create(cn, 1, CV_64F, -1, true); + Mat dst = _dst.getMat(); + int dcn = (int)dst.total(); + CV_Assert( dst.type() == CV_64F && dst.isContinuous() && + (dst.cols == 1 || dst.rows == 1) && dcn >= cn ); + double* dptr = dst.ptr(); + for( k = 0; k < cn; k++ ) + dptr[k] = sptr[k]; + for( ; k < dcn; k++ ) + dptr[k] = 0; + } + + return true; +} + +} void cv::meanStdDev( InputArray _src, OutputArray _mean, OutputArray _sdv, InputArray _mask ) { + if (ocl::useOpenCL() && _src.isUMat() && _mask.empty() && ocl_meanStdDev(_src, _mean, _sdv)) + return; + Mat src = _src.getMat(), mask = _mask.getMat(); CV_Assert( mask.empty() || mask.type() == CV_8U ); @@ -2602,9 +2758,8 @@ void cv::findNonZero( InputArray _src, OutputArray _idx ) double cv::PSNR(InputArray _src1, InputArray _src2) { - Mat src1 = _src1.getMat(), src2 = _src2.getMat(); - CV_Assert( src1.depth() == CV_8U ); - double diff = std::sqrt(norm(src1, src2, NORM_L2SQR)/(src1.total()*src1.channels())); + CV_Assert( _src1.depth() == CV_8U ); + double diff = std::sqrt(norm(_src1, _src2, NORM_L2SQR)/(_src1.total()*_src1.channels())); return 20*log10(255./(diff+DBL_EPSILON)); } diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp index c574d004f3..10cec7bc00 100644 --- a/modules/core/test/ocl/test_arithm.cpp +++ b/modules/core/test/ocl/test_arithm.cpp @@ -119,7 +119,6 @@ PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool) bool use_roi; cv::Scalar val; - // declare Mat + UMat mirrors TEST_DECLARE_INPUT_PARAMETER(src1) TEST_DECLARE_INPUT_PARAMETER(src2) TEST_DECLARE_INPUT_PARAMETER(mask) @@ -281,6 +280,614 @@ OCL_TEST_P(Subtract, Scalar_Mask) } } +//////////////////////////////// Mul ///////////////////////////////////////////////// + +typedef ArithmTestBase Mul; + +OCL_TEST_P(Mul, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::multiply(src1_roi, src2_roi, dst1_roi)); + OCL_ON(cv::multiply(usrc1_roi, usrc2_roi, udst1_roi)); + Near(0); + } +} + +OCL_TEST_P(Mul, DISABLED_Scalar) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::multiply(src1_roi, val, dst1_roi)); + OCL_ON(cv::multiply(val, usrc1_roi, udst1_roi)); + + Near(udst1_roi.depth() >= CV_32F ? 1e-3 : 1); + } +} + +OCL_TEST_P(Mul, DISABLED_Mat_Scale) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::multiply(src1_roi, src2_roi, dst1_roi, val[0])); + OCL_ON(cv::multiply(usrc1_roi, usrc2_roi, udst1_roi, val[0])); + + Near(udst1_roi.depth() >= CV_32F ? 1e-3 : 1); + } +} + +//////////////////////////////// Div ///////////////////////////////////////////////// + +typedef ArithmTestBase Div; + +OCL_TEST_P(Div, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::divide(src1_roi, src2_roi, dst1_roi)); + OCL_ON(cv::divide(usrc1_roi, usrc2_roi, udst1_roi)); + Near(1); + } +} + +OCL_TEST_P(Div, DISABLED_Scalar) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::divide(val, src1_roi, dst1_roi)); + OCL_ON(cv::divide(val, usrc1_roi, udst1_roi)); + + Near(udst1_roi.depth() >= CV_32F ? 1e-3 : 1); + } +} + +OCL_TEST_P(Div, Mat_Scale) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + 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 ? 4e-3 : 1); + } +} + + +OCL_TEST_P(Div, DISABLED_Mat_Scalar_Scale) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::divide(src1_roi, val, dst1_roi, val[0])); + OCL_ON(cv::divide(usrc1_roi, val, udst1_roi, val[0])); + + Near(udst1_roi.depth() >= CV_32F ? 4e-3 : 1); + } +} + +//////////////////////////////// Min/Max ///////////////////////////////////////////////// + +typedef ArithmTestBase Min; + +OCL_TEST_P(Min, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::max(src1_roi, src2_roi, dst1_roi)); + OCL_ON(cv::max(usrc1_roi, usrc2_roi, udst1_roi)); + Near(0); + } +} + +typedef ArithmTestBase Max; + +OCL_TEST_P(Max, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::min(src1_roi, src2_roi, dst1_roi)); + OCL_ON(cv::min(usrc1_roi, usrc2_roi, udst1_roi)); + Near(0); + } +} + +//////////////////////////////// Absdiff ///////////////////////////////////////////////// + +typedef ArithmTestBase Absdiff; + +OCL_TEST_P(Absdiff, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::absdiff(src1_roi, src2_roi, dst1_roi)); + OCL_ON(cv::absdiff(usrc1_roi, usrc2_roi, udst1_roi)); + Near(0); + } +} + +OCL_TEST_P(Absdiff, Scalar) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::absdiff(src1_roi, val, dst1_roi)); + OCL_ON(cv::absdiff(usrc1_roi, val, udst1_roi)); + Near(1e-5); + } +} + +//////////////////////////////// CartToPolar ///////////////////////////////////////////////// + +typedef ArithmTestBase CartToPolar; + +OCL_TEST_P(CartToPolar, angleInDegree) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::cartToPolar(src1_roi, src2_roi, dst1_roi, dst2_roi, true)); + OCL_ON(cv::cartToPolar(usrc1_roi, usrc2_roi, udst1_roi, udst2_roi, true)); + Near(0.5); + Near1(0.5); + } +} + +OCL_TEST_P(CartToPolar, angleInRadians) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::cartToPolar(src1_roi, src2_roi, dst1_roi, dst2_roi)); + OCL_ON(cv::cartToPolar(usrc1_roi, usrc2_roi, udst1_roi, udst2_roi)); + Near(0.5); + Near1(0.5); + } +} + +//////////////////////////////// PolarToCart ///////////////////////////////////////////////// + +typedef ArithmTestBase PolarToCart; + +OCL_TEST_P(PolarToCart, angleInDegree) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::polarToCart(src1_roi, src2_roi, dst1_roi, dst2_roi, true)); + OCL_ON(cv::polarToCart(usrc1_roi, usrc2_roi, udst1_roi, udst2_roi, true)); + + Near(0.5); + Near1(0.5); + } +} + +OCL_TEST_P(PolarToCart, angleInRadians) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::polarToCart(src1_roi, src2_roi, dst1_roi, dst2_roi)); + OCL_ON(cv::polarToCart(usrc1_roi, usrc2_roi, udst1_roi, udst2_roi)); + + Near(0.5); + Near1(0.5); + } +} + +//////////////////////////////// Transpose ///////////////////////////////////////////////// + +typedef ArithmTestBase Transpose; + +OCL_TEST_P(Transpose, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::transpose(src1_roi, dst1_roi)); + OCL_ON(cv::transpose(usrc1_roi, udst1_roi)); + + Near(1e-5); + } +} + +OCL_TEST_P(Transpose, SquareInplace) +{ + const int type = CV_MAKE_TYPE(depth, cn); + + for (int j = 0; j < test_loop_times; j++) + { + Size roiSize = randomSize(1, MAX_VALUE); + roiSize.height = roiSize.width; // make it square + + Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src1, src1_roi, roiSize, srcBorder, type, 5, 16); + + UMAT_UPLOAD_OUTPUT_PARAMETER(src1); + + OCL_OFF(cv::transpose(src1_roi, src1_roi)); + OCL_ON(cv::transpose(usrc1_roi, usrc1_roi)); + + EXPECT_MAT_NEAR(src1, usrc1, 0.0); + EXPECT_MAT_NEAR(src1_roi, usrc1_roi, 0.0); + } +} + +//////////////////////////////// Bitwise_and ///////////////////////////////////////////////// + +typedef ArithmTestBase Bitwise_and; + +OCL_TEST_P(Bitwise_and, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_and(src1_roi, src2_roi, dst1_roi)); + OCL_ON(cv::bitwise_and(usrc1_roi, usrc2_roi, udst1_roi)); + Near(0); + } +} + +OCL_TEST_P(Bitwise_and, Mat_Mask) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_and(src1_roi, src2_roi, dst1_roi, mask_roi)); + OCL_ON(cv::bitwise_and(usrc1_roi, usrc2_roi, udst1_roi, umask_roi)); + Near(0); + } +} + +OCL_TEST_P(Bitwise_and, Scalar) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_and(src1_roi, val, dst1_roi)); + OCL_ON(cv::bitwise_and(usrc1_roi, val, udst1_roi)); + Near(1e-5); + } +} + +OCL_TEST_P(Bitwise_and, Scalar_Mask) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_and(src1_roi, val, dst1_roi, mask_roi)); + OCL_ON(cv::bitwise_and(usrc1_roi, val, udst1_roi, umask_roi)); + Near(1e-5); + } +} + +//////////////////////////////// Bitwise_or ///////////////////////////////////////////////// + +typedef ArithmTestBase Bitwise_or; + +OCL_TEST_P(Bitwise_or, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_or(src1_roi, src2_roi, dst1_roi)); + OCL_ON(cv::bitwise_or(usrc1_roi, usrc2_roi, udst1_roi)); + Near(0); + } +} + +OCL_TEST_P(Bitwise_or, Mat_Mask) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_or(src1_roi, src2_roi, dst1_roi, mask_roi)); + OCL_ON(cv::bitwise_or(usrc1_roi, usrc2_roi, udst1_roi, umask_roi)); + Near(0); + } +} + +OCL_TEST_P(Bitwise_or, Scalar) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_or(src1_roi, val, dst1_roi)); + OCL_ON(cv::bitwise_or(usrc1_roi, val, udst1_roi)); + Near(1e-5); + } +} + +OCL_TEST_P(Bitwise_or, Scalar_Mask) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_or(src1_roi, val, dst1_roi, mask_roi)); + OCL_ON(cv::bitwise_or(val, usrc1_roi, udst1_roi, umask_roi)); + Near(1e-5); + } +} + +//////////////////////////////// Bitwise_xor ///////////////////////////////////////////////// + +typedef ArithmTestBase Bitwise_xor; + +OCL_TEST_P(Bitwise_xor, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_xor(src1_roi, src2_roi, dst1_roi)); + OCL_ON(cv::bitwise_xor(usrc1_roi, usrc2_roi, udst1_roi)); + Near(0); + } +} + +OCL_TEST_P(Bitwise_xor, Mat_Mask) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_xor(src1_roi, src2_roi, dst1_roi, mask_roi)); + OCL_ON(cv::bitwise_xor(usrc1_roi, usrc2_roi, udst1_roi, umask_roi)); + Near(0); + } +} + +OCL_TEST_P(Bitwise_xor, Scalar) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_xor(src1_roi, val, dst1_roi)); + OCL_ON(cv::bitwise_xor(usrc1_roi, val, udst1_roi)); + Near(1e-5); + } +} + +OCL_TEST_P(Bitwise_xor, Scalar_Mask) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_xor(src1_roi, val, dst1_roi, mask_roi)); + OCL_ON(cv::bitwise_xor(usrc1_roi, val, udst1_roi, umask_roi)); + Near(1e-5); + } +} + +//////////////////////////////// Bitwise_not ///////////////////////////////////////////////// + +typedef ArithmTestBase Bitwise_not; + +OCL_TEST_P(Bitwise_not, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::bitwise_not(src1_roi, dst1_roi)); + OCL_ON(cv::bitwise_not(usrc1_roi, udst1_roi)); + Near(0); + } +} + +//////////////////////////////// Compare ///////////////////////////////////////////////// + +typedef ArithmTestBase Compare; + +OCL_TEST_P(Compare, Mat) +{ + int cmp_codes[] = { CMP_EQ, CMP_GT, CMP_GE, CMP_LT, CMP_LE, CMP_NE }; + int cmp_num = sizeof(cmp_codes) / sizeof(int); + + for (int i = 0; i < cmp_num; ++i) + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::compare(src1_roi, src2_roi, dst1_roi, cmp_codes[i])); + OCL_ON(cv::compare(usrc1_roi, usrc2_roi, udst1_roi, cmp_codes[i])); + + Near(0); + } +} + +//////////////////////////////// Pow ///////////////////////////////////////////////// + +typedef ArithmTestBase Pow; + +OCL_TEST_P(Pow, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + double p = 4.5; + + OCL_OFF(cv::pow(src1_roi, p, dst1_roi)); + OCL_ON(cv::pow(usrc1_roi, p, udst1_roi)); + + Near(1); + } +} + +//////////////////////////////// AddWeighted ///////////////////////////////////////////////// + +typedef ArithmTestBase AddWeighted; + +OCL_TEST_P(AddWeighted, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + const double alpha = 2.0, beta = 1.0, gama = 3.0; + + OCL_OFF(cv::addWeighted(src1_roi, alpha, src2_roi, beta, gama, dst1_roi)); + OCL_ON(cv::addWeighted(usrc1_roi, alpha, usrc2_roi, beta, gama, udst1_roi)); + + Near(3e-4); + } +} + +//////////////////////////////// setIdentity ///////////////////////////////////////////////// + +typedef ArithmTestBase SetIdentity; + +OCL_TEST_P(SetIdentity, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + OCL_OFF(cv::setIdentity(dst1_roi, val)); + OCL_ON(cv::setIdentity(udst1_roi, val)); + + Near(0); + } +} + +//// Repeat + +struct RepeatTestCase : + public ArithmTestBase +{ + int nx, ny; + + virtual void generateTestData() + { + const int type = CV_MAKE_TYPE(depth, cn); + + nx = 2;//randomInt(1, 4); + ny = 2;//randomInt(1, 4); + + Size srcRoiSize = randomSize(1, MAX_VALUE); + Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src1, src1_roi, srcRoiSize, srcBorder, type, 2, 11); + + Size dstRoiSize(srcRoiSize.width * nx, srcRoiSize.height * ny); + Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst1, dst1_roi, dstRoiSize, dst1Border, type, 5, 16); + + UMAT_UPLOAD_INPUT_PARAMETER(src1) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst1) + } +}; + +typedef RepeatTestCase Repeat; + +OCL_TEST_P(Repeat, DISABLED_Mat) +{ + for (int i = 0; i < test_loop_times; ++i) + { + generateTestData(); + + OCL_OFF(cv::repeat(src1_roi, ny, nx, dst1_roi)); + OCL_ON(cv::repeat(usrc1_roi, ny, nx, udst1_roi)); + + Near(); + } +} + +//////////////////////////////// CountNonZero ///////////////////////////////////////////////// + +typedef ArithmTestBase CountNonZero; + +OCL_TEST_P(CountNonZero, MAT) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + int cpures, gpures; + OCL_OFF(cpures = cv::countNonZero(src1_roi)); + OCL_ON(gpures = cv::countNonZero(usrc1_roi)); + + EXPECT_EQ(cpures, gpures); + } +} + +//////////////////////////////// Sum ///////////////////////////////////////////////// + +typedef ArithmTestBase Sum; + +OCL_TEST_P(Sum, MAT) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + Scalar cpures, gpures; + OCL_OFF(cpures = cv::sum(src1_roi)); + OCL_ON(gpures = cv::sum(usrc1_roi)); + + for (int i = 0; i < cn; ++i) + EXPECT_NEAR(cpures[i], gpures[i], 0.1); + } +} + +//////////////////////////////// meanStdDev ///////////////////////////////////////////////// + +typedef ArithmTestBase MeanStdDev; + +OCL_TEST_P(MeanStdDev, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + Scalar cpu_mean, cpu_stddev; + Scalar gpu_mean, gpu_stddev; + + OCL_OFF(cv::meanStdDev(src1_roi, cpu_mean, cpu_stddev)); + OCL_ON(cv::meanStdDev(usrc1_roi, gpu_mean, gpu_stddev)); + + for (int i = 0; i < cn; ++i) + { + EXPECT_NEAR(cpu_mean[i], gpu_mean[i], 0.1); + EXPECT_NEAR(cpu_stddev[i], gpu_stddev[i], 0.1); + } + } +} + + //////////////////////////////////////// Log ///////////////////////////////////////// typedef ArithmTestBase Log; @@ -359,13 +966,33 @@ OCL_TEST_P(Magnitude, Mat) //////////////////////////////////////// Instantiation ///////////////////////////////////////// -OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, ::testing::Values(1, 2, 3, 4), Bool(), 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())); -OCL_INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(::testing::Values(CV_32F, CV_64F), ::testing::Values(1, 2, 3, 4), Bool())); -OCL_INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(::testing::Values(CV_32F, CV_64F), ::testing::Values(1, 2, 3, 4), Bool())); -OCL_INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(::testing::Values(CV_32F, CV_64F), ::testing::Values(1, 2, 3, 4), Bool())); -OCL_INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(::testing::Values(CV_32F, CV_64F), ::testing::Values(1, 2, 3, 4), Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Subtract, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Mul, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Div, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Min, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Max, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Absdiff, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, CartToPolar, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, PolarToCart, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Transpose, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_and, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_not, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_xor, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +//OCL_INSTANTIATE_TEST_CASE_P(Arithm, Bitwise_or, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Pow, Combine(testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Compare, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, SetIdentity, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Repeat, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, CountNonZero, Combine(OCL_ALL_DEPTHS, testing::Values(Channels(1)), Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Sum, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, MeanStdDev, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Log, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Exp, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); } } // namespace cvtest::ocl