added cv::cartToPolar to T-API

This commit is contained in:
Ilya Lavrenov 2013-11-30 20:05:54 +04:00
parent 435028ccee
commit ba850f0b64
4 changed files with 80 additions and 7 deletions

View File

@ -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();

View File

@ -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

View File

@ -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)

View File

@ -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;
}