diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp index 2f146c8559..d216589e70 100644 --- a/modules/core/src/mathfuncs.cpp +++ b/modules/core/src/mathfuncs.cpp @@ -497,10 +497,50 @@ 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; + + 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(); + + 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" : "")); + + 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)*/) + { + CV_Assert(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)); @@ -1970,7 +2010,7 @@ static bool ocl_pow(InputArray _src, double power, OutputArray _dst) 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" : "")); + doubleSupport ? " -D DOUBLE_SUPPORT" : "")); ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src), dstarg = ocl::KernelArg::WriteOnly(dst, cn); @@ -1986,8 +2026,11 @@ static bool ocl_pow(InputArray _src, double power, OutputArray _dst) void pow( InputArray _src, double power, OutputArray _dst ) { - if (ocl::useOpenCL() && _dst.isUMat() && ocl_pow(_src, power, _dst)) + if (ocl::useOpenCL() && _dst.isUMat() /*&& ocl_pow(_src, power, _dst)*/) + { + CV_Assert(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 94ce7aff96..b2b164e45f 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -2647,7 +2647,7 @@ void cv::transpose( InputArray _src, OutputArray _dst ) { TransposeInplaceFunc func = transposeInplaceTab[esz]; CV_Assert( func != 0 ); -// CV_Assert( dst.cols == dst.rows ); + 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..4ef3fdad2a 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -1893,7 +1893,7 @@ Context2& Context2::getDefault() // First, try to retrieve existing context of the same type. // In its turn, Platform::getContext() may call Context2::create() // if there is no such context. - ctx.create(Device::TYPE_ACCELERATOR); + ctx.create(Device::TYPE_CPU); if(!ctx.p) ctx.create(Device::TYPE_DGPU); if(!ctx.p) diff --git a/modules/core/src/opencl/arithm.cl b/modules/core/src/opencl/arithm.cl index 3b0fb99cb8..48d5f6fb43 100644 --- a/modules/core/src/opencl/arithm.cl +++ b/modules/core/src/opencl/arithm.cl @@ -63,12 +63,15 @@ #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 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) @@ -193,10 +197,35 @@ 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 + #else #error "unknown op type" #endif +#if defined OP_CTP_AD || defined OP_CTP_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 || \ @@ -222,6 +251,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; }