added cv::cartToPolar to T-API

diff --git a/modules/core/src/mathfuncs.cpp b/modules/core/src/mathfuncs.cpp
index 2f146c8..d216589 100644 (file)
--- 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 94ce7af..b2b164e 100644 (file)
--- 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 64460ef..4ef3fda 100644 (file)
--- 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 3b0fb99..48d5f6f 100644 (file)
--- 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;
     }