diff --git a/modules/cudaarithm/include/opencv2/cudaarithm.hpp b/modules/cudaarithm/include/opencv2/cudaarithm.hpp index bb74ea18918..a16c271881e 100644 --- a/modules/cudaarithm/include/opencv2/cudaarithm.hpp +++ b/modules/cudaarithm/include/opencv2/cudaarithm.hpp @@ -433,6 +433,17 @@ CV_EXPORTS_W void magnitudeSqr(InputArray x, InputArray y, OutputArray magnitude */ CV_EXPORTS_W void phase(InputArray x, InputArray y, OutputArray angle, bool angleInDegrees = false, Stream& stream = Stream::Null()); +/** @brief Computes polar angles of complex matrix elements. + +@param xy Source matrix containing real and imaginary components ( CV_32FC2 ). +@param angle Destination matrix of angles ( CV_32FC1 ). +@param angleInDegrees Flag for angles that must be evaluated in degrees. +@param stream Stream for the asynchronous version. + +@sa phase +*/ +CV_EXPORTS_W void phase(InputArray xy, OutputArray angle, bool angleInDegrees = false, Stream& stream = Stream::Null()); + /** @brief Converts Cartesian coordinates into polar. @param x Source matrix containing real components ( CV_32FC1 ). @@ -446,6 +457,29 @@ CV_EXPORTS_W void phase(InputArray x, InputArray y, OutputArray angle, bool angl */ CV_EXPORTS_W void cartToPolar(InputArray x, InputArray y, OutputArray magnitude, OutputArray angle, bool angleInDegrees = false, Stream& stream = Stream::Null()); +/** @brief Converts Cartesian coordinates into polar. + +@param xy Source matrix containing real and imaginary components ( CV_32FC2 ). +@param magnitude Destination matrix of float magnitudes ( CV_32FC1 ). +@param angle Destination matrix of angles ( CV_32FC1 ). +@param angleInDegrees Flag for angles that must be evaluated in degrees. +@param stream Stream for the asynchronous version. + +@sa cartToPolar +*/ +CV_EXPORTS_W void cartToPolar(InputArray xy, OutputArray magnitude, OutputArray angle, bool angleInDegrees = false, Stream& stream = Stream::Null()); + +/** @brief Converts Cartesian coordinates into polar. + +@param xy Source matrix containing real and imaginary components ( CV_32FC2 ). +@param magnitudeAngle Destination matrix of float magnitudes and angles ( CV_32FC2 ). +@param angleInDegrees Flag for angles that must be evaluated in degrees. +@param stream Stream for the asynchronous version. + +@sa cartToPolar +*/ +CV_EXPORTS_W void cartToPolar(InputArray xy, OutputArray magnitudeAngle, bool angleInDegrees = false, Stream& stream = Stream::Null()); + /** @brief Converts polar coordinates into Cartesian. @param magnitude Source matrix containing magnitudes ( CV_32FC1 or CV_64FC1 ). @@ -457,6 +491,25 @@ CV_EXPORTS_W void cartToPolar(InputArray x, InputArray y, OutputArray magnitude, */ CV_EXPORTS_W void polarToCart(InputArray magnitude, InputArray angle, OutputArray x, OutputArray y, bool angleInDegrees = false, Stream& stream = Stream::Null()); +/** @brief Converts polar coordinates into Cartesian. + +@param magnitude Source matrix containing magnitudes ( CV_32FC1 or CV_64FC1 ). +@param angle Source matrix containing angles ( same type as magnitude ). +@param xy Destination matrix of real and imaginary components ( same depth as magnitude, i.e. CV_32FC2 or CV_64FC2 ). +@param angleInDegrees Flag that indicates angles in degrees. +@param stream Stream for the asynchronous version. +*/ +CV_EXPORTS_W void polarToCart(InputArray magnitude, InputArray angle, OutputArray xy, bool angleInDegrees = false, Stream& stream = Stream::Null()); + +/** @brief Converts polar coordinates into Cartesian. + +@param magnitudeAngle Source matrix containing magnitudes and angles ( CV_32FC2 or CV_64FC2 ). +@param xy Destination matrix of real and imaginary components ( same depth as source ). +@param angleInDegrees Flag that indicates angles in degrees. +@param stream Stream for the asynchronous version. +*/ +CV_EXPORTS_W void polarToCart(InputArray magnitudeAngle, OutputArray xy, bool angleInDegrees = false, Stream& stream = Stream::Null()); + //! @} cudaarithm_elem //! @addtogroup cudaarithm_core diff --git a/modules/cudaarithm/src/cuda/polar_cart.cu b/modules/cudaarithm/src/cuda/polar_cart.cu index 12980e424ff..725f5741d81 100644 --- a/modules/cudaarithm/src/cuda/polar_cart.cu +++ b/modules/cudaarithm/src/cuda/polar_cart.cu @@ -52,8 +52,10 @@ #include "opencv2/cudev.hpp" #include "opencv2/core/private.cuda.hpp" -using namespace cv; -using namespace cv::cuda; +//do not use implicit cv::cuda to avoid clash of tuples from ::cuda::std +/*using namespace cv; +using namespace cv::cuda;*/ + using namespace cv::cudev; void cv::cuda::magnitude(InputArray _x, InputArray _y, OutputArray _dst, Stream& stream) @@ -66,11 +68,7 @@ void cv::cuda::magnitude(InputArray _x, InputArray _y, OutputArray _dst, Stream& GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream); - GpuMat_ xc(x.reshape(1)); - GpuMat_ yc(y.reshape(1)); - GpuMat_ magc(dst.reshape(1)); - - gridTransformBinary(xc, yc, magc, magnitude_func(), stream); + gridTransformBinary(globPtr(x), globPtr(y), globPtr(dst), magnitude_func(), stream); syncOutput(dst, _dst, stream); } @@ -85,11 +83,7 @@ void cv::cuda::magnitudeSqr(InputArray _x, InputArray _y, OutputArray _dst, Stre GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream); - GpuMat_ xc(x.reshape(1)); - GpuMat_ yc(y.reshape(1)); - GpuMat_ magc(dst.reshape(1)); - - gridTransformBinary(xc, yc, magc, magnitude_sqr_func(), stream); + gridTransformBinary(globPtr(x), globPtr(y), globPtr(dst), magnitude_sqr_func(), stream); syncOutput(dst, _dst, stream); } @@ -104,14 +98,26 @@ void cv::cuda::phase(InputArray _x, InputArray _y, OutputArray _dst, bool angleI GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream); - GpuMat_ xc(x.reshape(1)); - GpuMat_ yc(y.reshape(1)); - GpuMat_ anglec(dst.reshape(1)); + if (angleInDegrees) + gridTransformBinary(globPtr(x), globPtr(y), globPtr(dst), direction_func(), stream); + else + gridTransformBinary(globPtr(x), globPtr(y), globPtr(dst), direction_func(), stream); + + syncOutput(dst, _dst, stream); +} + +void cv::cuda::phase(InputArray _xy, OutputArray _dst, bool angleInDegrees, Stream& stream) +{ + GpuMat xy = getInputMat(_xy, stream); + + CV_Assert( xy.type() == CV_32FC2 ); + + GpuMat dst = getOutputMat(_dst, xy.size(), CV_32FC1, stream); if (angleInDegrees) - gridTransformBinary(xc, yc, anglec, direction_func(), stream); + gridTransformUnary(globPtr(xy), globPtr(dst), direction_interleaved_func(), stream); else - gridTransformBinary(xc, yc, anglec, direction_func(), stream); + gridTransformUnary(globPtr(xy), globPtr(dst), direction_interleaved_func(), stream); syncOutput(dst, _dst, stream); } @@ -127,10 +133,10 @@ void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, Outpu GpuMat mag = getOutputMat(_mag, x.size(), CV_32FC1, stream); GpuMat angle = getOutputMat(_angle, x.size(), CV_32FC1, stream); - GpuMat_ xc(x.reshape(1)); - GpuMat_ yc(y.reshape(1)); - GpuMat_ magc(mag.reshape(1)); - GpuMat_ anglec(angle.reshape(1)); + GpuMat_ xc(x); + GpuMat_ yc(y); + GpuMat_ magc(mag); + GpuMat_ anglec(angle); if (angleInDegrees) gridTransformBinary(xc, yc, magc, anglec, magnitude_func(), direction_func(), stream); @@ -141,6 +147,69 @@ void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, Outpu syncOutput(angle, _angle, stream); } +void cv::cuda::cartToPolar(InputArray _xy, OutputArray _mag, OutputArray _angle, bool angleInDegrees, Stream& stream) +{ + GpuMat xy = getInputMat(_xy, stream); + + CV_Assert( xy.type() == CV_32FC2 ); + + GpuMat mag = getOutputMat(_mag, xy.size(), CV_32FC1, stream); + GpuMat angle = getOutputMat(_angle, xy.size(), CV_32FC1, stream); + + GpuMat_ magc(mag); + GpuMat_ anglec(angle); + + if (angleInDegrees) + { + auto f1 = magnitude_interleaved_func(); + auto f2 = direction_interleaved_func(); + cv::cudev::tuple f12 = cv::cudev::make_tuple(f1, f2); + gridTransformTuple(globPtr(xy), + tie(magc, anglec), + f12, + stream); + } + else + { + auto f1 = magnitude_interleaved_func(); + auto f2 = direction_interleaved_func(); + cv::cudev::tuple f12 = cv::cudev::make_tuple(f1, f2); + gridTransformTuple(globPtr(xy), + tie(magc, anglec), + f12, + stream); + } + + syncOutput(mag, _mag, stream); + syncOutput(angle, _angle, stream); +} + +void cv::cuda::cartToPolar(InputArray _xy, OutputArray _magAngle, bool angleInDegrees, Stream& stream) +{ + GpuMat xy = getInputMat(_xy, stream); + + CV_Assert( xy.type() == CV_32FC2 ); + + GpuMat magAngle = getOutputMat(_magAngle, xy.size(), CV_32FC2, stream); + + if (angleInDegrees) + { + gridTransformUnary(globPtr(xy), + globPtr(magAngle), + magnitude_direction_interleaved_func(), + stream); + } + else + { + gridTransformUnary(globPtr(xy), + globPtr(magAngle), + magnitude_direction_interleaved_func(), + stream); + } + + syncOutput(magAngle, _magAngle, stream); +} + namespace { template struct sincos_op @@ -159,12 +228,12 @@ namespace }; template - __global__ void polarToCartImpl_(const GlobPtr mag, const GlobPtr angle, GlobPtr xmat, GlobPtr ymat, const T scale, const int rows, const int cols) + __global__ void polarToCartImpl_(const PtrStep mag, const PtrStepSz angle, PtrStep xmat, PtrStep ymat, const T scale) { const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; - if (x >= cols || y >= rows) + if (x >= angle.cols || y >= angle.rows) return; const T mag_val = useMag ? mag(y, x) : static_cast(1.0); @@ -178,23 +247,90 @@ namespace ymat(y, x) = mag_val * sin_a; } + template + __global__ void polarToCartDstInterleavedImpl_(const PtrStep mag, const PtrStepSz angle, PtrStep::type > xymat, const T scale) + { + typedef typename MakeVec::type T2; + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (x >= angle.cols || y >= angle.rows) + return; + + const T mag_val = useMag ? mag(y, x) : static_cast(1.0); + const T angle_val = angle(y, x); + + T sin_a, cos_a; + sincos_op op; + op(scale * angle_val, &sin_a, &cos_a); + + const T2 xy = {mag_val * cos_a, mag_val * sin_a}; + xymat(y, x) = xy; + } + + template + __global__ void polarToCartInterleavedImpl_(const PtrStepSz::type > magAngle, PtrStep::type > xymat, const T scale) + { + typedef typename MakeVec::type T2; + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (x >= magAngle.cols || y >= magAngle.rows) + return; + + const T2 magAngle_val = magAngle(y, x); + const T mag_val = magAngle_val.x; + const T angle_val = magAngle_val.y; + + T sin_a, cos_a; + sincos_op op; + op(scale * angle_val, &sin_a, &cos_a); + + const T2 xy = {mag_val * cos_a, mag_val * sin_a}; + xymat(y, x) = xy; + } + 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(angle.cols, block.x), divUp(angle.rows, block.y)); + + const T scale = angleInDegrees ? static_cast(CV_PI / 180.0) : static_cast(1.0); + + if (mag.empty()) + polarToCartImpl_ << > >(mag, angle, x, y, scale); + else + polarToCartImpl_ << > >(mag, angle, x, y, scale); + } + + template + void polarToCartDstInterleavedImpl(const GpuMat& mag, const GpuMat& angle, GpuMat& xy, bool angleInDegrees, cudaStream_t& stream) + { + typedef typename MakeVec::type T2; const dim3 block(32, 8); - const dim3 grid(divUp(anglec.cols, block.x), divUp(anglec.rows, block.y)); + const dim3 grid(divUp(angle.cols, block.x), divUp(angle.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); + if (mag.empty()) + polarToCartDstInterleavedImpl_ << > >(mag, angle, xy, scale); else - polarToCartImpl_ << > >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols); + polarToCartDstInterleavedImpl_ << > >(mag, angle, xy, scale); + } + + template + void polarToCartInterleavedImpl(const GpuMat& magAngle, GpuMat& xy, bool angleInDegrees, cudaStream_t& stream) + { + typedef typename MakeVec::type T2; + + const dim3 block(32, 8); + const dim3 grid(divUp(magAngle.cols, block.x), divUp(magAngle.rows, block.y)); + + const T scale = angleInDegrees ? static_cast(CV_PI / 180.0) : static_cast(1.0); + + polarToCartInterleavedImpl_ << > >(magAngle, xy, scale); } } @@ -223,4 +359,48 @@ void cv::cuda::polarToCart(InputArray _mag, InputArray _angle, OutputArray _x, O CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); } +void cv::cuda::polarToCart(InputArray _mag, InputArray _angle, OutputArray _xy, bool angleInDegrees, Stream& _stream) +{ + typedef void(*func_t)(const GpuMat& mag, const GpuMat& angle, GpuMat& xy, bool angleInDegrees, cudaStream_t& stream); + static const func_t funcs[7] = { 0, 0, 0, 0, 0, polarToCartDstInterleavedImpl, polarToCartDstInterleavedImpl }; + + GpuMat mag = getInputMat(_mag, _stream); + GpuMat angle = getInputMat(_angle, _stream); + + CV_Assert(angle.depth() == CV_32F || angle.depth() == CV_64F); + CV_Assert( mag.empty() || (mag.type() == angle.type() && mag.size() == angle.size()) ); + + GpuMat xy = getOutputMat(_xy, angle.size(), CV_MAKETYPE(angle.depth(), 2), _stream); + + cudaStream_t stream = StreamAccessor::getStream(_stream); + funcs[angle.depth()](mag, angle, xy, angleInDegrees, stream); + CV_CUDEV_SAFE_CALL( cudaGetLastError() ); + + syncOutput(xy, _xy, _stream); + + if (stream == 0) + CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); +} + +void cv::cuda::polarToCart(InputArray _magAngle, OutputArray _xy, bool angleInDegrees, Stream& _stream) +{ + typedef void(*func_t)(const GpuMat& magAngle, GpuMat& xy, bool angleInDegrees, cudaStream_t& stream); + static const func_t funcs[7] = { 0, 0, 0, 0, 0, polarToCartInterleavedImpl, polarToCartInterleavedImpl }; + + GpuMat magAngle = getInputMat(_magAngle, _stream); + + CV_Assert(magAngle.type() == CV_32FC2 || magAngle.type() == CV_64FC2); + + GpuMat xy = getOutputMat(_xy, magAngle.size(), magAngle.type(), _stream); + + cudaStream_t stream = StreamAccessor::getStream(_stream); + funcs[magAngle.depth()](magAngle, xy, angleInDegrees, stream); + CV_CUDEV_SAFE_CALL( cudaGetLastError() ); + + syncOutput(xy, _xy, _stream); + + if (stream == 0) + CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); +} + #endif diff --git a/modules/cudaarithm/test/test_element_operations.cpp b/modules/cudaarithm/test/test_element_operations.cpp index d2e314b10d9..a15ad7b3ee5 100644 --- a/modules/cudaarithm/test/test_element_operations.cpp +++ b/modules/cudaarithm/test/test_element_operations.cpp @@ -2765,6 +2765,47 @@ INSTANTIATE_TEST_CASE_P(CUDA_Arithm, Phase, testing::Combine( testing::Values(AngleInDegrees(false), AngleInDegrees(true)), WHOLE_SUBMAT)); +PARAM_TEST_CASE(PhaseInterleaved, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, UseRoi) +{ + cv::cuda::DeviceInfo devInfo; + cv::Size size; + bool angleInDegrees; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + angleInDegrees = GET_PARAM(2); + useRoi = GET_PARAM(3); + + cv::cuda::setDevice(devInfo.deviceID()); + } +}; + +CUDA_TEST_P(PhaseInterleaved, Accuracy) +{ + cv::Mat x = randomMat(size, CV_32FC1); + cv::Mat y = randomMat(size, CV_32FC1); + cv::Mat xy; + std::vector xyChannels = {x, y}; + cv::merge(xyChannels, xy); + + cv::cuda::GpuMat dst = createMat(size, CV_32FC1, useRoi); + cv::cuda::phase(loadMat(xy, useRoi), dst, angleInDegrees); + + cv::Mat dst_gold; + cv::phase(x, y, dst_gold, angleInDegrees); + + EXPECT_MAT_NEAR(dst_gold, dst, angleInDegrees ? 1e-2 : 1e-3); +} + +INSTANTIATE_TEST_CASE_P(CUDA_Arithm, PhaseInterleaved, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(AngleInDegrees(false), AngleInDegrees(true)), + WHOLE_SUBMAT)); + //////////////////////////////////////////////////////////////////////////////// // CartToPolar @@ -2809,6 +2850,97 @@ INSTANTIATE_TEST_CASE_P(CUDA_Arithm, CartToPolar, testing::Combine( testing::Values(AngleInDegrees(false), AngleInDegrees(true)), WHOLE_SUBMAT)); +PARAM_TEST_CASE(CartToPolarInterleavedXY, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, UseRoi) +{ + cv::cuda::DeviceInfo devInfo; + cv::Size size; + bool angleInDegrees; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + angleInDegrees = GET_PARAM(2); + useRoi = GET_PARAM(3); + + cv::cuda::setDevice(devInfo.deviceID()); + } +}; + +CUDA_TEST_P(CartToPolarInterleavedXY, Accuracy) +{ + cv::Mat x = randomMat(size, CV_32FC1); + cv::Mat y = randomMat(size, CV_32FC1); + cv::Mat xy; + std::vector xyChannels = {x, y}; + cv::merge(xyChannels, xy); + + cv::cuda::GpuMat mag = createMat(size, CV_32FC1, useRoi); + cv::cuda::GpuMat angle = createMat(size, CV_32FC1, useRoi); + cv::cuda::cartToPolar(loadMat(xy, useRoi), mag, angle, angleInDegrees); + + cv::Mat mag_gold; + cv::Mat angle_gold; + cv::cartToPolar(x, y, mag_gold, angle_gold, angleInDegrees); + + EXPECT_MAT_NEAR(mag_gold, mag, 1e-4); + EXPECT_MAT_NEAR(angle_gold, angle, angleInDegrees ? 1e-2 : 1e-3); +} + +INSTANTIATE_TEST_CASE_P(CUDA_Arithm, CartToPolarInterleavedXY, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(AngleInDegrees(false), AngleInDegrees(true)), + WHOLE_SUBMAT)); + +PARAM_TEST_CASE(CartToPolarInterleavedXYMagAngle, cv::cuda::DeviceInfo, cv::Size, AngleInDegrees, UseRoi) +{ + cv::cuda::DeviceInfo devInfo; + cv::Size size; + bool angleInDegrees; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + angleInDegrees = GET_PARAM(2); + useRoi = GET_PARAM(3); + + cv::cuda::setDevice(devInfo.deviceID()); + } +}; + +CUDA_TEST_P(CartToPolarInterleavedXYMagAngle, Accuracy) +{ + cv::Mat x = randomMat(size, CV_32FC1); + cv::Mat y = randomMat(size, CV_32FC1); + cv::Mat xy; + std::vector xyChannels = {x, y}; + cv::merge(xyChannels, xy); + + cv::cuda::GpuMat magAngle = createMat(size, CV_32FC2, useRoi); + cv::cuda::cartToPolar(loadMat(xy, useRoi), magAngle, angleInDegrees); + std::vector magAngleChannels; + cv::cuda::split(magAngle, magAngleChannels); + cv::cuda::GpuMat& mag = magAngleChannels[0]; + cv::cuda::GpuMat& angle = magAngleChannels[1]; + + cv::Mat mag_gold; + cv::Mat angle_gold; + cv::cartToPolar(x, y, mag_gold, angle_gold, angleInDegrees); + + EXPECT_MAT_NEAR(mag_gold, mag, 1e-4); + EXPECT_MAT_NEAR(angle_gold, angle, angleInDegrees ? 1e-2 : 1e-3); +} + +INSTANTIATE_TEST_CASE_P(CUDA_Arithm, CartToPolarInterleavedXYMagAngle, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(AngleInDegrees(false), AngleInDegrees(true)), + WHOLE_SUBMAT)); + //////////////////////////////////////////////////////////////////////////////// // polarToCart @@ -2857,5 +2989,104 @@ INSTANTIATE_TEST_CASE_P(CUDA_Arithm, PolarToCart, testing::Combine( testing::Values(AngleInDegrees(false), AngleInDegrees(true)), WHOLE_SUBMAT)); +PARAM_TEST_CASE(PolarToCartInterleaveXY, cv::cuda::DeviceInfo, cv::Size, MatType, AngleInDegrees, UseRoi) +{ + cv::cuda::DeviceInfo devInfo; + cv::Size size; + int type; + bool angleInDegrees; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + type = GET_PARAM(2); + angleInDegrees = GET_PARAM(3); + useRoi = GET_PARAM(4); + + cv::cuda::setDevice(devInfo.deviceID()); + } +}; + +CUDA_TEST_P(PolarToCartInterleaveXY, Accuracy) +{ + 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.47); + + cv::cuda::GpuMat xy = createMat(size, CV_MAKETYPE(CV_MAT_DEPTH(type), 2), useRoi); + cv::cuda::polarToCart(loadMat(magnitude, useRoi), loadMat(angle, useRoi), xy, angleInDegrees); + std::vector xyChannels; + cv::cuda::split(xy, xyChannels); + cv::cuda::GpuMat& x = xyChannels[0]; + cv::cuda::GpuMat& y = xyChannels[1]; + + cv::Mat x_gold; + cv::Mat y_gold; + cv::polarToCart(magnitude, angle, x_gold, y_gold, angleInDegrees); + + EXPECT_MAT_NEAR(x_gold, x, tol); + EXPECT_MAT_NEAR(y_gold, y, tol); +} + +INSTANTIATE_TEST_CASE_P(CUDA_Arithm, PolarToCartInterleaveXY, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(CV_32FC1, CV_64FC1), + testing::Values(AngleInDegrees(false), AngleInDegrees(true)), + WHOLE_SUBMAT)); + +PARAM_TEST_CASE(PolarToCartInterleaveMagAngleXY, cv::cuda::DeviceInfo, cv::Size, MatType, AngleInDegrees, UseRoi) +{ + cv::cuda::DeviceInfo devInfo; + cv::Size size; + int type; + bool angleInDegrees; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + type = GET_PARAM(2); + angleInDegrees = GET_PARAM(3); + useRoi = GET_PARAM(4); + + cv::cuda::setDevice(devInfo.deviceID()); + } +}; + +CUDA_TEST_P(PolarToCartInterleaveMagAngleXY, Accuracy) +{ + cv::Mat magnitude = randomMat(size, type); + cv::Mat angle = randomMat(size, type); + std::vector magAngleChannels = {magnitude, angle}; + cv::Mat magAngle; + cv::merge(magAngleChannels, magAngle); + const double tol = (type == CV_32FC1 ? 1.6e-4 : 1e-4) * (angleInDegrees ? 1.0 : 19.47); + + cv::cuda::GpuMat xy = createMat(size, CV_MAKETYPE(CV_MAT_DEPTH(type), 2), useRoi); + cv::cuda::polarToCart(loadMat(magAngle, useRoi), xy, angleInDegrees); + std::vector xyChannels; + cv::cuda::split(xy, xyChannels); + cv::cuda::GpuMat& x = xyChannels[0]; + cv::cuda::GpuMat& y = xyChannels[1]; + + cv::Mat x_gold; + cv::Mat y_gold; + cv::polarToCart(magnitude, angle, x_gold, y_gold, angleInDegrees); + + EXPECT_MAT_NEAR(x_gold, x, tol); + EXPECT_MAT_NEAR(y_gold, y, tol); +} + +INSTANTIATE_TEST_CASE_P(CUDA_Arithm, PolarToCartInterleaveMagAngleXY, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(CV_32FC1, CV_64FC1), + testing::Values(AngleInDegrees(false), AngleInDegrees(true)), + WHOLE_SUBMAT)); + }} // namespace #endif // HAVE_CUDA diff --git a/modules/cudawarping/include/opencv2/cudawarping.hpp b/modules/cudawarping/include/opencv2/cudawarping.hpp index 45cca1ccf86..b9ca957358e 100644 --- a/modules/cudawarping/include/opencv2/cudawarping.hpp +++ b/modules/cudawarping/include/opencv2/cudawarping.hpp @@ -70,6 +70,8 @@ namespace cv { namespace cuda { @param ymap Y values. Only CV_32FC1 type is supported. @param interpolation Interpolation method (see resize ). INTER_NEAREST , INTER_LINEAR and INTER_CUBIC are supported for now. +The extra flag WARP_RELATIVE_MAP can be ORed to the interpolation method +(e.g. INTER_LINEAR | WARP_RELATIVE_MAP) @param borderMode Pixel extrapolation method (see borderInterpolate ). BORDER_REFLECT101 , BORDER_REPLICATE , BORDER_CONSTANT , BORDER_REFLECT and BORDER_WRAP are supported for now. @param borderValue Value used in case of a constant border. By default, it is 0. @@ -79,6 +81,10 @@ The function transforms the source image using the specified map: \f[\texttt{dst} (x,y) = \texttt{src} (xmap(x,y), ymap(x,y))\f] +with the WARP_RELATIVE_MAP flag : + +\f[\texttt{dst} (x,y) = \texttt{src} (x+map_x(x,y),y+map_y(x,y))\f] + Values of pixels with non-integer coordinates are computed using the bilinear interpolation. @sa remap diff --git a/modules/cudawarping/src/cuda/remap.cu b/modules/cudawarping/src/cuda/remap.cu index 38edf19ae24..8f698a2ffb4 100644 --- a/modules/cudawarping/src/cuda/remap.cu +++ b/modules/cudawarping/src/cuda/remap.cu @@ -68,9 +68,23 @@ namespace cv { namespace cuda { namespace device } } + template __global__ void remap_relative(const Ptr2D src, const PtrStepf mapx, const PtrStepf mapy, PtrStepSz dst) + { + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (x < dst.cols && y < dst.rows) + { + const float xcoo = x+mapx.ptr(y)[x]; + const float ycoo = y+mapy.ptr(y)[x]; + + dst.ptr(y)[x] = saturate_cast(src(ycoo, xcoo)); + } + } + template