Merge pull request #12576 from cv3d:fix/cuda/polar_to_cart

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.
  */

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

modules/cudaarithm/src/cuda/polar_cart.cu

@@ -157,8 +157,23 @@ void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, Outpu
 
 namespace
 {
-    template <bool useMag>
-    __global__ void polarToCartImpl(const GlobPtr<float> mag, const GlobPtr<float> angle, GlobPtr<float> xmat, GlobPtr<float> ymat, const float scale, const int rows, const int cols)
+    template <typename T> struct sincos_op
+    {
+        __device__ __forceinline__ void operator()(T a, T *sptr, T *cptr) const
+        {
+            ::sincos(a, sptr, cptr);
+        }
+    };
+    template <> struct sincos_op<float>
+    {
+        __device__ __forceinline__ void operator()(float a, float *sptr, float *cptr) const
+        {
+            ::sincosf(a, sptr, cptr);
+        }
+    };
+
+    template <typename T, bool useMag>
+    __global__ void polarToCartImpl_(const GlobPtr<T> mag, const GlobPtr<T> angle, GlobPtr<T> xmat, GlobPtr<T> 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<T>(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<T> op;
+        op(scale * angle_val, &sin_a, &cos_a);
 
         xmat(y, x) = mag_val * cos_a;
         ymat(y, x) = mag_val * sin_a;
     }
+
+    template <typename T>
+    void polarToCartImpl(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t& stream)
+    {
+        GpuMat_<T> xc(x.reshape(1));
+        GpuMat_<T> yc(y.reshape(1));
+        GpuMat_<T> magc(mag.reshape(1));
+        GpuMat_<T> 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<T>(CV_PI / 180.0) : static_cast<T>(1.0);
+
+        if (magc.empty())
+            polarToCartImpl_<T, false> << <grid, block, 0, stream >> >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols);
+        else
+            polarToCartImpl_<T, true> << <grid, block, 0, stream >> >(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<float>, polarToCartImpl<double> };
+
     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_<float> xc(x.reshape(1));
-    GpuMat_<float> yc(y.reshape(1));
-    GpuMat_<float> magc(mag.reshape(1));
-    GpuMat_<float> 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<false><<<grid, block, 0, stream>>>(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols);
-    else
-        polarToCartImpl<true><<<grid, block, 0, stream>>>(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);

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