Merge pull request #2180 from cudawarped:example_fix_for_cudacodec

OpenCV fix 14774 breaks cudacodec

* Example to describe comilation issue

* Added required changes, builds with -DWITH_FFMPEG=ON|OFF

* Working with standard ffmpeg cap.

* Changed cudacodec to use new retrieveRaw() function, to retrieve the raw encoded bitstream, from the videoio module  instead of its own implementation.

* Fix cv::cudacodec::VideoWriter

* Update to use VideoContainer

* Remove header used in testing

* Remove warning

* remove dependence on redundant ffmpeg codecs

* cudacodec: use .set(CAP_PROP_FORMAT, -1) to extract RAW streams

* whitespace

* addressed alalek's comment

modules/cudacodec/include/opencv2/cudacodec.hpp

@@ -59,6 +59,8 @@
 
 namespace cv { namespace cudacodec {
 
+using namespace cuda;  // Stream
+
 //! @addtogroup cudacodec
 //! @{
 
@@ -253,6 +255,7 @@ enum Codec
     HEVC,
     VP8,
     VP9,
+    NumCodecs,
 
     Uncompressed_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V')),   //!< Y,U,V (4:2:0)
     Uncompressed_YV12   = (('Y'<<24)|('V'<<16)|('1'<<8)|('2')),   //!< Y,V,U (4:2:0)
@@ -268,7 +271,8 @@ enum ChromaFormat
     Monochrome = 0,
     YUV420,
     YUV422,
-    YUV444
+    YUV444,
+    NumFormats
 };
 
 /** @brief Struct providing information about video file format. :
@@ -298,7 +302,7 @@ public:
     If no frames has been grabbed (there are no more frames in video file), the methods return false .
     The method throws Exception if error occurs.
      */
-    CV_WRAP virtual bool nextFrame(OutputArray frame) = 0;
+    CV_WRAP virtual bool nextFrame(OutputArray frame, Stream &stream = Stream::Null()) = 0;
 
     /** @brief Returns information about video file format.
     */
@@ -318,9 +322,8 @@ public:
 
     @param data Pointer to frame data.
     @param size Size in bytes of current frame.
-    @param endOfFile Indicates that it is end of stream.
      */
-    virtual bool getNextPacket(unsigned char** data, int* size, bool* endOfFile) = 0;
+    virtual bool getNextPacket(unsigned char** data, size_t* size) = 0;
 
     /** @brief Returns information about video file format.
     */

modules/cudacodec/perf/perf_video.cpp

@@ -46,14 +46,21 @@
 
 namespace opencv_test { namespace {
 
+#if defined(HAVE_NVCUVID)
+
+#if defined(HAVE_FFMPEG_WRAPPER) // should this be set in preprocessor or in cvconfig.h
+#define VIDEO_SRC Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi")
+#else
+// CUDA demuxer has to fall back to ffmpeg to process "gpu/video/768x576.avi"
+#define VIDEO_SRC Values( "gpu/video/1920x1080.avi")
+#endif
+
 DEF_PARAM_TEST_1(FileName, string);
 
 //////////////////////////////////////////////////////
 // VideoReader
 
-#if defined(HAVE_NVCUVID)
-
-PERF_TEST_P(FileName, VideoReader, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
+PERF_TEST_P(FileName, VideoReader, VIDEO_SRC)
 {
     declare.time(20);
 
@@ -89,7 +96,7 @@ PERF_TEST_P(FileName, VideoReader, Values("gpu/video/768x576.avi", "gpu/video/19
 
 #if defined(HAVE_NVCUVID) && defined(_WIN32)
 
-PERF_TEST_P(FileName, VideoWriter, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
+PERF_TEST_P(FileName, VideoWriter, VIDEO_SRC)
 {
     declare.time(30);
 

modules/cudacodec/src/cuda/nv12_to_rgb.cu

@@ -60,7 +60,7 @@
 using namespace cv;
 using namespace cv::cudev;
 
-void videoDecPostProcessFrame(const GpuMat& decodedFrame, OutputArray _outFrame, int width, int height);
+void videoDecPostProcessFrame(const GpuMat& decodedFrame, OutputArray _outFrame, int width, int height, cudaStream_t stream);
 
 namespace
 {
@@ -186,7 +186,7 @@ namespace
     }
 }
 
-void videoDecPostProcessFrame(const GpuMat& decodedFrame, OutputArray _outFrame, int width, int height)
+void videoDecPostProcessFrame(const GpuMat& decodedFrame, OutputArray _outFrame, int width, int height, cudaStream_t stream)
 {
     // Final Stage: NV12toARGB color space conversion
 
@@ -196,12 +196,13 @@ void videoDecPostProcessFrame(const GpuMat& decodedFrame, OutputArray _outFrame,
     dim3 block(32, 8);
     dim3 grid(divUp(width, 2 * block.x), divUp(height, block.y));
 
-    NV12_to_RGB<<<grid, block>>>(decodedFrame.ptr<uchar>(), decodedFrame.step,
+    NV12_to_RGB<<<grid, block, 0, stream>>>(decodedFrame.ptr<uchar>(), decodedFrame.step,
                                  outFrame.ptr<uint>(), outFrame.step,
                                  width, height);
 
     CV_CUDEV_SAFE_CALL( cudaGetLastError() );
-    CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
+    if (stream == 0)
+      CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() );
 }
 
 #endif

modules/cudacodec/src/cuvid_video_source.cpp

@@ -68,6 +68,7 @@ cv::cudacodec::detail::CuvidVideoSource::CuvidVideoSource(const String& fname)
     CUVIDEOFORMAT vidfmt;
     cuSafeCall( cuvidGetSourceVideoFormat(videoSource_, &vidfmt, 0) );
 
+    CV_Assert(Codec::NumCodecs == cudaVideoCodec::cudaVideoCodec_NumCodecs);
     format_.codec = static_cast<Codec>(vidfmt.codec);
     format_.chromaFormat = static_cast<ChromaFormat>(vidfmt.chroma_format);
     format_.nBitDepthMinus8 = vidfmt.bit_depth_luma_minus8;

modules/cudacodec/src/ffmpeg_video_source.cpp

@@ -44,82 +44,100 @@
 #include "precomp.hpp"
 
 #ifdef HAVE_NVCUVID
-
 using namespace cv;
 using namespace cv::cudacodec;
 using namespace cv::cudacodec::detail;
 
-namespace
+#ifndef CV_FOURCC_MACRO
+#define CV_FOURCC_MACRO(c1, c2, c3, c4) (((c1) & 255) + (((c2) & 255) << 8) + (((c3) & 255) << 16) + (((c4) & 255) << 24))
+#endif
+
+static std::string fourccToString(int fourcc)
 {
-    Create_InputMediaStream_FFMPEG_Plugin create_InputMediaStream_FFMPEG_p = 0;
-    Release_InputMediaStream_FFMPEG_Plugin release_InputMediaStream_FFMPEG_p = 0;
-    Read_InputMediaStream_FFMPEG_Plugin read_InputMediaStream_FFMPEG_p = 0;
+    union {
+        int u32;
+        unsigned char c[4];
+    } i32_c;
+    i32_c.u32 = fourcc;
+    return cv::format("%c%c%c%c",
+        (i32_c.c[0] >= ' ' && i32_c.c[0] < 128) ? i32_c.c[0] : '?',
+        (i32_c.c[1] >= ' ' && i32_c.c[1] < 128) ? i32_c.c[1] : '?',
+        (i32_c.c[2] >= ' ' && i32_c.c[2] < 128) ? i32_c.c[2] : '?',
+        (i32_c.c[3] >= ' ' && i32_c.c[3] < 128) ? i32_c.c[3] : '?');
+}
 
-    bool init_MediaStream_FFMPEG()
+static
+Codec FourccToCodec(int codec)
+{
+    switch (codec)
     {
-        static bool initialized = 0;
-
-        if (!initialized)
-        {
-            #if defined _WIN32
-                const char* module_name = "opencv_ffmpeg"
-                    CVAUX_STR(CV_VERSION_MAJOR) CVAUX_STR(CV_VERSION_MINOR) CVAUX_STR(CV_VERSION_REVISION)
-                #if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__)
-                    "_64"
-                #endif
-                    ".dll";
-
-                static HMODULE cvFFOpenCV = LoadLibrary(module_name);
-
-                if (cvFFOpenCV)
-                {
-                    create_InputMediaStream_FFMPEG_p =
-                        (Create_InputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "create_InputMediaStream_FFMPEG");
-                    release_InputMediaStream_FFMPEG_p =
-                        (Release_InputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "release_InputMediaStream_FFMPEG");
-                    read_InputMediaStream_FFMPEG_p =
-                        (Read_InputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "read_InputMediaStream_FFMPEG");
-
-                    initialized = create_InputMediaStream_FFMPEG_p != 0 && release_InputMediaStream_FFMPEG_p != 0 && read_InputMediaStream_FFMPEG_p != 0;
-                }
-            #elif defined HAVE_FFMPEG
-                create_InputMediaStream_FFMPEG_p = create_InputMediaStream_FFMPEG;
-                release_InputMediaStream_FFMPEG_p = release_InputMediaStream_FFMPEG;
-                read_InputMediaStream_FFMPEG_p = read_InputMediaStream_FFMPEG;
-
-                initialized = true;
-            #endif
-        }
-
-        return initialized;
+    case CV_FOURCC_MACRO('m', 'p', 'e', 'g'): // fallthru
+    case CV_FOURCC_MACRO('M', 'P', 'G', '1'): return MPEG1;
+    case CV_FOURCC_MACRO('M', 'P', 'G', '2'): return MPEG2;
+    case CV_FOURCC_MACRO('X', 'V', 'I', 'D'): // fallthru
+    case CV_FOURCC_MACRO('D', 'I', 'V', 'X'): return MPEG4;
+    case CV_FOURCC_MACRO('W', 'V', 'C', '1'): return VC1;
+    case CV_FOURCC_MACRO('H', '2', '6', '4'): // fallthru
+    case CV_FOURCC_MACRO('h', '2', '6', '4'): // fallthru
+    case CV_FOURCC_MACRO('a', 'v', 'c', '1'): return H264;
+    case CV_FOURCC_MACRO('H', '2', '6', '5'): // fallthru
+    case CV_FOURCC_MACRO('h', '2', '6', '5'): // fallthru
+    case CV_FOURCC_MACRO('h', 'e', 'v', 'c'): return HEVC;
+    case CV_FOURCC_MACRO('M', 'J', 'P', 'G'): return JPEG;
+    case CV_FOURCC_MACRO('V', 'P', '8', '0'): return VP8;
+    case CV_FOURCC_MACRO('V', 'P', '9', '0'): return VP9;
+    default:
+        break;
     }
+
+    std::string msg = cv::format("Unknown codec FOURCC: 0x%08X (%s)", codec, fourccToString(codec).c_str());
+    CV_LOG_WARNING(NULL, msg);
+    CV_Error(Error::StsUnsupportedFormat, msg);
 }
 
-cv::cudacodec::detail::FFmpegVideoSource::FFmpegVideoSource(const String& fname) :
-    stream_(0)
+static
+void FourccToChromaFormat(const int pixelFormat, ChromaFormat &chromaFormat, int & nBitDepthMinus8)
 {
-    CV_Assert( init_MediaStream_FFMPEG() );
+    switch (pixelFormat)
+    {
+    case CV_FOURCC_MACRO('I', '4', '2', '0'):
+        chromaFormat = YUV420;
+        nBitDepthMinus8 = 0;
+        break;
+    default:
+        CV_LOG_WARNING(NULL, cv::format("ChromaFormat not recognized: 0x%08X (%s). Assuming I420", pixelFormat, fourccToString(pixelFormat).c_str()));
+        chromaFormat = YUV420;
+        nBitDepthMinus8 = 0;
+        break;
+    }
+}
 
-    int codec;
-    int chroma_format;
-    int width;
-    int height;
+cv::cudacodec::detail::FFmpegVideoSource::FFmpegVideoSource(const String& fname)
+{
+    if (!videoio_registry::hasBackend(CAP_FFMPEG))
+        CV_Error(Error::StsNotImplemented, "FFmpeg backend not found");
 
-    stream_ = create_InputMediaStream_FFMPEG_p(fname.c_str(), &codec, &chroma_format, &width, &height);
-    if (!stream_)
+    cap.open(fname, CAP_FFMPEG);
+    if (!cap.isOpened())
         CV_Error(Error::StsUnsupportedFormat, "Unsupported video source");
 
-    format_.codec = static_cast<Codec>(codec);
-    format_.chromaFormat = static_cast<ChromaFormat>(chroma_format);
-    format_.nBitDepthMinus8 = -1;
-    format_.width = width;
-    format_.height = height;
+    if (!cap.set(CAP_PROP_FORMAT, -1))  // turn off video decoder (extract stream)
+        CV_Error(Error::StsUnsupportedFormat, "Fetching of RAW video streams is not supported");
+    CV_Assert(cap.get(CAP_PROP_FORMAT) == -1);
+
+    int codec = (int)cap.get(CAP_PROP_FOURCC);
+    int pixelFormat = (int)cap.get(CAP_PROP_CODEC_PIXEL_FORMAT);
+
+    format_.codec = FourccToCodec(codec);
+    format_.height = cap.get(CAP_PROP_FRAME_HEIGHT);
+    format_.width = cap.get(CAP_PROP_FRAME_WIDTH);
+    FourccToChromaFormat(pixelFormat, format_.chromaFormat, format_.nBitDepthMinus8);
 }
 
 cv::cudacodec::detail::FFmpegVideoSource::~FFmpegVideoSource()
 {
-    if (stream_)
-        release_InputMediaStream_FFMPEG_p(stream_);
+    if (cap.isOpened())
+        cap.release();
 }
 
 FormatInfo cv::cudacodec::detail::FFmpegVideoSource::format() const
@@ -127,14 +145,12 @@ FormatInfo cv::cudacodec::detail::FFmpegVideoSource::format() const
     return format_;
 }
 
-bool cv::cudacodec::detail::FFmpegVideoSource::getNextPacket(unsigned char** data, int* size, bool* bEndOfFile)
+bool cv::cudacodec::detail::FFmpegVideoSource::getNextPacket(unsigned char** data, size_t* size)
 {
-    int endOfFile;
-
-    int res = read_InputMediaStream_FFMPEG_p(stream_, data, size, &endOfFile);
-
-    *bEndOfFile = (endOfFile != 0);
-    return res != 0;
+    cap >> rawFrame;
+    *data = rawFrame.data;
+    *size = rawFrame.total();
+    return *size != 0;
 }
 
 #endif // HAVE_CUDA

modules/cudacodec/src/ffmpeg_video_source.hpp

@@ -46,8 +46,6 @@
 
 #include "opencv2/cudacodec.hpp"
 
-struct InputMediaStream_FFMPEG;
-
 namespace cv { namespace cudacodec { namespace detail {
 
 class FFmpegVideoSource : public RawVideoSource
@@ -56,14 +54,14 @@ public:
     FFmpegVideoSource(const String& fname);
     ~FFmpegVideoSource();
 
-    bool getNextPacket(unsigned char** data, int* size, bool* endOfFile) CV_OVERRIDE;
+    bool getNextPacket(unsigned char** data, size_t* size) CV_OVERRIDE;
 
     FormatInfo format() const CV_OVERRIDE;
 
 private:
     FormatInfo format_;
-
-    InputMediaStream_FFMPEG* stream_;
+    VideoCapture cap;
+    Mat rawFrame;
 };
 
 }}}

modules/cudacodec/src/precomp.hpp

@@ -52,8 +52,10 @@
 #include <iostream>
 
 #include "opencv2/cudacodec.hpp"
-
+#include "opencv2/videoio.hpp"
+#include "opencv2/videoio/registry.hpp"
 #include "opencv2/core/private.cuda.hpp"
+#include <opencv2/core/utils/logger.hpp>
 
 #ifdef HAVE_NVCUVID
     #if CUDA_VERSION >= 9000 && CUDA_VERSION < 10000
@@ -81,7 +83,6 @@
     #include "video_decoder.hpp"
     #include "video_parser.hpp"
 
-    #include "../src/cap_ffmpeg_api.hpp"
 #endif
 
 #endif /* OPENCV_PRECOMP_H */

modules/cudacodec/src/video_decoder.cpp

@@ -47,7 +47,8 @@
 
 void cv::cudacodec::detail::VideoDecoder::create(const FormatInfo& videoFormat)
 {
-    release();
+    if (videoFormat.nBitDepthMinus8 > 0 || videoFormat.chromaFormat != YUV420)
+        CV_Error(Error::StsUnsupportedFormat, "NV12 output requires 8 bit YUV420");
 
     cudaVideoCodec _codec = static_cast<cudaVideoCodec>(videoFormat.codec);
     cudaVideoChromaFormat _chromaFormat = static_cast<cudaVideoChromaFormat>(videoFormat.chromaFormat);
@@ -89,20 +90,22 @@ void cv::cudacodec::detail::VideoDecoder::create(const FormatInfo& videoFormat)
 
 #if (CUDART_VERSION >= 9000)
     // Check video format is supported by GPU's hardware video decoder
-    if (videoFormat.nBitDepthMinus8 != -1) { // info not available call to create CuvidVideoSource() failed
-        CUVIDDECODECAPS decodeCaps = {};
-        decodeCaps.eCodecType = _codec;
-        decodeCaps.eChromaFormat = _chromaFormat;
-        decodeCaps.nBitDepthMinus8 = videoFormat.nBitDepthMinus8;
-        cuSafeCall(cuvidGetDecoderCaps(&decodeCaps));
-        if (!decodeCaps.bIsSupported)
-            CV_Error(Error::StsUnsupportedFormat, "Video source is not supported by hardware video decoder");
-
-        CV_Assert(videoFormat.width >= decodeCaps.nMinWidth &&
-            videoFormat.height >= decodeCaps.nMinHeight &&
-            videoFormat.width <= decodeCaps.nMaxWidth &&
-            videoFormat.height <= decodeCaps.nMaxHeight);
-    }
+    CUVIDDECODECAPS decodeCaps = {};
+    decodeCaps.eCodecType = _codec;
+    decodeCaps.eChromaFormat = _chromaFormat;
+    decodeCaps.nBitDepthMinus8 = videoFormat.nBitDepthMinus8;
+    cuSafeCall(cuCtxPushCurrent(ctx_));
+    cuSafeCall(cuvidGetDecoderCaps(&decodeCaps));
+    cuSafeCall(cuCtxPopCurrent(NULL));
+    if (!decodeCaps.bIsSupported)
+        CV_Error(Error::StsUnsupportedFormat, "Video source is not supported by hardware video decoder");
+
+    CV_Assert(videoFormat.width >= decodeCaps.nMinWidth &&
+        videoFormat.height >= decodeCaps.nMinHeight &&
+        videoFormat.width <= decodeCaps.nMaxWidth &&
+        videoFormat.height <= decodeCaps.nMaxHeight);
+
+    CV_Assert((videoFormat.width >> 4)* (videoFormat.height >> 4) <= decodeCaps.nMaxMBCount);
 #endif
 
     // Fill the decoder-create-info struct from the given video-format struct.
@@ -113,11 +116,6 @@ void cv::cudacodec::detail::VideoDecoder::create(const FormatInfo& videoFormat)
     createInfo_.ulWidth             = videoFormat.width;
     createInfo_.ulHeight            = videoFormat.height;
     createInfo_.ulNumDecodeSurfaces = FrameQueue::MaximumSize;
-
-    // Limit decode memory to 24MB (16M pixels at 4:2:0 = 24M bytes)
-    while (createInfo_.ulNumDecodeSurfaces * videoFormat.width * videoFormat.height > 16 * 1024 * 1024)
-        createInfo_.ulNumDecodeSurfaces--;
-
     createInfo_.ChromaFormat    = _chromaFormat;
     createInfo_.OutputFormat    = cudaVideoSurfaceFormat_NV12;
     createInfo_.DeinterlaceMode = cudaVideoDeinterlaceMode_Adaptive;
@@ -131,8 +129,9 @@ void cv::cudacodec::detail::VideoDecoder::create(const FormatInfo& videoFormat)
     createInfo_.ulCreationFlags     = videoCreateFlags;
     createInfo_.vidLock = lock_;
 
-    // create the decoder
-    cuSafeCall( cuvidCreateDecoder(&decoder_, &createInfo_) );
+    cuSafeCall(cuCtxPushCurrent(ctx_));
+    cuSafeCall(cuvidCreateDecoder(&decoder_, &createInfo_));
+    cuSafeCall(cuCtxPopCurrent(NULL));
 }
 
 void cv::cudacodec::detail::VideoDecoder::release()

modules/cudacodec/src/video_decoder.hpp

@@ -59,7 +59,7 @@ namespace cv { namespace cudacodec { namespace detail
 class VideoDecoder
 {
 public:
-    VideoDecoder(const FormatInfo& videoFormat, CUvideoctxlock lock) : lock_(lock), decoder_(0)
+    VideoDecoder(const FormatInfo& videoFormat, CUcontext ctx, CUvideoctxlock lock) : ctx_(ctx), lock_(lock), decoder_(0)
     {
         create(videoFormat);
     }
@@ -83,6 +83,7 @@ public:
     unsigned long targetHeight() const { return createInfo_.ulTargetHeight; }
 
     cudaVideoChromaFormat chromaFormat() const { return createInfo_.ChromaFormat; }
+    int nBitDepthMinus8() const { return createInfo_.bitDepthMinus8; }
 
     bool decodePicture(CUVIDPICPARAMS* picParams)
     {
@@ -96,6 +97,7 @@ public:
 
         cuSafeCall( cuvidMapVideoFrame(decoder_, picIdx, &ptr, &pitch, &videoProcParams) );
 
+
         return cuda::GpuMat(targetHeight() * 3 / 2, targetWidth(), CV_8UC1, (void*) ptr, pitch);
     }
 
@@ -107,6 +109,7 @@ public:
 
 private:
     CUvideoctxlock lock_;
+    CUcontext ctx_;
     CUVIDDECODECREATEINFO createInfo_;
     CUvideodecoder        decoder_;
 };

modules/cudacodec/src/video_parser.cpp

@@ -80,7 +80,7 @@ bool cv::cudacodec::detail::VideoParser::parseVideoData(const unsigned char* dat
         return false;
     }
 
-    const int maxUnparsedPackets = 15;
+    const int maxUnparsedPackets = 20;
 
     ++unparsedPackets_;
     if (unparsedPackets_ > maxUnparsedPackets)
@@ -105,7 +105,8 @@ int CUDAAPI cv::cudacodec::detail::VideoParser::HandleVideoSequence(void* userDa
     if (format->codec         != thiz->videoDecoder_->codec()       ||
         format->coded_width   != thiz->videoDecoder_->frameWidth()  ||
         format->coded_height  != thiz->videoDecoder_->frameHeight() ||
-        format->chroma_format != thiz->videoDecoder_->chromaFormat())
+        format->chroma_format != thiz->videoDecoder_->chromaFormat()||
+        format->bit_depth_luma_minus8 != thiz->videoDecoder_->nBitDepthMinus8())
     {
         FormatInfo newFormat;
 
@@ -113,6 +114,7 @@ int CUDAAPI cv::cudacodec::detail::VideoParser::HandleVideoSequence(void* userDa
         newFormat.chromaFormat = static_cast<ChromaFormat>(format->chroma_format);
         newFormat.width = format->coded_width;
         newFormat.height = format->coded_height;
+        newFormat.nBitDepthMinus8 = format->bit_depth_luma_minus8;
 
         try
         {

modules/cudacodec/src/video_reader.cpp

@@ -53,7 +53,7 @@ Ptr<VideoReader> cv::cudacodec::createVideoReader(const Ptr<RawVideoSource>&) {
 
 #else // HAVE_NVCUVID
 
-void videoDecPostProcessFrame(const GpuMat& decodedFrame, OutputArray _outFrame, int width, int height);
+void videoDecPostProcessFrame(const GpuMat& decodedFrame, OutputArray _outFrame, int width, int height, cudaStream_t stream);
 
 using namespace cv::cudacodec::detail;
 
@@ -65,7 +65,7 @@ namespace
         explicit VideoReaderImpl(const Ptr<VideoSource>& source);
         ~VideoReaderImpl();
 
-        bool nextFrame(OutputArray frame) CV_OVERRIDE;
+        bool nextFrame(OutputArray frame, Stream& stream) CV_OVERRIDE;
 
         FormatInfo format() const CV_OVERRIDE;
 
@@ -99,7 +99,7 @@ namespace
         cuSafeCall( cuvidCtxLockCreate(&lock_, ctx) );
 
         frameQueue_.reset(new FrameQueue);
-        videoDecoder_.reset(new VideoDecoder(videoSource_->format(), lock_));
+        videoDecoder_.reset(new VideoDecoder(videoSource_->format(), ctx, lock_));
         videoParser_.reset(new VideoParser(videoDecoder_, frameQueue_));
 
         videoSource_->setVideoParser(videoParser_);
@@ -122,14 +122,11 @@ namespace
         CUvideoctxlock m_lock;
     };
 
-    bool VideoReaderImpl::nextFrame(OutputArray frame)
+    bool VideoReaderImpl::nextFrame(OutputArray frame, Stream& stream)
     {
         if (videoSource_->hasError() || videoParser_->hasError())
             CV_Error(Error::StsUnsupportedFormat, "Unsupported video source");
 
-        if (!videoSource_->isStarted() || frameQueue_->isEndOfDecode())
-            return false;
-
         if (frames_.empty())
         {
             CUVIDPARSERDISPINFO displayInfo;
@@ -180,7 +177,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());
+            videoDecPostProcessFrame(decodedFrame, frame, videoDecoder_->targetWidth(), videoDecoder_->targetHeight(), StreamAccessor::getStream(stream));
 
             // unmap video frame
             // unmapFrame() synchronizes with the VideoDecode API (ensures the frame has finished decoding)
@@ -203,12 +200,13 @@ Ptr<VideoReader> cv::cudacodec::createVideoReader(const String& filename)
 
     try
     {
-        videoSource.reset(new CuvidVideoSource(filename));
+        // prefer ffmpeg to cuvidGetSourceVideoFormat() which doesn't always return the corrct raw pixel format
+        Ptr<RawVideoSource> source(new FFmpegVideoSource(filename));
+        videoSource.reset(new RawVideoSourceWrapper(source));
     }
     catch (...)
     {
-        Ptr<RawVideoSource> source(new FFmpegVideoSource(filename));
-        videoSource.reset(new RawVideoSourceWrapper(source));
+        videoSource.reset(new CuvidVideoSource(filename));
     }
 
     return makePtr<VideoReaderImpl>(videoSource);

modules/cudacodec/src/video_source.cpp

@@ -96,12 +96,11 @@ void cv::cudacodec::detail::RawVideoSourceWrapper::readLoop(void* userData)
     for (;;)
     {
         unsigned char* data;
-        int size;
-        bool endOfFile;
+        size_t size;
 
-        if (!thiz->source_->getNextPacket(&data, &size, &endOfFile))
+        if (!thiz->source_->getNextPacket(&data, &size))
         {
-            thiz->hasError_ = !endOfFile;
+            thiz->hasError_ = false;
             break;
         }
 

modules/cudacodec/src/video_writer.cpp

@@ -54,11 +54,11 @@ cv::cudacodec::EncoderParams::EncoderParams(const String&) { throw_no_cuda(); }
 void cv::cudacodec::EncoderParams::load(const String&) { throw_no_cuda(); }
 void cv::cudacodec::EncoderParams::save(const String&) const { throw_no_cuda(); }
 
-Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const String&, Size, double, SurfaceFormat) { throw_no_cuda(); return Ptr<VideoWriter>(); }
-Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const String&, Size, double, const EncoderParams&, SurfaceFormat) { throw_no_cuda(); return Ptr<VideoWriter>(); }
+Ptr<cv::cudacodec::VideoWriter> cv::cudacodec::createVideoWriter(const String&, Size, double, SurfaceFormat) { throw_no_cuda(); return Ptr<cv::cudacodec::VideoWriter>(); }
+Ptr<cv::cudacodec::VideoWriter> cv::cudacodec::createVideoWriter(const String&, Size, double, const EncoderParams&, SurfaceFormat) { throw_no_cuda(); return Ptr<cv::cudacodec::VideoWriter>(); }
 
-Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>&, Size, double, SurfaceFormat) { throw_no_cuda(); return Ptr<VideoWriter>(); }
-Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>&, Size, double, const EncoderParams&, SurfaceFormat) { throw_no_cuda(); return Ptr<VideoWriter>(); }
+Ptr<cv::cudacodec::VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>&, Size, double, SurfaceFormat) { throw_no_cuda(); return Ptr<cv::cudacodec::VideoWriter>(); }
+Ptr<cv::cudacodec::VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>&, Size, double, const EncoderParams&, SurfaceFormat) { throw_no_cuda(); return Ptr<cv::cudacodec::VideoWriter>(); }
 
 #else // !defined HAVE_NVCUVENC || !defined _WIN32
 
@@ -913,4 +913,4 @@ Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>& en
     return makePtr<VideoWriterImpl>(encoderCallback, frameSize, fps, params, format);
 }
 
-#endif // !defined HAVE_NVCUVENC || !defined _WIN32
+#endif // !defined HAVE_NVCUVENC || !defined _WIN32 || defined HAVE_FFMPEG_WRAPPER

modules/cudacodec/test/test_precomp.hpp

@@ -43,6 +43,7 @@
 #define OPENCV_TEST_PRECOMP_HPP
 
 #include "opencv2/ts.hpp"
+#include "opencv2/videoio/registry.hpp"
 #include "opencv2/ts/cuda_test.hpp"
 
 #include "opencv2/cudacodec.hpp"

modules/cudacodec/test/test_video.cpp

@@ -41,15 +41,15 @@
 //M*/
 
 #include "test_precomp.hpp"
+namespace opencv_test {
+    namespace {
 
-namespace opencv_test { namespace {
-
-#ifdef HAVE_NVCUVID
-
+#if defined(HAVE_NVCUVID) || defined(HAVE_NVCUVENC)
 PARAM_TEST_CASE(Video, cv::cuda::DeviceInfo, std::string)
 {
 };
 
+#if defined(HAVE_NVCUVID)
 //////////////////////////////////////////////////////
 // VideoReader
 
@@ -57,24 +57,26 @@ CUDA_TEST_P(Video, Reader)
 {
     cv::cuda::setDevice(GET_PARAM(0).deviceID());
 
-    const std::string inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + GET_PARAM(1);
+    // CUDA demuxer has to fall back to ffmpeg to process "gpu/video/768x576.avi"
+    if (GET_PARAM(1) == "gpu/video/768x576.avi" && !videoio_registry::hasBackend(CAP_FFMPEG))
+        throw SkipTestException("FFmpeg backend not found");
 
+    std::string inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "../" + GET_PARAM(1);
     cv::Ptr<cv::cudacodec::VideoReader> reader = cv::cudacodec::createVideoReader(inputFile);
 
     cv::cuda::GpuMat frame;
-
-    for (int i = 0; i < 10; ++i)
+    for (int i = 0; i < 100; i++)
     {
         ASSERT_TRUE(reader->nextFrame(frame));
         ASSERT_FALSE(frame.empty());
     }
 }
+#endif // HAVE_NVCUVID
 
+#if defined(_WIN32) && defined(HAVE_NVCUVENC)
 //////////////////////////////////////////////////////
 // VideoWriter
 
-#ifdef _WIN32
-
 CUDA_TEST_P(Video, Writer)
 {
     cv::cuda::setDevice(GET_PARAM(0).deviceID());
@@ -118,11 +120,14 @@ CUDA_TEST_P(Video, Writer)
     }
 }
 
-#endif // _WIN32
+#endif // _WIN32, HAVE_NVCUVENC
 
+#define VIDEO_SRC "gpu/video/768x576.avi", "gpu/video/1920x1080.avi", "highgui/video/big_buck_bunny.avi", \
+    "highgui/video/big_buck_bunny.h264", "highgui/video/big_buck_bunny.h265", "highgui/video/big_buck_bunny.mpg", \
+    "highgui/video/big_buck_bunny.mpg"
 INSTANTIATE_TEST_CASE_P(CUDA_Codec, Video, testing::Combine(
     ALL_DEVICES,
-    testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi"))));
+    testing::Values(VIDEO_SRC)));
 
-#endif // HAVE_NVCUVID
+#endif // HAVE_NVCUVID || HAVE_NVCUVENC
 }} // namespace