Skip to content

add interleaved versions of phase/cartToPolar/polarToCart #3607

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

Merged
asmorkalov 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 2 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
208 changes: 208 additions & 0 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 @@ -116,6 +116,25 @@ void cv::cuda::phase(InputArray _x, InputArray _y, OutputArray _dst, bool angleI
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);

GpuMat_<float2> xyc(xy.reshape(2));
Copy link
Contributor

@cudawarped cudawarped Jan 1, 2024

Choose a reason for hiding this comment

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

I know the existing functions use reshape to convert a GpuMat to a GpuMat_ before being passed to gridTransformxxx but I find this confusing because nothing is being reshaped, would it not be better to use globPtr<> directly on the GpuMats. e.g.

    if (angleInDegrees)
        gridTransformUnary(globPtr<float2>(xy), globPtr<float>(dst), direction_interleaved_func<float2, true>(), stream);

If so the existing routines could be updated to remove the bloat.

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

if (angleInDegrees)
gridTransformUnary(xyc, anglec, direction_interleaved_func<float2, true>(), stream);
else
gridTransformUnary(xyc, anglec, direction_interleaved_func<float2, false>(), stream);

syncOutput(dst, _dst, stream);
}

void cv::cuda::cartToPolar(InputArray _x, InputArray _y, OutputArray _mag, OutputArray _angle, bool angleInDegrees, Stream& stream)
{
GpuMat x = getInputMat(_x, stream);
Expand Down Expand Up @@ -155,6 +174,71 @@ 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_<float2> xyc(xy.reshape(2));
Copy link
Contributor

@cudawarped cudawarped Jan 1, 2024

Choose a reason for hiding this comment

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

Same here.

GpuMat_<float> magc(mag.reshape(1));
Copy link
Contributor

@cudawarped cudawarped Jan 9, 2024

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(xyc,
tie(magc, anglec),
make_tuple(
magnitude_interleaved_func<float2>(),
direction_interleaved_func<float2, true>()),
stream);
}
else
{
gridTransformTuple(xyc,
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);

GpuMat_<float2> xyc(xy.reshape(2));
Copy link
Contributor

@cudawarped cudawarped Jan 1, 2024

Choose a reason for hiding this comment

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

And here.

GpuMat_<float2> magAnglec(magAngle.reshape(2));

if (angleInDegrees)
{
gridTransformUnary(xyc,
magAnglec,
magnitude_direction_interleaved_func<float2, true>(),
stream);
}
else
{
gridTransformUnary(xyc,
magAnglec,
magnitude_direction_interleaved_func<float2, false>(),
stream);
}

syncOutput(magAngle, _magAngle, stream);
}

namespace
{
template <typename T> struct sincos_op
Expand Down Expand Up @@ -192,6 +276,49 @@ namespace
ymat(y, x) = mag_val * sin_a;
}

template <typename T, bool useMag>
__global__ void polarToCartDstInterleavedImpl_(const GlobPtr<T> mag, const GlobPtr<T> angle, GlobPtr<typename MakeVec<T, 2>::type > xymat, const T scale, const int rows, const int cols)
Copy link
Contributor

@cudawarped cudawarped Jan 1, 2024

Choose a reason for hiding this comment

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

If you use PtrStep<T> for mag, angle and xymat then you pass them directly to this function instead of using the intermediate GpuMat_<T> with reshape. Additionaly if angle is a PtrStepSz<T> then you don't need to pass cols and rows seperately. e.g.

    __global__ void polarToCartDstInterleavedImpl_(const PtrStep<T> mag, const PtrStepSz<T> angle, PtrStep<typename MakeVec<T, 2>::type > xymat, const T scale, const int rows, const int cols)
    {
        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 >= angle.cols || y >= angle.rows)
            return;

You can also make this adjustment to all the other polarToCart kernel calls including the existing one.

{
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 >= cols || y >= rows)
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, bool useMag>
__global__ void polarToCartInterleavedImpl_(const GlobPtr<typename MakeVec<T, 2>::type > magAngle, GlobPtr<typename MakeVec<T, 2>::type > xymat, const T scale, const int rows, const int cols)
{
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 >= cols || y >= rows)
return;

const T2 magAngle_val = magAngle(y, x);
const T mag_val = useMag ? magAngle_val.x : static_cast<T>(1.0);
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)
{
Expand All @@ -210,6 +337,43 @@ namespace
else
polarToCartImpl_<T, true> << <grid, block, 0, stream >> >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xc), shrinkPtr(yc), scale, anglec.rows, anglec.cols);
}

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;
GpuMat_<T2> xyc(xy.reshape(2));
Copy link
Contributor

@cudawarped cudawarped Jan 1, 2024

Choose a reason for hiding this comment

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

If you switch to PtrStep and PtrStepSz inside polarToCartDstInterleavedImpl_ then this can be simplified to

    template <typename T>
    void polarToCartDstInterleavedImpl(const GpuMat& mag, const GpuMat& angle, GpuMat& xy, bool angleInDegrees, cudaStream_t& stream)
    {
        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())
            polarToCartDstInterleavedImpl_<T, false> << <grid, block, 0, stream >> >(mag, angle, xy, scale, angle.rows, angle.cols);
        else
            polarToCartDstInterleavedImpl_<T, true> << <grid, block, 0, stream >> >(mag, angle, xy, scale, angle.rows, angle.cols);
    }

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())
polarToCartDstInterleavedImpl_<T, false> << <grid, block, 0, stream >> >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xyc), scale, anglec.rows, anglec.cols);
else
polarToCartDstInterleavedImpl_<T, true> << <grid, block, 0, stream >> >(shrinkPtr(magc), shrinkPtr(anglec), shrinkPtr(xyc), scale, anglec.rows, anglec.cols);
}

template <typename T>
void polarToCartInterleavedImpl(const GpuMat& magAngle, GpuMat& xy, bool angleInDegrees, cudaStream_t& stream)
{
typedef typename MakeVec<T, 2>::type T2;
GpuMat_<T2> xyc(xy.reshape(2));
GpuMat_<T2> magAnglec(magAngle.reshape(2));

const dim3 block(32, 8);
const dim3 grid(divUp(magAnglec.cols, block.x), divUp(magAnglec.rows, block.y));

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

if (magAnglec.empty())
polarToCartInterleavedImpl_<T, false> << <grid, block, 0, stream >> >(shrinkPtr(magAnglec), shrinkPtr(xyc), scale, magAnglec.rows, magAnglec.cols);
else
polarToCartInterleavedImpl_<T, true> << <grid, block, 0, stream >> >(shrinkPtr(magAnglec), shrinkPtr(xyc), scale, magAnglec.rows, magAnglec.cols);
}
}

void cv::cuda::polarToCart(InputArray _mag, InputArray _angle, OutputArray _x, OutputArray _y, bool angleInDegrees, Stream& _stream)
Expand Down Expand Up @@ -237,4 +401,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