From 7c8c836a7bfc57ff534141e1c26ec391137ad2e1 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 26 Aug 2013 10:34:04 +0400 Subject: [PATCH] switched to new device layer in polar <-> cart --- modules/cudaarithm/src/cuda/polar_cart.cu | 336 ++++++++++----------- modules/cudaarithm/src/element_operations.cpp | 106 ------- .../opencv2/cudev/functional/functional.hpp | 24 ++ 3 files changed, 189 insertions(+), 277 deletions(-) diff --git a/modules/cudaarithm/src/cuda/polar_cart.cu b/modules/cudaarithm/src/cuda/polar_cart.cu index ce2143b..200b79c 100644 --- a/modules/cudaarithm/src/cuda/polar_cart.cu +++ b/modules/cudaarithm/src/cuda/polar_cart.cu @@ -40,178 +40,172 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" +#ifndef HAVE_OPENCV_CUDEV -namespace cv { namespace cuda { namespace device +#error "opencv_cudev is required" + +#else + +#include "opencv2/cudaarithm.hpp" +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void cv::cuda::magnitude(InputArray _x, InputArray _y, OutputArray _dst, Stream& stream) +{ + GpuMat x = _x.getGpuMat(); + GpuMat y = _y.getGpuMat(); + + CV_DbgAssert( x.depth() == CV_32F ); + CV_DbgAssert( y.type() == x.type() && y.size() == x.size() ); + + _dst.create(x.size(), CV_32FC1); + GpuMat dst = _dst.getGpuMat(); + + GpuMat_ xc(x.reshape(1)); + GpuMat_ yc(y.reshape(1)); + GpuMat_ magc(dst.reshape(1)); + + gridTransformBinary(xc, yc, magc, magnitude_func(), stream); +} + +void cv::cuda::magnitudeSqr(InputArray _x, InputArray _y, OutputArray _dst, Stream& stream) +{ + GpuMat x = _x.getGpuMat(); + GpuMat y = _y.getGpuMat(); + + CV_DbgAssert( x.depth() == CV_32F ); + CV_DbgAssert( y.type() == x.type() && y.size() == x.size() ); + + _dst.create(x.size(), CV_32FC1); + GpuMat dst = _dst.getGpuMat(); + + GpuMat_ xc(x.reshape(1)); + GpuMat_ yc(y.reshape(1)); + GpuMat_ magc(dst.reshape(1)); + + gridTransformBinary(xc, yc, magc, magnitude_sqr_func(), stream); +} + +void cv::cuda::phase(InputArray _x, InputArray _y, OutputArray _dst, bool angleInDegrees, Stream& stream) +{ + GpuMat x = _x.getGpuMat(); + GpuMat y = _y.getGpuMat(); + + CV_DbgAssert( x.depth() == CV_32F ); + CV_DbgAssert( y.type() == x.type() && y.size() == x.size() ); + + _dst.create(x.size(), CV_32FC1); + GpuMat dst = _dst.getGpuMat(); + + GpuMat_ xc(x.reshape(1)); + GpuMat_ yc(y.reshape(1)); + GpuMat_ anglec(dst.reshape(1)); + + if (angleInDegrees) + gridTransformBinary(xc, yc, anglec, direction_func(), stream); + else + gridTransformBinary(xc, yc, anglec, direction_func(), stream); +} + +void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, OutputArray _angle, bool angleInDegrees, Stream& stream) { - namespace mathfunc + GpuMat x = _x.getGpuMat(); + GpuMat y = _y.getGpuMat(); + + CV_DbgAssert( x.depth() == CV_32F ); + CV_DbgAssert( y.type() == x.type() && y.size() == x.size() ); + + _mag.create(x.size(), CV_32FC1); + GpuMat mag = _mag.getGpuMat(); + + _angle.create(x.size(), CV_32FC1); + GpuMat angle = _angle.getGpuMat(); + + GpuMat_ xc(x.reshape(1)); + GpuMat_ yc(y.reshape(1)); + GpuMat_ magc(mag.reshape(1)); + GpuMat_ anglec(angle.reshape(1)); + + if (angleInDegrees) + { + gridTransformTuple(zipPtr(xc, yc), + tie(magc, anglec), + make_tuple( + binaryTupleAdapter<0, 1>(magnitude_func()), + binaryTupleAdapter<0, 1>(direction_func())), + stream); + } + else + { + gridTransformTuple(zipPtr(xc, yc), + tie(magc, anglec), + make_tuple( + binaryTupleAdapter<0, 1>(magnitude_func()), + binaryTupleAdapter<0, 1>(direction_func())), + stream); + } +} + +namespace +{ + template + __global__ void polarToCartImpl(const GlobPtr mag, const GlobPtr angle, GlobPtr xmat, GlobPtr ymat, const float scale, const int rows, const int cols) { - ////////////////////////////////////////////////////////////////////////////////////// - // Cart <-> Polar - - struct Nothing - { - static __device__ __forceinline__ void calc(int, int, float, float, float*, size_t, float) - { - } - }; - struct Magnitude - { - static __device__ __forceinline__ 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__ __forceinline__ 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__ __forceinline__ void calc(int x, int y, float x_data, float y_data, float* dst, size_t dst_step, float scale) - { - float angle = ::atan2f(y_data, x_data); - angle += (angle < 0) * 2.0f * CV_PI_F; - dst[y * dst_step + x] = scale * angle; - } - }; - 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__ __forceinline__ float get(const float* mag, size_t mag_step, int x, int y) - { - return mag[y * mag_step + x]; - } - }; - struct EmptyMag - { - static __device__ __forceinline__ 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; - } - } - - template - void cartToPolar_caller(PtrStepSzf x, PtrStepSzf y, PtrStepSzf mag, PtrStepSzf angle, bool angleInDegrees, cudaStream_t stream) - { - dim3 threads(32, 8, 1); - dim3 grid(1, 1, 1); - - grid.x = divUp(x.cols, threads.x); - grid.y = divUp(x.rows, threads.y); - - const float scale = angleInDegrees ? (180.0f / CV_PI_F) : 1.f; - - cartToPolar<<>>( - x.data, x.step/x.elemSize(), y.data, y.step/y.elemSize(), - mag.data, mag.step/mag.elemSize(), angle.data, angle.step/angle.elemSize(), scale, x.cols, x.rows); - cudaSafeCall( cudaGetLastError() ); - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - - void cartToPolar_gpu(PtrStepSzf x, PtrStepSzf y, PtrStepSzf mag, bool magSqr, PtrStepSzf angle, bool angleInDegrees, cudaStream_t stream) - { - typedef void (*caller_t)(PtrStepSzf x, PtrStepSzf y, PtrStepSzf mag, PtrStepSzf 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.data == 0][magSqr][angle.data == 0](x, y, mag, angle, angleInDegrees, stream); - } - - template - void polarToCart_caller(PtrStepSzf mag, PtrStepSzf angle, PtrStepSzf x, PtrStepSzf y, bool angleInDegrees, cudaStream_t stream) - { - dim3 threads(32, 8, 1); - dim3 grid(1, 1, 1); - - grid.x = divUp(mag.cols, threads.x); - grid.y = divUp(mag.rows, threads.y); - - const float scale = angleInDegrees ? (CV_PI_F / 180.0f) : 1.0f; - - polarToCart<<>>(mag.data, mag.step/mag.elemSize(), - angle.data, angle.step/angle.elemSize(), scale, x.data, x.step/x.elemSize(), y.data, y.step/y.elemSize(), mag.cols, mag.rows); - cudaSafeCall( cudaGetLastError() ); - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - - void polarToCart_gpu(PtrStepSzf mag, PtrStepSzf angle, PtrStepSzf x, PtrStepSzf y, bool angleInDegrees, cudaStream_t stream) - { - typedef void (*caller_t)(PtrStepSzf mag, PtrStepSzf angle, PtrStepSzf x, PtrStepSzf y, bool angleInDegrees, cudaStream_t stream); - static const caller_t callers[2] = - { - polarToCart_caller, - polarToCart_caller - }; - - callers[mag.data == 0](mag, angle, x, y, angleInDegrees, stream); - } - } // namespace mathfunc -}}} // namespace cv { namespace cuda { namespace cudev - -#endif /* CUDA_DISABLER */ + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (x >= cols || y >= rows) + return; + + const float mag_val = useMag ? mag(y, x) : 1.0f; + const float angle_val = angle(y, x); + + float sin_a, cos_a; + ::sincosf(scale * angle_val, &sin_a, &cos_a); + + xmat(y, x) = mag_val * cos_a; + ymat(y, x) = mag_val * sin_a; + } +} + +void cv::cuda::polarToCart(InputArray _mag, InputArray _angle, OutputArray _x, OutputArray _y, bool angleInDegrees, Stream& _stream) +{ + GpuMat mag = _mag.getGpuMat(); + GpuMat angle = _angle.getGpuMat(); + + CV_DbgAssert( angle.depth() == CV_32F ); + CV_DbgAssert( mag.empty() || (mag.type() == angle.type() && mag.size() == angle.size()) ); + + _x.create(angle.size(), CV_32FC1); + GpuMat x = _x.getGpuMat(); + + _y.create(angle.size(), CV_32FC1); + GpuMat y = _y.getGpuMat(); + + 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; + + 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); + + CV_CUDEV_SAFE_CALL( cudaGetLastError() ); + + if (stream == 0) + CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); +} + +#endif diff --git a/modules/cudaarithm/src/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index a188cc9..795d7ff 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -493,110 +493,4 @@ void cv::cuda::magnitudeSqr(InputArray _src, OutputArray _dst, Stream& stream) npp_magnitude(src, dst, nppiMagnitudeSqr_32fc32f_C1R, StreamAccessor::getStream(stream)); } -//////////////////////////////////////////////////////////////////////// -// Polar <-> Cart - -namespace cv { namespace cuda { namespace device -{ - namespace mathfunc - { - void cartToPolar_gpu(PtrStepSzf x, PtrStepSzf y, PtrStepSzf mag, bool magSqr, PtrStepSzf angle, bool angleInDegrees, cudaStream_t stream); - void polarToCart_gpu(PtrStepSzf mag, PtrStepSzf angle, PtrStepSzf x, PtrStepSzf y, bool angleInDegrees, cudaStream_t stream); - } -}}} - -namespace -{ - void cartToPolar_caller(const GpuMat& x, const GpuMat& y, GpuMat* mag, bool magSqr, GpuMat* angle, bool angleInDegrees, cudaStream_t stream) - { - using namespace ::cv::cuda::device::mathfunc; - - CV_Assert(x.size() == y.size() && x.type() == y.type()); - CV_Assert(x.depth() == CV_32F); - - GpuMat x1cn = x.reshape(1); - GpuMat y1cn = y.reshape(1); - GpuMat mag1cn = mag ? mag->reshape(1) : GpuMat(); - GpuMat angle1cn = angle ? angle->reshape(1) : GpuMat(); - - cartToPolar_gpu(x1cn, y1cn, mag1cn, magSqr, angle1cn, angleInDegrees, stream); - } - - void polarToCart_caller(const GpuMat& mag, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, cudaStream_t stream) - { - using namespace ::cv::cuda::device::mathfunc; - - CV_Assert((mag.empty() || mag.size() == angle.size()) && mag.type() == angle.type()); - CV_Assert(mag.depth() == CV_32F); - - GpuMat mag1cn = mag.reshape(1); - GpuMat angle1cn = angle.reshape(1); - GpuMat x1cn = x.reshape(1); - GpuMat y1cn = y.reshape(1); - - polarToCart_gpu(mag1cn, angle1cn, x1cn, y1cn, angleInDegrees, stream); - } -} - -void cv::cuda::magnitude(InputArray _x, InputArray _y, OutputArray _dst, Stream& stream) -{ - GpuMat x = _x.getGpuMat(); - GpuMat y = _y.getGpuMat(); - - _dst.create(x.size(), CV_32FC1); - GpuMat dst = _dst.getGpuMat(); - - cartToPolar_caller(x, y, &dst, false, 0, false, StreamAccessor::getStream(stream)); -} - -void cv::cuda::magnitudeSqr(InputArray _x, InputArray _y, OutputArray _dst, Stream& stream) -{ - GpuMat x = _x.getGpuMat(); - GpuMat y = _y.getGpuMat(); - - _dst.create(x.size(), CV_32FC1); - GpuMat dst = _dst.getGpuMat(); - - cartToPolar_caller(x, y, &dst, true, 0, false, StreamAccessor::getStream(stream)); -} - -void cv::cuda::phase(InputArray _x, InputArray _y, OutputArray _dst, bool angleInDegrees, Stream& stream) -{ - GpuMat x = _x.getGpuMat(); - GpuMat y = _y.getGpuMat(); - - _dst.create(x.size(), CV_32FC1); - GpuMat dst = _dst.getGpuMat(); - - cartToPolar_caller(x, y, 0, false, &dst, angleInDegrees, StreamAccessor::getStream(stream)); -} - -void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, OutputArray _angle, bool angleInDegrees, Stream& stream) -{ - GpuMat x = _x.getGpuMat(); - GpuMat y = _y.getGpuMat(); - - _mag.create(x.size(), CV_32FC1); - GpuMat mag = _mag.getGpuMat(); - - _angle.create(x.size(), CV_32FC1); - GpuMat angle = _angle.getGpuMat(); - - cartToPolar_caller(x, y, &mag, false, &angle, angleInDegrees, StreamAccessor::getStream(stream)); -} - -void cv::cuda::polarToCart(InputArray _mag, InputArray _angle, OutputArray _x, OutputArray _y, bool angleInDegrees, Stream& stream) -{ - GpuMat mag = _mag.getGpuMat(); - GpuMat angle = _angle.getGpuMat(); - - _x.create(mag.size(), CV_32FC1); - GpuMat x = _x.getGpuMat(); - - _y.create(mag.size(), CV_32FC1); - GpuMat y = _y.getGpuMat(); - - polarToCart_caller(mag, angle, x, y, angleInDegrees, StreamAccessor::getStream(stream)); -} - #endif diff --git a/modules/cudev/include/opencv2/cudev/functional/functional.hpp b/modules/cudev/include/opencv2/cudev/functional/functional.hpp index 3ac5328..7934f78 100644 --- a/modules/cudev/include/opencv2/cudev/functional/functional.hpp +++ b/modules/cudev/include/opencv2/cudev/functional/functional.hpp @@ -616,6 +616,30 @@ template struct magnitude_func : binary_function struct magnitude_sqr_func : binary_function::type> +{ + __device__ __forceinline__ typename functional_detail::FloatType::type operator ()(typename TypeTraits::parameter_type a, typename TypeTraits::parameter_type b) const + { + return a * a + b * b; + } +}; + +template struct direction_func : binary_function +{ + __device__ T operator ()(T x, T y) const + { + atan2_func f; + typename atan2_func::result_type angle = f(y, x); + + angle += (angle < 0) * (2.0f * CV_PI_F); + + if (angleInDegrees) + angle *= (180.0f / CV_PI_F); + + return saturate_cast(angle); + } +}; + template struct pow_func : binary_function { __device__ __forceinline__ float operator ()(T val, float power) const -- 2.7.4