Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

add interleaved versions of phase/cartToPolar/polarToCart #3607

Merged
merged 13 commits into from
Jul 22, 2024
Merged

add interleaved versions of phase/cartToPolar/polarToCart #3607

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
13 commits
Select commit Hold shift + click to select a range
f19a582
add interleaved versions of phase/cartToPolar/polarToCart
chacha21 Dec 12, 2023
b330b6c
fixed compilation
chacha21 Dec 26, 2023
7e1435b
simplifications as suggested
chacha21 Jan 8, 2024
997927f
Merge branch 'cuda_phase_interleaved' of https://github.com/chacha21/…
chacha21 Jan 8, 2024
0552aed
more simplifications as suggested
chacha21 Jan 8, 2024
094d517
fixed bug
chacha21 Jan 9, 2024
19c772f
modifications as suggested
chacha21 Jan 9, 2024
9b4b9dd
Merge branch '4.x' into cuda_phase_interleaved
chacha21 Jan 9, 2024
63b1f29
Merge remote-tracking branch 'upstream/4.x' into cuda_phase_interleaved
chacha21 May 2, 2024
d4e341e
Merge remote-tracking branch 'upstream/4.x' into cuda_phase_interleaved
chacha21 Jul 18, 2024
65f75dc
disambiguification of tuple
chacha21 Jul 18, 2024
0b65a8b
new attempt to fix compiler error under Ubuntu
chacha21 Jul 18, 2024
cf284c8
new attempt to please both msvc and gcc
chacha21 Jul 19, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
53 changes: 53 additions & 0 deletions modules/cudaarithm/include/opencv2/cudaarithm.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 ).
Expand All @@ -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 ).
Expand All @@ -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
Expand Down
226 changes: 201 additions & 25 deletions modules/cudaarithm/src/cuda/polar_cart.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,11 +66,7 @@ void cv::cuda::magnitude(InputArray _x, InputArray _y, OutputArray _dst, Stream&

GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream);

GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
GpuMat_<float> magc(dst.reshape(1));

gridTransformBinary(xc, yc, magc, magnitude_func<float>(), stream);
gridTransformBinary(globPtr<float>(x), globPtr<float>(y), globPtr<float>(dst), magnitude_func<float>(), stream);

syncOutput(dst, _dst, stream);
}
Expand All @@ -85,11 +81,7 @@ void cv::cuda::magnitudeSqr(InputArray _x, InputArray _y, OutputArray _dst, Stre

GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream);

GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
GpuMat_<float> magc(dst.reshape(1));

gridTransformBinary(xc, yc, magc, magnitude_sqr_func<float>(), stream);
gridTransformBinary(globPtr<float>(x), globPtr<float>(y), globPtr<float>(dst), magnitude_sqr_func<float>(), stream);

syncOutput(dst, _dst, stream);
}
Expand All @@ -104,14 +96,26 @@ void cv::cuda::phase(InputArray _x, InputArray _y, OutputArray _dst, bool angleI

GpuMat dst = getOutputMat(_dst, x.size(), CV_32FC1, stream);

GpuMat_<float> xc(x.reshape(1));
GpuMat_<float> yc(y.reshape(1));
GpuMat_<float> anglec(dst.reshape(1));
if (angleInDegrees)
gridTransformBinary(globPtr<float>(x), globPtr<float>(y), globPtr<float>(dst), direction_func<float, true>(), stream);
else
gridTransformBinary(globPtr<float>(x), globPtr<float>(y), globPtr<float>(dst), direction_func<float, false>(), 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<float, true>(), stream);
gridTransformUnary(globPtr<float2>(xy), globPtr<float>(dst), direction_interleaved_func<float2, true>(), stream);
else
gridTransformBinary(xc, yc, anglec, direction_func<float, false>(), stream);
gridTransformUnary(globPtr<float2>(xy), globPtr<float>(dst), direction_interleaved_func<float2, false>(), stream);

syncOutput(dst, _dst, stream);
}
Expand Down Expand Up @@ -155,6 +159,67 @@ 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_<float> magc(mag.reshape(1));
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can we remove the reshape completely, looking at it again it doesn't do anything? i.e.
GpuMat_ magc(mag);

GpuMat_<float> anglec(angle.reshape(1));

if (angleInDegrees)
{
gridTransformTuple(globPtr<float2>(xy),
tie(magc, anglec),
make_tuple(
magnitude_interleaved_func<float2>(),
direction_interleaved_func<float2, true>()),
stream);
}
else
{
gridTransformTuple(globPtr<float2>(xy),
tie(magc, anglec),
make_tuple(
magnitude_interleaved_func<float2>(),
direction_interleaved_func<float2, false>()),
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<float2>(xy),
globPtr<float2>(magAngle),
magnitude_direction_interleaved_func<float2, true>(),
stream);
}
else
{
gridTransformUnary(globPtr<float2>(xy),
globPtr<float2>(magAngle),
magnitude_direction_interleaved_func<float2, false>(),
stream);
}

syncOutput(magAngle, _magAngle, stream);
}

namespace
{
template <typename T> struct sincos_op
Expand All @@ -173,12 +238,12 @@ namespace
};

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)
__global__ void polarToCartImpl_(const PtrStepSz<T> mag, const PtrStepSz<T> angle, PtrStepSz<T> xmat, PtrStepSz<T> 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<T>(1.0);
Expand All @@ -192,23 +257,90 @@ namespace
ymat(y, x) = mag_val * sin_a;
}

template <typename T, bool useMag>
__global__ void polarToCartDstInterleavedImpl_(const PtrStepSz<T> mag, const PtrStepSz<T> angle, PtrStepSz<typename MakeVec<T, 2>::type > xymat, const T scale)
{
typedef typename MakeVec<T, 2>::type T2;
const int x = blockDim.x * blockIdx.x + threadIdx.x;
const int y = blockDim.y * blockIdx.y + threadIdx.y;

if (x >= xymat.cols || y >= xymat.rows)
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Try to keep the out of range check consistent. I realize mag can be empty but your using xymat, angle and magAngle. Maybe stick with angle and magAngle.

Then you only need to use PtrStepSz for angle/magAngle the other inputs can be PtrStep.

return;

const T mag_val = useMag ? mag(y, x) : static_cast<T>(1.0);
const T angle_val = angle(y, x);

T sin_a, cos_a;
sincos_op<T> 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 <typename T>
__global__ void polarToCartInterleavedImpl_(const PtrStepSz<typename MakeVec<T, 2>::type > magAngle, PtrStepSz<typename MakeVec<T, 2>::type > xymat, const T scale)
{
typedef typename MakeVec<T, 2>::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<T> 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 <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(angle.cols, block.x), divUp(angle.rows, block.y));

const T scale = angleInDegrees ? static_cast<T>(CV_PI / 180.0) : static_cast<T>(1.0);

if (mag.empty())
polarToCartImpl_<T, false> << <grid, block, 0, stream >> >(mag, angle, x, y, scale);
else
polarToCartImpl_<T, true> << <grid, block, 0, stream >> >(mag, angle, x, y, scale);
}

template <typename T>
void polarToCartDstInterleavedImpl(const GpuMat& mag, const GpuMat& angle, GpuMat& xy, bool angleInDegrees, cudaStream_t& stream)
{
typedef typename MakeVec<T, 2>::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<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);
if (mag.empty())
polarToCartDstInterleavedImpl_<T, false> << <grid, block, 0, stream >> >(mag, angle, xy, scale);
else
polarToCartImpl_<T, true> << <grid, block, 0, stream >> >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols);
polarToCartDstInterleavedImpl_<T, true> << <grid, block, 0, stream >> >(mag, angle, xy, scale);
}

template <typename T>
void polarToCartInterleavedImpl(const GpuMat& magAngle, GpuMat& xy, bool angleInDegrees, cudaStream_t& stream)
{
typedef typename MakeVec<T, 2>::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<T>(CV_PI / 180.0) : static_cast<T>(1.0);

polarToCartInterleavedImpl_<T> << <grid, block, 0, stream >> >(magAngle, xy, scale);
}
}

Expand Down Expand Up @@ -237,4 +369,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<float>, polarToCartDstInterleavedImpl<double> };

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<float>, polarToCartInterleavedImpl<double> };

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
Loading