From c897bc93e0e475b21456b411b4bfcc180a90c3e5 Mon Sep 17 00:00:00 2001 From: cudawarped <12133430+cudawarped@users.noreply.github.com> Date: Sun, 8 May 2022 12:15:04 +0100 Subject: [PATCH 1/3] Force decoding of all supported YUV inputs to NV12 and log warning to indicate this is taking place. Add YUV output. --- .../cudacodec/include/opencv2/cudacodec.hpp | 9 ++++-- modules/cudacodec/src/video_decoder.cpp | 31 ++++++++++++++++--- modules/cudacodec/src/video_reader.cpp | 3 ++ modules/cudacodec/test/test_video.cpp | 4 ++- 4 files changed, 40 insertions(+), 7 deletions(-) diff --git a/modules/cudacodec/include/opencv2/cudacodec.hpp b/modules/cudacodec/include/opencv2/cudacodec.hpp index 552010c8bba..24f939af4b3 100644 --- a/modules/cudacodec/include/opencv2/cudacodec.hpp +++ b/modules/cudacodec/include/opencv2/cudacodec.hpp @@ -326,12 +326,13 @@ enum class VideoReaderProps { #endif }; -/** @brief ColorFormat for the frame returned by the decoder. +/** @brief ColorFormat for the frame returned by nextFrame()/retrieve(). */ enum class ColorFormat { BGRA = 1, BGR = 2, GRAY = 3, + YUV = 4, #ifndef CV_DOXYGEN PROP_NOT_SUPPORTED #endif @@ -394,7 +395,11 @@ class CV_EXPORTS_W VideoReader */ CV_WRAP virtual bool set(const VideoReaderProps propertyId, const double propertyVal) = 0; - CV_WRAP virtual void set(const ColorFormat _colorFormat) = 0; + /** @brief Set the desired ColorFormat for the frame returned by nextFrame()/retrieve(). + + @param colorFormat Value of the ColorFormat. + */ + CV_WRAP virtual void set(const ColorFormat colorFormat) = 0; /** @brief Returns the specified VideoReader property diff --git a/modules/cudacodec/src/video_decoder.cpp b/modules/cudacodec/src/video_decoder.cpp index e5e9353e4fa..c05d0fd7305 100644 --- a/modules/cudacodec/src/video_decoder.cpp +++ b/modules/cudacodec/src/video_decoder.cpp @@ -45,14 +45,36 @@ #ifdef HAVE_NVCUVID +static const char* GetVideoChromaFormatString(cudaVideoChromaFormat eChromaFormat) { + static struct { + cudaVideoChromaFormat eChromaFormat; + const char* name; + } aChromaFormatName[] = { + { cudaVideoChromaFormat_Monochrome, "YUV 400 (Monochrome)" }, + { cudaVideoChromaFormat_420, "YUV 420" }, + { cudaVideoChromaFormat_422, "YUV 422" }, + { cudaVideoChromaFormat_444, "YUV 444" }, + }; + + if (eChromaFormat >= 0 && eChromaFormat < sizeof(aChromaFormatName) / sizeof(aChromaFormatName[0])) { + return aChromaFormatName[eChromaFormat].name; + } + return "Unknown"; +} + void cv::cudacodec::detail::VideoDecoder::create(const FormatInfo& videoFormat) { - if (videoFormat.nBitDepthMinus8 > 0 || videoFormat.chromaFormat == YUV444) - CV_Error(Error::StsUnsupportedFormat, "NV12 output currently supported for 8 bit YUV420, YUV422 and Monochrome inputs."); - videoFormat_ = videoFormat; const cudaVideoCodec _codec = static_cast(videoFormat.codec); const cudaVideoChromaFormat _chromaFormat = static_cast(videoFormat.chromaFormat); + if (videoFormat.nBitDepthMinus8 > 0) { + std::ostringstream warning; + warning << "NV12 (8 bit luma, 4 bit chroma) is currently the only supported decoder output format. Video input is " << videoFormat.nBitDepthMinus8 + 8 << " bit " \ + << std::string(GetVideoChromaFormatString(_chromaFormat)) << ". Truncating luma to 8 bits"; + if (videoFormat.chromaFormat != YUV420) + warning << " and chroma to 4 bits"; + CV_LOG_WARNING(NULL, warning.str()); + } const cudaVideoCreateFlags videoCreateFlags = (_codec == cudaVideoCodec_JPEG || _codec == cudaVideoCodec_MPEG2) ? cudaVideoCreate_PreferCUDA : cudaVideoCreate_PreferCUVID; @@ -98,7 +120,7 @@ void cv::cudacodec::detail::VideoDecoder::create(const FormatInfo& videoFormat) cuSafeCall(cuCtxPushCurrent(ctx_)); cuSafeCall(cuvidGetDecoderCaps(&decodeCaps)); cuSafeCall(cuCtxPopCurrent(NULL)); - if (!decodeCaps.bIsSupported) + if (!(decodeCaps.bIsSupported && (decodeCaps.nOutputFormatMask & (1 << cudaVideoSurfaceFormat_NV12)))) CV_Error(Error::StsUnsupportedFormat, "Video source is not supported by hardware video decoder"); CV_Assert(videoFormat.ulWidth >= decodeCaps.nMinWidth && @@ -115,6 +137,7 @@ void cv::cudacodec::detail::VideoDecoder::create(const FormatInfo& videoFormat) createInfo_.ulHeight = videoFormat.ulHeight; createInfo_.ulNumDecodeSurfaces = videoFormat.ulNumDecodeSurfaces; createInfo_.ChromaFormat = _chromaFormat; + createInfo_.bitDepthMinus8 = videoFormat.nBitDepthMinus8; createInfo_.OutputFormat = cudaVideoSurfaceFormat_NV12; createInfo_.DeinterlaceMode = static_cast(videoFormat.deinterlaceMode); createInfo_.ulTargetWidth = videoFormat.width; diff --git a/modules/cudacodec/src/video_reader.cpp b/modules/cudacodec/src/video_reader.cpp index b096d2a0f7e..998efa959a3 100644 --- a/modules/cudacodec/src/video_reader.cpp +++ b/modules/cudacodec/src/video_reader.cpp @@ -74,6 +74,9 @@ void videoDecPostProcessFrame(const GpuMat& decodedFrame, GpuMat& outFrame, int outFrame.create(height, width, CV_8UC1); cudaMemcpy2DAsync(outFrame.ptr(), outFrame.step, decodedFrame.ptr(), decodedFrame.step, width, height, cudaMemcpyDeviceToDevice, stream); } + else if (colorFormat == ColorFormat::YUV) { + decodedFrame.copyTo(outFrame); + } } using namespace cv::cudacodec::detail; diff --git a/modules/cudacodec/test/test_video.cpp b/modules/cudacodec/test/test_video.cpp index 980f393302a..0a9a3c5362c 100644 --- a/modules/cudacodec/test/test_video.cpp +++ b/modules/cudacodec/test/test_video.cpp @@ -187,6 +187,7 @@ CUDA_TEST_P(Video, Reader) {cudacodec::ColorFormat::GRAY,1}, {cudacodec::ColorFormat::BGR,3}, {cudacodec::ColorFormat::BGRA,4}, + {cudacodec::ColorFormat::YUV,1} }; std::string inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "../" + GET_PARAM(1); @@ -201,7 +202,8 @@ CUDA_TEST_P(Video, Reader) ASSERT_TRUE(reader->nextFrame(frame)); if(!fmt.valid) fmt = reader->format(); - ASSERT_TRUE(frame.cols == fmt.width && frame.rows == fmt.height); + const int height = formatToChannels.first == cudacodec::ColorFormat::YUV ? 1.5 * fmt.height : fmt.height; + ASSERT_TRUE(frame.cols == fmt.width && frame.rows == height); ASSERT_FALSE(frame.empty()); ASSERT_TRUE(frame.channels() == formatToChannels.second); } From 4ea24bfa407907cd90d66cef24ec6f37944369f9 Mon Sep 17 00:00:00 2001 From: cudawarped <12133430+cudawarped@users.noreply.github.com> Date: Wed, 25 May 2022 22:19:14 +0100 Subject: [PATCH 2/3] Update to include missing CUDA stream argument to raw frame copy. --- modules/cudacodec/src/video_reader.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/modules/cudacodec/src/video_reader.cpp b/modules/cudacodec/src/video_reader.cpp index 998efa959a3..604e4e2004f 100644 --- a/modules/cudacodec/src/video_reader.cpp +++ b/modules/cudacodec/src/video_reader.cpp @@ -56,10 +56,10 @@ Ptr cv::cudacodec::createVideoReader(const Ptr&, co void nv12ToBgra(const GpuMat& decodedFrame, GpuMat& outFrame, int width, int height, cudaStream_t stream); void videoDecPostProcessFrame(const GpuMat& decodedFrame, GpuMat& outFrame, int width, int height, const ColorFormat colorFormat, - cudaStream_t stream) + Stream stream) { if (colorFormat == ColorFormat::BGRA) { - nv12ToBgra(decodedFrame, outFrame, width, height, stream); + nv12ToBgra(decodedFrame, outFrame, width, height, StreamAccessor::getStream(stream)); } else if (colorFormat == ColorFormat::BGR) { outFrame.create(height, width, CV_8UC3); @@ -67,15 +67,15 @@ void videoDecPostProcessFrame(const GpuMat& decodedFrame, GpuMat& outFrame, int NppiSize oSizeROI = { width,height }; NppStreamContext nppStreamCtx; nppSafeCall(nppGetStreamContext(&nppStreamCtx)); - nppStreamCtx.hStream = stream; + nppStreamCtx.hStream = StreamAccessor::getStream(stream); nppSafeCall(nppiNV12ToBGR_8u_P2C3R_Ctx(pSrc, decodedFrame.step, outFrame.data, outFrame.step, oSizeROI, nppStreamCtx)); } else if (colorFormat == ColorFormat::GRAY) { outFrame.create(height, width, CV_8UC1); - cudaMemcpy2DAsync(outFrame.ptr(), outFrame.step, decodedFrame.ptr(), decodedFrame.step, width, height, cudaMemcpyDeviceToDevice, stream); + cudaMemcpy2DAsync(outFrame.ptr(), outFrame.step, decodedFrame.ptr(), decodedFrame.step, width, height, cudaMemcpyDeviceToDevice, StreamAccessor::getStream(stream)); } else if (colorFormat == ColorFormat::YUV) { - decodedFrame.copyTo(outFrame); + decodedFrame.copyTo(outFrame, stream); } } @@ -220,7 +220,7 @@ namespace // perform post processing on the CUDA surface (performs colors space conversion and post processing) // comment this out if we include the line of code seen above - videoDecPostProcessFrame(decodedFrame, frame, videoDecoder_->targetWidth(), videoDecoder_->targetHeight(), colorFormat, StreamAccessor::getStream(stream)); + videoDecPostProcessFrame(decodedFrame, frame, videoDecoder_->targetWidth(), videoDecoder_->targetHeight(), colorFormat, stream); // unmap video frame // unmapFrame() synchronizes with the VideoDecode API (ensures the frame has finished decoding) From 75c1eb0da36cf5b859c3980bd6d1a1dd6edcd0ec Mon Sep 17 00:00:00 2001 From: cudawarped <12133430+cudawarped@users.noreply.github.com> Date: Thu, 26 May 2022 12:27:49 +0100 Subject: [PATCH 3/3] Fix copy paste oversight. (cherry picked from commit 8f6820cdeb1e5ac5ec48d2d6387f5695517e5506) --- modules/cudacodec/src/video_reader.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/cudacodec/src/video_reader.cpp b/modules/cudacodec/src/video_reader.cpp index 604e4e2004f..5a21357189f 100644 --- a/modules/cudacodec/src/video_reader.cpp +++ b/modules/cudacodec/src/video_reader.cpp @@ -99,7 +99,7 @@ namespace bool set(const VideoReaderProps propertyId, const double propertyVal) CV_OVERRIDE; - void VideoReaderImpl::set(const ColorFormat _colorFormat) CV_OVERRIDE; + void set(const ColorFormat _colorFormat) CV_OVERRIDE; bool get(const VideoReaderProps propertyId, double& propertyVal) const CV_OVERRIDE;