//! supports only CV_32FC1 type\r
CV_EXPORTS void log(const GpuMat& a, GpuMat& b);\r
\r
- //! computes magnitude of each (x(i), y(i)) vector\r
- //! supports only CV_32FC1 type\r
- CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude);\r
//! computes magnitude of complex (x(i).re, x(i).im) vector\r
//! supports only CV_32FC2 type\r
CV_EXPORTS void magnitude(const GpuMat& x, GpuMat& magnitude);\r
\r
- //! computes squared magnitude of each (x(i), y(i)) vector\r
- //! supports only CV_32FC1 type\r
- CV_EXPORTS void magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& magnitude);\r
//! computes squared magnitude of complex (x(i).re, x(i).im) vector\r
//! supports only CV_32FC2 type\r
CV_EXPORTS void magnitudeSqr(const GpuMat& x, GpuMat& magnitude);\r
\r
+ //! computes magnitude of each (x(i), y(i)) vector\r
+ //! supports only floating-point source\r
+ CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude);\r
+ //! Acync version\r
+ CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, const Stream& stream);\r
+ \r
+ //! computes squared magnitude of each (x(i), y(i)) vector\r
+ //! supports only floating-point source\r
+ CV_EXPORTS void magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& magnitude);\r
+ //! Acync version\r
+ CV_EXPORTS void magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, const Stream& stream);\r
+\r
+ //! computes angle (angle(i)) of each (x(i), y(i)) vector\r
+ //! supports only floating-point source\r
+ CV_EXPORTS void phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees = false);\r
+ //! Acync version\r
+ CV_EXPORTS void phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees, const Stream& stream);\r
+\r
+ //! converts Cartesian coordinates to polar\r
+ //! supports only floating-point source\r
+ CV_EXPORTS void cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, GpuMat& angle, bool angleInDegrees = false);\r
+ //! Acync version\r
+ CV_EXPORTS void cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, GpuMat& angle, bool angleInDegrees, const Stream& stream);\r
+\r
+ //! converts polar coordinates to Cartesian\r
+ //! supports only floating-point source\r
+ CV_EXPORTS void polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees = false);\r
+ //! Acync version\r
+ CV_EXPORTS void polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, const Stream& stream);\r
+ \r
////////////////////////////// Image processing //////////////////////////////\r
\r
//! DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation.\r
void cv::gpu::LUT(const GpuMat&, const Mat&, GpuMat&) { throw_nogpu(); }\r
void cv::gpu::exp(const GpuMat&, GpuMat&) { throw_nogpu(); }\r
void cv::gpu::log(const GpuMat&, GpuMat&) { throw_nogpu(); }\r
-void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
void cv::gpu::magnitude(const GpuMat&, GpuMat&) { throw_nogpu(); }\r
-void cv::gpu::magnitudeSqr(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
void cv::gpu::magnitudeSqr(const GpuMat&, GpuMat&) { throw_nogpu(); }\r
+void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
+void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&, const Stream&) { throw_nogpu(); }\r
+void cv::gpu::magnitudeSqr(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }\r
+void cv::gpu::magnitudeSqr(const GpuMat&, const GpuMat&, GpuMat&, const Stream&) { throw_nogpu(); }\r
+void cv::gpu::phase(const GpuMat&, const GpuMat&, GpuMat&, bool) { throw_nogpu(); }\r
+void cv::gpu::phase(const GpuMat&, const GpuMat&, GpuMat&, bool, const Stream&) { throw_nogpu(); }\r
+void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool) { throw_nogpu(); }\r
+void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, const Stream&) { throw_nogpu(); }\r
+void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool) { throw_nogpu(); }\r
+void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, const Stream&) { throw_nogpu(); }\r
\r
#else /* !defined (HAVE_CUDA) */\r
\r
}\r
\r
////////////////////////////////////////////////////////////////////////\r
-// magnitude\r
+// NPP magnitide\r
+\r
+namespace\r
+{\r
+ typedef NppStatus (*nppMagnitude_t)(const Npp32fc* pSrc, int nSrcStep, Npp32f* pDst, int nDstStep, NppiSize oSizeROI);\r
+\r
+ inline void npp_magnitude(const GpuMat& src, GpuMat& dst, nppMagnitude_t func)\r
+ {\r
+ CV_Assert(src.type() == CV_32FC2);\r
+\r
+ dst.create(src.size(), CV_32FC1);\r
+\r
+ NppiSize sz;\r
+ sz.width = src.cols;\r
+ sz.height = src.rows;\r
+\r
+ nppSafeCall( func(src.ptr<Npp32fc>(), src.step, dst.ptr<Npp32f>(), dst.step, sz) );\r
+ }\r
+}\r
\r
void cv::gpu::magnitude(const GpuMat& src, GpuMat& dst)\r
{\r
- CV_Assert(src.type() == CV_32FC2);\r
+ ::npp_magnitude(src, dst, nppiMagnitude_32fc32f_C1R);\r
+}\r
\r
- dst.create(src.size(), CV_32FC1);\r
+void cv::gpu::magnitudeSqr(const GpuMat& src, GpuMat& dst)\r
+{\r
+ ::npp_magnitude(src, dst, nppiMagnitudeSqr_32fc32f_C1R);\r
+}\r
\r
- NppiSize sz;\r
- sz.width = src.cols;\r
- sz.height = src.rows;\r
+////////////////////////////////////////////////////////////////////////\r
+// Polar <-> Cart\r
\r
- nppSafeCall( nppiMagnitude_32fc32f_C1R(src.ptr<Npp32fc>(), src.step, dst.ptr<Npp32f>(), dst.step, sz) );\r
+namespace cv { namespace gpu { namespace mathfunc \r
+{\r
+ void cartToPolar_gpu(const DevMem2Df& x, const DevMem2Df& y, const DevMem2Df& mag, bool magSqr, const DevMem2Df& angle, bool angleInDegrees, cudaStream_t stream);\r
+ void polarToCart_gpu(const DevMem2Df& mag, const DevMem2Df& angle, const DevMem2Df& x, const DevMem2Df& y, bool angleInDegrees, cudaStream_t stream);\r
+}}}\r
+\r
+namespace\r
+{\r
+ inline void cartToPolar_caller(const GpuMat& x, const GpuMat& y, GpuMat& mag, bool magSqr, GpuMat& angle, bool angleInDegrees, cudaStream_t stream)\r
+ {\r
+ CV_DbgAssert(x.size() == y.size() && x.type() == y.type());\r
+ CV_Assert(x.depth() == CV_32F);\r
+\r
+ mag.create(x.size(), x.type());\r
+ angle.create(x.size(), x.type());\r
+\r
+ GpuMat x1cn = x.reshape(1);\r
+ GpuMat y1cn = y.reshape(1);\r
+ GpuMat mag1cn = mag.reshape(1);\r
+ GpuMat angle1cn = angle.reshape(1);\r
+\r
+ mathfunc::cartToPolar_gpu(x1cn, y1cn, mag1cn, magSqr, angle1cn, angleInDegrees, stream);\r
+ }\r
+\r
+ inline void polarToCart_caller(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t stream)\r
+ {\r
+ CV_DbgAssert((mag.empty() || mag.size() == angle.size()) && mag.type() == angle.type());\r
+ CV_Assert(mag.depth() == CV_32F);\r
+\r
+ x.create(mag.size(), mag.type());\r
+ y.create(mag.size(), mag.type());\r
+\r
+ GpuMat mag1cn = mag.reshape(1);\r
+ GpuMat angle1cn = angle.reshape(1);\r
+ GpuMat x1cn = x.reshape(1);\r
+ GpuMat y1cn = y.reshape(1);\r
+\r
+ mathfunc::polarToCart_gpu(mag1cn, angle1cn, x1cn, y1cn, angleInDegrees, stream);\r
+ }\r
}\r
\r
-void cv::gpu::magnitude(const GpuMat& src1, const GpuMat& src2, GpuMat& dst)\r
+void cv::gpu::magnitude(const GpuMat& x, const GpuMat& y, GpuMat& dst)\r
{\r
- CV_DbgAssert(src1.type() == src2.type() && src1.size() == src2.size());\r
- CV_Assert(src1.type() == CV_32FC1);\r
+ ::cartToPolar_caller(x, y, dst, false, GpuMat(), false, 0);\r
+}\r
\r
- GpuMat src(src1.size(), CV_32FC2);\r
- GpuMat srcs[] = {src1, src2};\r
- cv::gpu::merge(srcs, 2, src);\r
+void cv::gpu::magnitude(const GpuMat& x, const GpuMat& y, GpuMat& dst, const Stream& stream)\r
+{\r
+ ::cartToPolar_caller(x, y, dst, false, GpuMat(), false, StreamAccessor::getStream(stream));\r
+}\r
\r
- cv::gpu::magnitude(src, dst);\r
+void cv::gpu::magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& dst)\r
+{\r
+ ::cartToPolar_caller(x, y, dst, true, GpuMat(), false, 0);\r
}\r
\r
-void cv::gpu::magnitudeSqr(const GpuMat& src, GpuMat& dst)\r
+void cv::gpu::magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& dst, const Stream& stream)\r
{\r
- CV_Assert(src.type() == CV_32FC2);\r
+ ::cartToPolar_caller(x, y, dst, true, GpuMat(), false, StreamAccessor::getStream(stream));\r
+}\r
\r
- dst.create(src.size(), CV_32FC1);\r
+void cv::gpu::phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees)\r
+{\r
+ ::cartToPolar_caller(x, y, GpuMat(), false, angle, angleInDegrees, 0);\r
+}\r
\r
- NppiSize sz;\r
- sz.width = src.cols;\r
- sz.height = src.rows;\r
+void cv::gpu::phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees, const Stream& stream)\r
+{ \r
+ ::cartToPolar_caller(x, y, GpuMat(), false, angle, angleInDegrees, StreamAccessor::getStream(stream));\r
+}\r
\r
- nppSafeCall( nppiMagnitudeSqr_32fc32f_C1R(src.ptr<Npp32fc>(), src.step, dst.ptr<Npp32f>(), dst.step, sz) );\r
+void cv::gpu::cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& mag, GpuMat& angle, bool angleInDegrees)\r
+{\r
+ ::cartToPolar_caller(x, y, mag, false, angle, angleInDegrees, 0);\r
}\r
\r
-void cv::gpu::magnitudeSqr(const GpuMat& src1, const GpuMat& src2, GpuMat& dst)\r
+void cv::gpu::cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& mag, GpuMat& angle, bool angleInDegrees, const Stream& stream)\r
{\r
- CV_DbgAssert(src1.type() == src2.type() && src1.size() == src2.size());\r
- CV_Assert(src1.type() == CV_32FC1);\r
+ ::cartToPolar_caller(x, y, mag, false, angle, angleInDegrees, StreamAccessor::getStream(stream));\r
+}\r
\r
- GpuMat src(src1.size(), CV_32FC2);\r
- GpuMat srcs[] = {src1, src2};\r
- cv::gpu::merge(srcs, 2, src);\r
+void cv::gpu::polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees)\r
+{\r
+ ::polarToCart_caller(magnitude, angle, x, y, angleInDegrees, 0);\r
+}\r
\r
- cv::gpu::magnitudeSqr(src, dst);\r
+void cv::gpu::polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, const Stream& stream)\r
+{\r
+ ::polarToCart_caller(magnitude, angle, x, y, angleInDegrees, StreamAccessor::getStream(stream));\r
}\r
\r
-#endif /* !defined (HAVE_CUDA) */
\ No newline at end of file
+#endif /* !defined (HAVE_CUDA) */\r
--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "cuda_shared.hpp"\r
+\r
+using namespace cv::gpu;\r
+\r
+#ifndef CV_PI\r
+#define CV_PI 3.1415926535897932384626433832795f\r
+#endif\r
+\r
+namespace mathfunc_krnls \r
+{\r
+ struct Nothing\r
+ {\r
+ static __device__ void calc(int, int, float, float, float*, size_t, float)\r
+ {\r
+ }\r
+ };\r
+ struct Magnitude\r
+ {\r
+ static __device__ void calc(int x, int y, float x_data, float y_data, float* dst, size_t dst_step, float)\r
+ {\r
+ dst[y * dst_step + x] = sqrtf(x_data * x_data + y_data * y_data);\r
+ }\r
+ };\r
+ struct MagnitudeSqr\r
+ {\r
+ static __device__ void calc(int x, int y, float x_data, float y_data, float* dst, size_t dst_step, float)\r
+ {\r
+ dst[y * dst_step + x] = x_data * x_data + y_data * y_data;\r
+ }\r
+ };\r
+ struct Atan2\r
+ {\r
+ static __device__ void calc(int x, int y, float x_data, float y_data, float* dst, size_t dst_step, float scale)\r
+ {\r
+ dst[y * dst_step + x] = scale * atan2f(y_data, x_data);\r
+ }\r
+ };\r
+ template <typename Mag, typename Angle>\r
+ __global__ void cartToPolar(const float* xptr, size_t x_step, const float* yptr, size_t y_step, \r
+ float* mag, size_t mag_step, float* angle, size_t angle_step, float scale, int width, int height)\r
+ {\r
+ const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+ const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+\r
+ if (x < width && y < height)\r
+ {\r
+ float x_data = xptr[y * x_step + x];\r
+ float y_data = yptr[y * y_step + x];\r
+\r
+ Mag::calc(x, y, x_data, y_data, mag, mag_step, scale);\r
+ Angle::calc(x, y, x_data, y_data, angle, angle_step, scale);\r
+ }\r
+ }\r
+\r
+ struct NonEmptyMag\r
+ {\r
+ static __device__ float get(const float* mag, size_t mag_step, int x, int y)\r
+ {\r
+ return mag[y * mag_step + x];\r
+ }\r
+ };\r
+ struct EmptyMag\r
+ {\r
+ static __device__ float get(const float*, size_t, int, int)\r
+ {\r
+ return 1.0f;\r
+ }\r
+ };\r
+ template <typename Mag>\r
+ __global__ void polarToCart(const float* mag, size_t mag_step, const float* angle, size_t angle_step, float scale,\r
+ float* xptr, size_t x_step, float* yptr, size_t y_step, int width, int height)\r
+ {\r
+ const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+ const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+\r
+ if (x < width && y < height)\r
+ {\r
+ float mag_data = Mag::get(mag, mag_step, x, y);\r
+ float angle_data = angle[y * angle_step + x];\r
+ float sin_a, cos_a;\r
+\r
+ sincosf(scale * angle_data, &sin_a, &cos_a);\r
+\r
+ xptr[y * x_step + x] = mag_data * cos_a;\r
+ yptr[y * y_step + x] = mag_data * sin_a;\r
+ }\r
+ }\r
+}\r
+\r
+namespace cv { namespace gpu { namespace mathfunc \r
+{\r
+ template <typename Mag, typename Angle>\r
+ void cartToPolar_caller(const DevMem2Df& x, const DevMem2Df& y, const DevMem2Df& mag, const DevMem2Df& angle, bool angleInDegrees, cudaStream_t stream)\r
+ {\r
+ dim3 threads(16, 16, 1);\r
+ dim3 grid(1, 1, 1);\r
+\r
+ grid.x = divUp(x.cols, threads.x);\r
+ grid.y = divUp(x.rows, threads.y);\r
+ \r
+ const float scale = angleInDegrees ? (float)(180.0f / CV_PI) : 1.f;\r
+\r
+ mathfunc_krnls::cartToPolar<Mag, Angle><<<grid, threads, 0, stream>>>(\r
+ x.ptr, x.step / sizeof(float), y.ptr, y.step / sizeof(float), \r
+ mag.ptr, mag.step / sizeof(float), angle.ptr, angle.step / sizeof(float), scale, x.cols, x.rows);\r
+\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaThreadSynchronize() );\r
+ }\r
+\r
+ void cartToPolar_gpu(const DevMem2Df& x, const DevMem2Df& y, const DevMem2Df& mag, bool magSqr, const DevMem2Df& angle, bool angleInDegrees, cudaStream_t stream)\r
+ {\r
+ typedef void (*caller_t)(const DevMem2Df& x, const DevMem2Df& y, const DevMem2Df& mag, const DevMem2Df& angle, bool angleInDegrees, cudaStream_t stream);\r
+ static const caller_t callers[2][2][2] = \r
+ {\r
+ {\r
+ {\r
+ cartToPolar_caller<mathfunc_krnls::Magnitude, mathfunc_krnls::Atan2>,\r
+ cartToPolar_caller<mathfunc_krnls::Magnitude, mathfunc_krnls::Nothing>\r
+ },\r
+ {\r
+ cartToPolar_caller<mathfunc_krnls::MagnitudeSqr, mathfunc_krnls::Atan2>,\r
+ cartToPolar_caller<mathfunc_krnls::MagnitudeSqr, mathfunc_krnls::Nothing>,\r
+ }\r
+ },\r
+ {\r
+ {\r
+ cartToPolar_caller<mathfunc_krnls::Nothing, mathfunc_krnls::Atan2>,\r
+ cartToPolar_caller<mathfunc_krnls::Nothing, mathfunc_krnls::Nothing>\r
+ },\r
+ {\r
+ cartToPolar_caller<mathfunc_krnls::Nothing, mathfunc_krnls::Atan2>,\r
+ cartToPolar_caller<mathfunc_krnls::Nothing, mathfunc_krnls::Nothing>,\r
+ }\r
+ }\r
+ };\r
+\r
+ callers[mag.ptr == 0][magSqr][angle.ptr == 0](x, y, mag, angle, angleInDegrees, stream);\r
+ }\r
+\r
+ template <typename Mag>\r
+ void polarToCart_caller(const DevMem2Df& mag, const DevMem2Df& angle, const DevMem2Df& x, const DevMem2Df& y, bool angleInDegrees, cudaStream_t stream)\r
+ {\r
+ dim3 threads(16, 16, 1);\r
+ dim3 grid(1, 1, 1);\r
+\r
+ grid.x = divUp(mag.cols, threads.x);\r
+ grid.y = divUp(mag.rows, threads.y);\r
+ \r
+ const float scale = angleInDegrees ? (float)(CV_PI / 180.0f) : 1.0f;\r
+\r
+ mathfunc_krnls::polarToCart<Mag><<<grid, threads, 0, stream>>>(mag.ptr, mag.step / sizeof(float), \r
+ angle.ptr, angle.step / sizeof(float), scale, x.ptr, x.step / sizeof(float), y.ptr, y.step / sizeof(float), mag.cols, mag.rows);\r
+\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaThreadSynchronize() );\r
+ }\r
+\r
+ void polarToCart_gpu(const DevMem2Df& mag, const DevMem2Df& angle, const DevMem2Df& x, const DevMem2Df& y, bool angleInDegrees, cudaStream_t stream)\r
+ {\r
+ typedef void (*caller_t)(const DevMem2Df& mag, const DevMem2Df& angle, const DevMem2Df& x, const DevMem2Df& y, bool angleInDegrees, cudaStream_t stream);\r
+ static const caller_t callers[2] = \r
+ {\r
+ polarToCart_caller<mathfunc_krnls::NonEmptyMag>,\r
+ polarToCart_caller<mathfunc_krnls::EmptyMag>\r
+ };\r
+\r
+ callers[mag.ptr == 0](mag, angle, x, y, angleInDegrees, stream);\r
+ }\r
+}}}\r
{\r
double ret = norm(m1, m2, NORM_INF);\r
\r
- if (ret < std::numeric_limits<double>::epsilon())\r
+ if (ret < 1e-5)\r
return CvTS::OK;\r
\r
ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);\r
{\r
double ret = ::fabs(d1 - d2);\r
\r
- if (ret < std::numeric_limits<double>::epsilon())\r
+ if (ret < 1e-5)\r
return CvTS::OK;\r
\r
ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);\r
}\r
};\r
\r
+////////////////////////////////////////////////////////////////////////////////\r
+// phase\r
+struct CV_GpuNppImagePhaseTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImagePhaseTest() : CV_GpuArithmTest( "GPU-NppImagePhase", "phase" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuRes;\r
+ cv::phase(mat1, mat2, cpuRes);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuRes;\r
+ cv::gpu::phase(gpu1, gpu2, gpuRes);\r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// cartToPolar\r
+struct CV_GpuNppImageCartToPolarTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageCartToPolarTest() : CV_GpuArithmTest( "GPU-NppImageCartToPolar", "cartToPolar" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuMag, cpuAngle;\r
+ cv::cartToPolar(mat1, mat2, cpuMag, cpuAngle);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuMag, gpuAngle;\r
+ cv::gpu::cartToPolar(gpu1, gpu2, gpuMag, gpuAngle);\r
+\r
+ int magRes = CheckNorm(cpuMag, gpuMag);\r
+ int angleRes = CheckNorm(cpuAngle, gpuAngle);\r
+\r
+ return magRes == CvTS::OK && angleRes == CvTS::OK ? CvTS::OK : CvTS::FAIL_GENERIC;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// polarToCart\r
+struct CV_GpuNppImagePolarToCartTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImagePolarToCartTest() : CV_GpuArithmTest( "GPU-NppImagePolarToCart", "polarToCart" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuX, cpuY;\r
+ cv::polarToCart(mat1, mat2, cpuX, cpuY);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuX, gpuY;\r
+ cv::gpu::polarToCart(gpu1, gpu2, gpuX, gpuY);\r
+\r
+ int xRes = CheckNorm(cpuX, gpuX);\r
+ int yRes = CheckNorm(cpuY, gpuY);\r
+\r
+ return xRes == CvTS::OK && yRes == CvTS::OK ? CvTS::OK : CvTS::FAIL_GENERIC;\r
+ }\r
+};\r
+\r
+\r
/////////////////////////////////////////////////////////////////////////////\r
/////////////////// tests registration /////////////////////////////////////\r
/////////////////////////////////////////////////////////////////////////////\r
CV_GpuNppImageExpTest CV_GpuNppImageExp_test;\r
CV_GpuNppImageLogTest CV_GpuNppImageLog_test;\r
CV_GpuNppImageMagnitudeTest CV_GpuNppImageMagnitude_test;\r
+CV_GpuNppImagePhaseTest CV_GpuNppImagePhase_test;\r
+CV_GpuNppImageCartToPolarTest CV_GpuNppImageCartToPolar_test;\r
+CV_GpuNppImagePolarToCartTest CV_GpuNppImagePolarToCart_test;\r
projects / profile / ivi / opencv.git / commitdiff