diff --git a/modules/cudaarithm/include/opencv2/cudaarithm.hpp b/modules/cudaarithm/include/opencv2/cudaarithm.hpp index a482b49fcf..d040bb11f0 100644 --- a/modules/cudaarithm/include/opencv2/cudaarithm.hpp +++ b/modules/cudaarithm/include/opencv2/cudaarithm.hpp @@ -415,10 +415,10 @@ CV_EXPORTS void cartToPolar(InputArray x, InputArray y, OutputArray magnitude, O /** @brief Converts polar coordinates into Cartesian. -@param magnitude Source matrix containing magnitudes ( CV_32FC1 ). -@param angle Source matrix containing angles ( CV_32FC1 ). -@param x Destination matrix of real components ( CV_32FC1 ). -@param y Destination matrix of imaginary components ( CV_32FC1 ). +@param magnitude Source matrix containing magnitudes ( CV_32FC1 or CV_64FC1 ). +@param angle Source matrix containing angles ( same type as magnitude ). +@param x Destination matrix of real components ( same type as magnitude ). +@param y Destination matrix of imaginary components ( same type as magnitude ). @param angleInDegrees Flag that indicates angles in degrees. @param stream Stream for the asynchronous version. */ diff --git a/modules/cudaarithm/perf/perf_element_operations.cpp b/modules/cudaarithm/perf/perf_element_operations.cpp index 02f412d994..9aa2d4e4e0 100644 --- a/modules/cudaarithm/perf/perf_element_operations.cpp +++ b/modules/cudaarithm/perf/perf_element_operations.cpp @@ -1346,6 +1346,7 @@ PERF_TEST_P(Sz, MagnitudeSqr, // Phase DEF_PARAM_TEST(Sz_AngleInDegrees, cv::Size, bool); +DEF_PARAM_TEST(Sz_Type_AngleInDegrees, cv::Size, MatType, bool); PERF_TEST_P(Sz_AngleInDegrees, Phase, Combine(CUDA_TYPICAL_MAT_SIZES, @@ -1423,17 +1424,19 @@ PERF_TEST_P(Sz_AngleInDegrees, CartToPolar, ////////////////////////////////////////////////////////////////////// // PolarToCart -PERF_TEST_P(Sz_AngleInDegrees, PolarToCart, +PERF_TEST_P(Sz_Type_AngleInDegrees, PolarToCart, Combine(CUDA_TYPICAL_MAT_SIZES, + testing::Values(CV_32FC1, CV_64FC1), Bool())) { const cv::Size size = GET_PARAM(0); - const bool angleInDegrees = GET_PARAM(1); + const int type = GET_PARAM(1); + const bool angleInDegrees = GET_PARAM(2); - cv::Mat magnitude(size, CV_32FC1); + cv::Mat magnitude(size, type); declare.in(magnitude, WARMUP_RNG); - cv::Mat angle(size, CV_32FC1); + cv::Mat angle(size, type); declare.in(angle, WARMUP_RNG); if (PERF_RUN_CUDA()) diff --git a/modules/cudaarithm/src/cuda/polar_cart.cu b/modules/cudaarithm/src/cuda/polar_cart.cu index 0a949b42ed..2fb1315e61 100644 --- a/modules/cudaarithm/src/cuda/polar_cart.cu +++ b/modules/cudaarithm/src/cuda/polar_cart.cu @@ -157,8 +157,23 @@ void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, Outpu namespace { - template - __global__ void polarToCartImpl(const GlobPtr mag, const GlobPtr angle, GlobPtr xmat, GlobPtr ymat, const float scale, const int rows, const int cols) + template struct sincos_op + { + __device__ __forceinline__ void operator()(T a, T *sptr, T *cptr) const + { + ::sincos(a, sptr, cptr); + } + }; + template <> struct sincos_op + { + __device__ __forceinline__ void operator()(float a, float *sptr, float *cptr) const + { + ::sincosf(a, sptr, cptr); + } + }; + + template + __global__ void polarToCartImpl_(const GlobPtr mag, const GlobPtr angle, GlobPtr xmat, GlobPtr ymat, const T scale, const int rows, const int cols) { const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; @@ -166,45 +181,53 @@ namespace if (x >= cols || y >= rows) return; - const float mag_val = useMag ? mag(y, x) : 1.0f; - const float angle_val = angle(y, x); + const T mag_val = useMag ? mag(y, x) : static_cast(1.0); + const T angle_val = angle(y, x); - float sin_a, cos_a; - ::sincosf(scale * angle_val, &sin_a, &cos_a); + T sin_a, cos_a; + sincos_op op; + op(scale * angle_val, &sin_a, &cos_a); xmat(y, x) = mag_val * cos_a; ymat(y, x) = mag_val * sin_a; } + + template + void polarToCartImpl(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t& stream) + { + GpuMat_ xc(x.reshape(1)); + GpuMat_ yc(y.reshape(1)); + GpuMat_ magc(mag.reshape(1)); + GpuMat_ anglec(angle.reshape(1)); + + const dim3 block(32, 8); + const dim3 grid(divUp(anglec.cols, block.x), divUp(anglec.rows, block.y)); + + const T scale = angleInDegrees ? static_cast(CV_PI / 180.0) : static_cast(1.0); + + if (magc.empty()) + polarToCartImpl_ << > >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols); + else + polarToCartImpl_ << > >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols); + } } void cv::cuda::polarToCart(InputArray _mag, InputArray _angle, OutputArray _x, OutputArray _y, bool angleInDegrees, Stream& _stream) { + typedef void(*func_t)(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t& stream); + static const func_t funcs[7] = { 0, 0, 0, 0, 0, polarToCartImpl, polarToCartImpl }; + GpuMat mag = getInputMat(_mag, _stream); GpuMat angle = getInputMat(_angle, _stream); - CV_Assert( angle.depth() == CV_32F ); + CV_Assert(angle.depth() == CV_32F || angle.depth() == CV_64F); CV_Assert( mag.empty() || (mag.type() == angle.type() && mag.size() == angle.size()) ); - GpuMat x = getOutputMat(_x, angle.size(), CV_32FC1, _stream); - GpuMat y = getOutputMat(_y, angle.size(), CV_32FC1, _stream); - - GpuMat_ xc(x.reshape(1)); - GpuMat_ yc(y.reshape(1)); - GpuMat_ magc(mag.reshape(1)); - GpuMat_ anglec(angle.reshape(1)); - - const dim3 block(32, 8); - const dim3 grid(divUp(anglec.cols, block.x), divUp(anglec.rows, block.y)); - - const float scale = angleInDegrees ? (CV_PI_F / 180.0f) : 1.0f; + GpuMat x = getOutputMat(_x, angle.size(), CV_MAKETYPE(angle.depth(), 1), _stream); + GpuMat y = getOutputMat(_y, angle.size(), CV_MAKETYPE(angle.depth(), 1), _stream); cudaStream_t stream = StreamAccessor::getStream(_stream); - - if (magc.empty()) - polarToCartImpl<<>>(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols); - else - polarToCartImpl<<>>(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols); - + funcs[angle.depth()](mag, angle, x, y, angleInDegrees, stream); CV_CUDEV_SAFE_CALL( cudaGetLastError() ); syncOutput(x, _x, _stream); diff --git a/modules/cudaarithm/test/test_element_operations.cpp b/modules/cudaarithm/test/test_element_operations.cpp index cf133024e4..9d20046df9 100644 --- a/modules/cudaarithm/test/test_element_operations.cpp +++ b/modules/cudaarithm/test/test_element_operations.cpp @@ -2754,10 +2754,11 @@ INSTANTIATE_TEST_CASE_P(CUDA_Arithm, CartToPolar, testing::Combine( //////////////////////////////////////////////////////////////////////////////// // polarToCart -PARAM_TEST_CASE(PolarToCart, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, UseRoi) +PARAM_TEST_CASE(PolarToCart, cv::cuda::DeviceInfo, cv::Size, MatType, AngleInDegrees, UseRoi) { cv::cuda::DeviceInfo devInfo; cv::Size size; + int type; bool angleInDegrees; bool useRoi; @@ -2765,8 +2766,9 @@ PARAM_TEST_CASE(PolarToCart, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, Use { devInfo = GET_PARAM(0); size = GET_PARAM(1); - angleInDegrees = GET_PARAM(2); - useRoi = GET_PARAM(3); + type = GET_PARAM(2); + angleInDegrees = GET_PARAM(3); + useRoi = GET_PARAM(4); cv::cuda::setDevice(devInfo.deviceID()); } @@ -2774,24 +2776,26 @@ PARAM_TEST_CASE(PolarToCart, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, Use CUDA_TEST_P(PolarToCart, Accuracy) { - cv::Mat magnitude = randomMat(size, CV_32FC1); - cv::Mat angle = randomMat(size, CV_32FC1); + cv::Mat magnitude = randomMat(size, type); + cv::Mat angle = randomMat(size, type); + const double tol = (type == CV_32FC1 ? 1.6e-4 : 1e-4) * (angleInDegrees ? 1.0 : 19.0); - cv::cuda::GpuMat x = createMat(size, CV_32FC1, useRoi); - cv::cuda::GpuMat y = createMat(size, CV_32FC1, useRoi); + cv::cuda::GpuMat x = createMat(size, type, useRoi); + cv::cuda::GpuMat y = createMat(size, type, useRoi); cv::cuda::polarToCart(loadMat(magnitude, useRoi), loadMat(angle, useRoi), x, y, angleInDegrees); cv::Mat x_gold; cv::Mat y_gold; cv::polarToCart(magnitude, angle, x_gold, y_gold, angleInDegrees); - EXPECT_MAT_NEAR(x_gold, x, 1e-4); - EXPECT_MAT_NEAR(y_gold, y, 1e-4); + EXPECT_MAT_NEAR(x_gold, x, tol); + EXPECT_MAT_NEAR(y_gold, y, tol); } INSTANTIATE_TEST_CASE_P(CUDA_Arithm, PolarToCart, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, + testing::Values(CV_32FC1, CV_64FC1), testing::Values(AngleInDegrees(false), AngleInDegrees(true)), WHOLE_SUBMAT));