From: Vladislav Vinogradov Date: Mon, 18 Oct 2010 11:12:14 +0000 (+0000) Subject: added gpu version of magnitude, magnitudeSqr, phase, cartToPolar, polarToCart X-Git-Tag: accepted/2.0/20130307.220821~4374 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=e12e4798e09b5ee5317ce51f2f79093b7c42053b;p=profile%2Fivi%2Fopencv.git added gpu version of magnitude, magnitudeSqr, phase, cartToPolar, polarToCart --- diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index ff6f279..93a7304 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -452,20 +452,44 @@ namespace cv //! supports only CV_32FC1 type CV_EXPORTS void log(const GpuMat& a, GpuMat& b); - //! computes magnitude of each (x(i), y(i)) vector - //! supports only CV_32FC1 type - CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude); //! computes magnitude of complex (x(i).re, x(i).im) vector //! supports only CV_32FC2 type CV_EXPORTS void magnitude(const GpuMat& x, GpuMat& magnitude); - //! computes squared magnitude of each (x(i), y(i)) vector - //! supports only CV_32FC1 type - CV_EXPORTS void magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& magnitude); //! computes squared magnitude of complex (x(i).re, x(i).im) vector //! supports only CV_32FC2 type CV_EXPORTS void magnitudeSqr(const GpuMat& x, GpuMat& magnitude); + //! computes magnitude of each (x(i), y(i)) vector + //! supports only floating-point source + CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude); + //! Acync version + CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, const Stream& stream); + + //! computes squared magnitude of each (x(i), y(i)) vector + //! supports only floating-point source + CV_EXPORTS void magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& magnitude); + //! Acync version + CV_EXPORTS void magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, const Stream& stream); + + //! computes angle (angle(i)) of each (x(i), y(i)) vector + //! supports only floating-point source + CV_EXPORTS void phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees = false); + //! Acync version + CV_EXPORTS void phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees, const Stream& stream); + + //! converts Cartesian coordinates to polar + //! supports only floating-point source + CV_EXPORTS void cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, GpuMat& angle, bool angleInDegrees = false); + //! Acync version + CV_EXPORTS void cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, GpuMat& angle, bool angleInDegrees, const Stream& stream); + + //! converts polar coordinates to Cartesian + //! supports only floating-point source + CV_EXPORTS void polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees = false); + //! Acync version + CV_EXPORTS void polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, const Stream& stream); + ////////////////////////////// Image processing ////////////////////////////// //! DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation. diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index 3d66a78..2ed3f43 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -69,10 +69,18 @@ void cv::gpu::minMax(const GpuMat&, double*, double*) { throw_nogpu(); } void cv::gpu::LUT(const GpuMat&, const Mat&, GpuMat&) { throw_nogpu(); } void cv::gpu::exp(const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::log(const GpuMat&, GpuMat&) { throw_nogpu(); } -void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::magnitude(const GpuMat&, GpuMat&) { throw_nogpu(); } -void cv::gpu::magnitudeSqr(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::magnitudeSqr(const GpuMat&, GpuMat&) { throw_nogpu(); } +void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } +void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&, const Stream&) { throw_nogpu(); } +void cv::gpu::magnitudeSqr(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } +void cv::gpu::magnitudeSqr(const GpuMat&, const GpuMat&, GpuMat&, const Stream&) { throw_nogpu(); } +void cv::gpu::phase(const GpuMat&, const GpuMat&, GpuMat&, bool) { throw_nogpu(); } +void cv::gpu::phase(const GpuMat&, const GpuMat&, GpuMat&, bool, const Stream&) { throw_nogpu(); } +void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool) { throw_nogpu(); } +void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, const Stream&) { throw_nogpu(); } +void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool) { throw_nogpu(); } +void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, const Stream&) { throw_nogpu(); } #else /* !defined (HAVE_CUDA) */ @@ -611,56 +619,128 @@ void cv::gpu::log(const GpuMat& src, GpuMat& dst) } //////////////////////////////////////////////////////////////////////// -// magnitude +// NPP magnitide + +namespace +{ + typedef NppStatus (*nppMagnitude_t)(const Npp32fc* pSrc, int nSrcStep, Npp32f* pDst, int nDstStep, NppiSize oSizeROI); + + inline void npp_magnitude(const GpuMat& src, GpuMat& dst, nppMagnitude_t func) + { + CV_Assert(src.type() == CV_32FC2); + + dst.create(src.size(), CV_32FC1); + + NppiSize sz; + sz.width = src.cols; + sz.height = src.rows; + + nppSafeCall( func(src.ptr(), src.step, dst.ptr(), dst.step, sz) ); + } +} void cv::gpu::magnitude(const GpuMat& src, GpuMat& dst) { - CV_Assert(src.type() == CV_32FC2); + ::npp_magnitude(src, dst, nppiMagnitude_32fc32f_C1R); +} - dst.create(src.size(), CV_32FC1); +void cv::gpu::magnitudeSqr(const GpuMat& src, GpuMat& dst) +{ + ::npp_magnitude(src, dst, nppiMagnitudeSqr_32fc32f_C1R); +} - NppiSize sz; - sz.width = src.cols; - sz.height = src.rows; +//////////////////////////////////////////////////////////////////////// +// Polar <-> Cart - nppSafeCall( nppiMagnitude_32fc32f_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz) ); +namespace cv { namespace gpu { namespace mathfunc +{ + void cartToPolar_gpu(const DevMem2Df& x, const DevMem2Df& y, const DevMem2Df& mag, bool magSqr, const DevMem2Df& angle, bool angleInDegrees, cudaStream_t stream); + void polarToCart_gpu(const DevMem2Df& mag, const DevMem2Df& angle, const DevMem2Df& x, const DevMem2Df& y, bool angleInDegrees, cudaStream_t stream); +}}} + +namespace +{ + inline void cartToPolar_caller(const GpuMat& x, const GpuMat& y, GpuMat& mag, bool magSqr, GpuMat& angle, bool angleInDegrees, cudaStream_t stream) + { + CV_DbgAssert(x.size() == y.size() && x.type() == y.type()); + CV_Assert(x.depth() == CV_32F); + + mag.create(x.size(), x.type()); + angle.create(x.size(), x.type()); + + GpuMat x1cn = x.reshape(1); + GpuMat y1cn = y.reshape(1); + GpuMat mag1cn = mag.reshape(1); + GpuMat angle1cn = angle.reshape(1); + + mathfunc::cartToPolar_gpu(x1cn, y1cn, mag1cn, magSqr, angle1cn, angleInDegrees, stream); + } + + inline void polarToCart_caller(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t stream) + { + CV_DbgAssert((mag.empty() || mag.size() == angle.size()) && mag.type() == angle.type()); + CV_Assert(mag.depth() == CV_32F); + + x.create(mag.size(), mag.type()); + y.create(mag.size(), mag.type()); + + GpuMat mag1cn = mag.reshape(1); + GpuMat angle1cn = angle.reshape(1); + GpuMat x1cn = x.reshape(1); + GpuMat y1cn = y.reshape(1); + + mathfunc::polarToCart_gpu(mag1cn, angle1cn, x1cn, y1cn, angleInDegrees, stream); + } } -void cv::gpu::magnitude(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) +void cv::gpu::magnitude(const GpuMat& x, const GpuMat& y, GpuMat& dst) { - CV_DbgAssert(src1.type() == src2.type() && src1.size() == src2.size()); - CV_Assert(src1.type() == CV_32FC1); + ::cartToPolar_caller(x, y, dst, false, GpuMat(), false, 0); +} - GpuMat src(src1.size(), CV_32FC2); - GpuMat srcs[] = {src1, src2}; - cv::gpu::merge(srcs, 2, src); +void cv::gpu::magnitude(const GpuMat& x, const GpuMat& y, GpuMat& dst, const Stream& stream) +{ + ::cartToPolar_caller(x, y, dst, false, GpuMat(), false, StreamAccessor::getStream(stream)); +} - cv::gpu::magnitude(src, dst); +void cv::gpu::magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& dst) +{ + ::cartToPolar_caller(x, y, dst, true, GpuMat(), false, 0); } -void cv::gpu::magnitudeSqr(const GpuMat& src, GpuMat& dst) +void cv::gpu::magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& dst, const Stream& stream) { - CV_Assert(src.type() == CV_32FC2); + ::cartToPolar_caller(x, y, dst, true, GpuMat(), false, StreamAccessor::getStream(stream)); +} - dst.create(src.size(), CV_32FC1); +void cv::gpu::phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees) +{ + ::cartToPolar_caller(x, y, GpuMat(), false, angle, angleInDegrees, 0); +} - NppiSize sz; - sz.width = src.cols; - sz.height = src.rows; +void cv::gpu::phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees, const Stream& stream) +{ + ::cartToPolar_caller(x, y, GpuMat(), false, angle, angleInDegrees, StreamAccessor::getStream(stream)); +} - nppSafeCall( nppiMagnitudeSqr_32fc32f_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz) ); +void cv::gpu::cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& mag, GpuMat& angle, bool angleInDegrees) +{ + ::cartToPolar_caller(x, y, mag, false, angle, angleInDegrees, 0); } -void cv::gpu::magnitudeSqr(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) +void cv::gpu::cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& mag, GpuMat& angle, bool angleInDegrees, const Stream& stream) { - CV_DbgAssert(src1.type() == src2.type() && src1.size() == src2.size()); - CV_Assert(src1.type() == CV_32FC1); + ::cartToPolar_caller(x, y, mag, false, angle, angleInDegrees, StreamAccessor::getStream(stream)); +} - GpuMat src(src1.size(), CV_32FC2); - GpuMat srcs[] = {src1, src2}; - cv::gpu::merge(srcs, 2, src); +void cv::gpu::polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees) +{ + ::polarToCart_caller(magnitude, angle, x, y, angleInDegrees, 0); +} - cv::gpu::magnitudeSqr(src, dst); +void cv::gpu::polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, const Stream& stream) +{ + ::polarToCart_caller(magnitude, angle, x, y, angleInDegrees, StreamAccessor::getStream(stream)); } -#endif /* !defined (HAVE_CUDA) */ \ No newline at end of file +#endif /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu new file mode 100644 index 0000000..fef25ac --- /dev/null +++ b/modules/gpu/src/cuda/mathfunc.cu @@ -0,0 +1,212 @@ +/*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) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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*/ + +#include "cuda_shared.hpp" + +using namespace cv::gpu; + +#ifndef CV_PI +#define CV_PI 3.1415926535897932384626433832795f +#endif + +namespace mathfunc_krnls +{ + struct Nothing + { + static __device__ void calc(int, int, float, float, float*, size_t, float) + { + } + }; + struct Magnitude + { + static __device__ void calc(int x, int y, float x_data, float y_data, float* dst, size_t dst_step, float) + { + dst[y * dst_step + x] = sqrtf(x_data * x_data + y_data * y_data); + } + }; + struct MagnitudeSqr + { + static __device__ void calc(int x, int y, float x_data, float y_data, float* dst, size_t dst_step, float) + { + dst[y * dst_step + x] = x_data * x_data + y_data * y_data; + } + }; + struct Atan2 + { + static __device__ void calc(int x, int y, float x_data, float y_data, float* dst, size_t dst_step, float scale) + { + dst[y * dst_step + x] = scale * atan2f(y_data, x_data); + } + }; + template + __global__ void cartToPolar(const float* xptr, size_t x_step, const float* yptr, size_t y_step, + float* mag, size_t mag_step, float* angle, size_t angle_step, float scale, int width, int height) + { + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (x < width && y < height) + { + float x_data = xptr[y * x_step + x]; + float y_data = yptr[y * y_step + x]; + + Mag::calc(x, y, x_data, y_data, mag, mag_step, scale); + Angle::calc(x, y, x_data, y_data, angle, angle_step, scale); + } + } + + struct NonEmptyMag + { + static __device__ float get(const float* mag, size_t mag_step, int x, int y) + { + return mag[y * mag_step + x]; + } + }; + struct EmptyMag + { + static __device__ float get(const float*, size_t, int, int) + { + return 1.0f; + } + }; + template + __global__ void polarToCart(const float* mag, size_t mag_step, const float* angle, size_t angle_step, float scale, + float* xptr, size_t x_step, float* yptr, size_t y_step, int width, int height) + { + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (x < width && y < height) + { + float mag_data = Mag::get(mag, mag_step, x, y); + float angle_data = angle[y * angle_step + x]; + float sin_a, cos_a; + + sincosf(scale * angle_data, &sin_a, &cos_a); + + xptr[y * x_step + x] = mag_data * cos_a; + yptr[y * y_step + x] = mag_data * sin_a; + } + } +} + +namespace cv { namespace gpu { namespace mathfunc +{ + template + void cartToPolar_caller(const DevMem2Df& x, const DevMem2Df& y, const DevMem2Df& mag, const DevMem2Df& angle, bool angleInDegrees, cudaStream_t stream) + { + dim3 threads(16, 16, 1); + dim3 grid(1, 1, 1); + + grid.x = divUp(x.cols, threads.x); + grid.y = divUp(x.rows, threads.y); + + const float scale = angleInDegrees ? (float)(180.0f / CV_PI) : 1.f; + + mathfunc_krnls::cartToPolar<<>>( + x.ptr, x.step / sizeof(float), y.ptr, y.step / sizeof(float), + mag.ptr, mag.step / sizeof(float), angle.ptr, angle.step / sizeof(float), scale, x.cols, x.rows); + + if (stream == 0) + cudaSafeCall( cudaThreadSynchronize() ); + } + + void cartToPolar_gpu(const DevMem2Df& x, const DevMem2Df& y, const DevMem2Df& mag, bool magSqr, const DevMem2Df& angle, bool angleInDegrees, cudaStream_t stream) + { + typedef void (*caller_t)(const DevMem2Df& x, const DevMem2Df& y, const DevMem2Df& mag, const DevMem2Df& angle, bool angleInDegrees, cudaStream_t stream); + static const caller_t callers[2][2][2] = + { + { + { + cartToPolar_caller, + cartToPolar_caller + }, + { + cartToPolar_caller, + cartToPolar_caller, + } + }, + { + { + cartToPolar_caller, + cartToPolar_caller + }, + { + cartToPolar_caller, + cartToPolar_caller, + } + } + }; + + callers[mag.ptr == 0][magSqr][angle.ptr == 0](x, y, mag, angle, angleInDegrees, stream); + } + + template + void polarToCart_caller(const DevMem2Df& mag, const DevMem2Df& angle, const DevMem2Df& x, const DevMem2Df& y, bool angleInDegrees, cudaStream_t stream) + { + dim3 threads(16, 16, 1); + dim3 grid(1, 1, 1); + + grid.x = divUp(mag.cols, threads.x); + grid.y = divUp(mag.rows, threads.y); + + const float scale = angleInDegrees ? (float)(CV_PI / 180.0f) : 1.0f; + + mathfunc_krnls::polarToCart<<>>(mag.ptr, mag.step / sizeof(float), + angle.ptr, angle.step / sizeof(float), scale, x.ptr, x.step / sizeof(float), y.ptr, y.step / sizeof(float), mag.cols, mag.rows); + + if (stream == 0) + cudaSafeCall( cudaThreadSynchronize() ); + } + + void polarToCart_gpu(const DevMem2Df& mag, const DevMem2Df& angle, const DevMem2Df& x, const DevMem2Df& y, bool angleInDegrees, cudaStream_t stream) + { + typedef void (*caller_t)(const DevMem2Df& mag, const DevMem2Df& angle, const DevMem2Df& x, const DevMem2Df& y, bool angleInDegrees, cudaStream_t stream); + static const caller_t callers[2] = + { + polarToCart_caller, + polarToCart_caller + }; + + callers[mag.ptr == 0](mag, angle, x, y, angleInDegrees, stream); + } +}}} diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp index 67e2de7..ddc2276 100644 --- a/tests/gpu/src/arithm.cpp +++ b/tests/gpu/src/arithm.cpp @@ -81,7 +81,7 @@ int CV_GpuArithmTest::CheckNorm(const Mat& m1, const Mat& m2) { double ret = norm(m1, m2, NORM_INF); - if (ret < std::numeric_limits::epsilon()) + if (ret < 1e-5) return CvTS::OK; ts->printf(CvTS::LOG, "\nNorm: %f\n", ret); @@ -99,7 +99,7 @@ int CV_GpuArithmTest::CheckNorm(double d1, double d2) { double ret = ::fabs(d1 - d2); - if (ret < std::numeric_limits::epsilon()) + if (ret < 1e-5) return CvTS::OK; ts->printf(CvTS::LOG, "\nNorm: %f\n", ret); @@ -605,6 +605,91 @@ struct CV_GpuNppImageMagnitudeTest : public CV_GpuArithmTest } }; +//////////////////////////////////////////////////////////////////////////////// +// phase +struct CV_GpuNppImagePhaseTest : public CV_GpuArithmTest +{ + CV_GpuNppImagePhaseTest() : CV_GpuArithmTest( "GPU-NppImagePhase", "phase" ) {} + + int test( const Mat& mat1, const Mat& mat2 ) + { + if (mat1.type() != CV_32FC1) + { + ts->printf(CvTS::LOG, "\nUnsupported type\n"); + return CvTS::OK; + } + + cv::Mat cpuRes; + cv::phase(mat1, mat2, cpuRes); + + GpuMat gpu1(mat1); + GpuMat gpu2(mat2); + GpuMat gpuRes; + cv::gpu::phase(gpu1, gpu2, gpuRes); + + return CheckNorm(cpuRes, gpuRes); + } +}; + +//////////////////////////////////////////////////////////////////////////////// +// cartToPolar +struct CV_GpuNppImageCartToPolarTest : public CV_GpuArithmTest +{ + CV_GpuNppImageCartToPolarTest() : CV_GpuArithmTest( "GPU-NppImageCartToPolar", "cartToPolar" ) {} + + int test( const Mat& mat1, const Mat& mat2 ) + { + if (mat1.type() != CV_32FC1) + { + ts->printf(CvTS::LOG, "\nUnsupported type\n"); + return CvTS::OK; + } + + cv::Mat cpuMag, cpuAngle; + cv::cartToPolar(mat1, mat2, cpuMag, cpuAngle); + + GpuMat gpu1(mat1); + GpuMat gpu2(mat2); + GpuMat gpuMag, gpuAngle; + cv::gpu::cartToPolar(gpu1, gpu2, gpuMag, gpuAngle); + + int magRes = CheckNorm(cpuMag, gpuMag); + int angleRes = CheckNorm(cpuAngle, gpuAngle); + + return magRes == CvTS::OK && angleRes == CvTS::OK ? CvTS::OK : CvTS::FAIL_GENERIC; + } +}; + +//////////////////////////////////////////////////////////////////////////////// +// polarToCart +struct CV_GpuNppImagePolarToCartTest : public CV_GpuArithmTest +{ + CV_GpuNppImagePolarToCartTest() : CV_GpuArithmTest( "GPU-NppImagePolarToCart", "polarToCart" ) {} + + int test( const Mat& mat1, const Mat& mat2 ) + { + if (mat1.type() != CV_32FC1) + { + ts->printf(CvTS::LOG, "\nUnsupported type\n"); + return CvTS::OK; + } + + cv::Mat cpuX, cpuY; + cv::polarToCart(mat1, mat2, cpuX, cpuY); + + GpuMat gpu1(mat1); + GpuMat gpu2(mat2); + GpuMat gpuX, gpuY; + cv::gpu::polarToCart(gpu1, gpu2, gpuX, gpuY); + + int xRes = CheckNorm(cpuX, gpuX); + int yRes = CheckNorm(cpuY, gpuY); + + return xRes == CvTS::OK && yRes == CvTS::OK ? CvTS::OK : CvTS::FAIL_GENERIC; + } +}; + + ///////////////////////////////////////////////////////////////////////////// /////////////////// tests registration ///////////////////////////////////// ///////////////////////////////////////////////////////////////////////////// @@ -629,3 +714,6 @@ CV_GpuNppImageLUTTest CV_GpuNppImageLUT_test; CV_GpuNppImageExpTest CV_GpuNppImageExp_test; CV_GpuNppImageLogTest CV_GpuNppImageLog_test; CV_GpuNppImageMagnitudeTest CV_GpuNppImageMagnitude_test; +CV_GpuNppImagePhaseTest CV_GpuNppImagePhase_test; +CV_GpuNppImageCartToPolarTest CV_GpuNppImageCartToPolar_test; +CV_GpuNppImagePolarToCartTest CV_GpuNppImagePolarToCart_test;