From 632ef5a2063ccf7bf8112a293e0c149687682e2c Mon Sep 17 00:00:00 2001 From: Jon Beniston Date: Sat, 8 Nov 2025 11:25:10 +0000 Subject: [PATCH] Add cv::cuda::cvtColorTwoPlane --- .../include/opencv2/cudaimgproc.hpp | 18 ++++ modules/cudaimgproc/src/color.cpp | 83 +++++++++++++++++++ modules/cudaimgproc/test/test_color.cpp | 61 ++++++++++++++ 3 files changed, 162 insertions(+) diff --git a/modules/cudaimgproc/include/opencv2/cudaimgproc.hpp b/modules/cudaimgproc/include/opencv2/cudaimgproc.hpp index 01e7c41ca9a..1f8eb900f08 100644 --- a/modules/cudaimgproc/include/opencv2/cudaimgproc.hpp +++ b/modules/cudaimgproc/include/opencv2/cudaimgproc.hpp @@ -90,6 +90,24 @@ performance. */ CV_EXPORTS_W void cvtColor(InputArray src, OutputArray dst, int code, int dcn = 0, Stream& stream = Stream::Null()); +/** @brief Converts an image from one color space to another where the source image is stored in two planes. + +This function currently only supports YUV420 in NV12 format to RGB/BGR conversion. The conversion used is: +- R=1.163999557*(Y-16) + 1.5959997177(V-128) +- G=1.163999557*(Y-16) -0.812999725(V-128) -0.390999794(U-128) +- B=1.163999557*(Y-16) + 2.017999649(U-128) + +@param src1 8-bit image (CV_8U) of the Y plane. +@param src2 image containing interleaved (CV_8UC2) U/V plane. +@param dst Destination image. +@param code Color space conversion code. It can take any of the following values: +- #COLOR_YUV2BGR_NV12 +- #COLOR_YUV2RGB_NV12 +@param stream Stream for the asynchronous version. + + */ +CV_EXPORTS_W void cvtColorTwoPlane(InputArray src1, InputArray src2, OutputArray dst, int code, Stream& stream = Stream::Null()); + enum DemosaicTypes { //! Bayer Demosaicing (Malvar, He, and Cutler) diff --git a/modules/cudaimgproc/src/color.cpp b/modules/cudaimgproc/src/color.cpp index 94ffe90fa2f..1c0597cad84 100644 --- a/modules/cudaimgproc/src/color.cpp +++ b/modules/cudaimgproc/src/color.cpp @@ -49,6 +49,8 @@ using namespace cv::cuda; void cv::cuda::cvtColor(InputArray, OutputArray, int, int, Stream&) { throw_no_cuda(); } +void cv::cuda::cvtColorTwoPlane(InputArray, InputArray, OutputArray, int, Stream&) { throw_no_cuda(); } + void cv::cuda::demosaicing(InputArray, OutputArray, int, int, Stream&) { throw_no_cuda(); } void cv::cuda::swapChannels(InputOutputArray, const int[4], Stream&) { throw_no_cuda(); } @@ -2112,6 +2114,87 @@ void cv::cuda::cvtColor(InputArray src, OutputArray dst, int code, int dcn, Stre func(src, dst, dcn, stream); } +//////////////////////////////////////////////////////////////////////// +// cvtColorTwoPlane + +#if NPP_VER_MAJOR <= 12 + +void cv::cuda::cvtColorTwoPlane(InputArray _src1, InputArray _src2, OutputArray _dst, int code, Stream& _stream) +{ + // See: https://forums.developer.nvidia.com/t/incorrect-offset-application-in-npp-colortwist-conversion/312180 + CV_Error(Error::StsBadFlag, "cvtColorTwoPlane not supported with NPP <= 12 as nppiNV12ToRGB_8u_ColorTwist32f_P2C3R_Ctx is buggy"); +} + +#else + +void cv::cuda::cvtColorTwoPlane(InputArray _src1, InputArray _src2, OutputArray _dst, int code, Stream& _stream) +{ + // Coeffs from opencl/color_yuv.cl c_YUV2RGBCoeffs_420 + static const Npp32f twistYUV2BGR[3][4] = { + {1.163999557f, 2.017999649f, 0.0f, -16.0f}, // B + {1.163999557f, -0.390999794f, -0.812999725f, -128.0f}, // G + {1.163999557f, 0.0f, 1.5959997177f, -128.0f} // R + }; + + static const Npp32f twistYUV2RGB[3][4] = { + {1.163999557f, 0.0f, 1.5959997177f, -16.0f}, // R + {1.163999557f, -0.390999794f, -0.812999725f, -128.0f}, // G + {1.163999557f, 2.017999649f, 0.0f, -128.0f}, // B + }; + + CV_Assert( !_src1.empty() ); + CV_Assert( !_src2.empty() ); + + const Npp32f (*twist)[3][4]; + + switch (code) + { + case COLOR_YUV2BGR_NV12: + twist = &twistYUV2BGR; + break; + case COLOR_YUV2RGB_NV12: + twist = &twistYUV2RGB; + break; + default: + CV_Error( cv::Error::StsBadFlag, "Unknown/unsupported color conversion code" ); + return; + } + + GpuMat src1 = _src1.getGpuMat(); + const int depth1 = _src1.depth(); + GpuMat src2 = _src2.getGpuMat(); + const int depth2 = _src2.depth(); + + CV_Assert( depth1 == CV_8U ); + CV_Assert( src1.channels() == 1 ); + CV_Assert( depth2 == CV_8U ); + CV_Assert( src2.channels() == 2 ); + + _dst.create(_src1.size(), CV_MAKE_TYPE(depth1, 3)); + GpuMat dst = _dst.getGpuMat(); + + cudaStream_t stream = StreamAccessor::getStream(_stream); + NppStreamHandler h(stream); + + NppiSize sz; + sz.width = src1.cols; + sz.height = src1.rows; + + const Npp8u *const pSrc[2] = {src1.ptr(), src2.ptr()}; + int aSrcStep[2] = {static_cast(src1.step), static_cast(src2.step)}; + +#if USE_NPP_STREAM_CTX + nppSafeCall( nppiNV12ToRGB_8u_ColorTwist32f_P2C3R_Ctx(pSrc, aSrcStep, dst.ptr(), static_cast(dst.step), sz, *twist, h) ); +#else + nppSafeCall( nppiNV12ToRGB_8u_ColorTwist32f_P2C3R_Ctx(pSrc, aSrcStep, dst.ptr(), static_cast(dst.step), sz, *twist) ); +#endif + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); +} + +#endif + //////////////////////////////////////////////////////////////////////// // demosaicing diff --git a/modules/cudaimgproc/test/test_color.cpp b/modules/cudaimgproc/test/test_color.cpp index e88bc8a1db6..a627ffba535 100644 --- a/modules/cudaimgproc/test/test_color.cpp +++ b/modules/cudaimgproc/test/test_color.cpp @@ -2291,6 +2291,67 @@ INSTANTIATE_TEST_CASE_P(CUDA_ImgProc, CvtColor, testing::Combine( testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)), WHOLE_SUBMAT)); +/////////////////////////////////////////////////////////////////////////////////////////////////////// +// cvtColorTwoPlane + +PARAM_TEST_CASE(CvtColorTwoPlane, cv::cuda::DeviceInfo, cv::Size, MatDepth, UseRoi) +{ + cv::cuda::DeviceInfo devInfo; + cv::Size size; + int depth; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + depth = GET_PARAM(2); + useRoi = GET_PARAM(3); + + cv::cuda::setDevice(devInfo.deviceID()); + } +}; + +CUDA_TEST_P(CvtColorTwoPlane, YUV2BGR_NV12) +{ + if ((depth != CV_8U) || useRoi) + return; + + cv::Mat src1 = randomMat(size, depth, 16.0, 235.0); + cv::Mat src2 = randomMat(size/2, CV_MAKE_TYPE(depth, 2), 16.0, 240.0); + + cv::cuda::GpuMat dst; + cv::cuda::cvtColorTwoPlane(loadMat(src1, useRoi), loadMat(src2, useRoi), dst, cv::COLOR_YUV2BGR_NV12); + + cv::Mat dst_gold; + cv::cvtColorTwoPlane(src1, src2, dst_gold, cv::COLOR_YUV2BGR_NV12); + + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 2); +} + +CUDA_TEST_P(CvtColorTwoPlane, YUV2RGB_NV12) +{ + if ((depth != CV_8U) || useRoi) + return; + + cv::Mat src1 = randomMat(size, depth, 16.0, 235.0); + cv::Mat src2 = randomMat(size/2, CV_MAKE_TYPE(depth, 2), 16.0, 240.0); + + cv::cuda::GpuMat dst; + cv::cuda::cvtColorTwoPlane(loadMat(src1, useRoi), loadMat(src2, useRoi), dst, cv::COLOR_YUV2RGB_NV12); + + cv::Mat dst_gold; + cv::cvtColorTwoPlane(src1, src2, dst_gold, cv::COLOR_YUV2RGB_NV12); + + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 2); +} + +INSTANTIATE_TEST_CASE_P(CUDA_ImgProc, CvtColorTwoPlane, testing::Combine( + ALL_DEVICES, + testing::Values(cv::Size(128, 128)), + testing::Values(MatDepth(CV_8U)), + WHOLE_SUBMAT)); + /////////////////////////////////////////////////////////////////////////////////////////////////////// // Demosaicing