added VideoReader_GPU

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -46,6 +46,7 @@
 #ifndef SKIP_INCLUDES
 #include <vector>
 #include <memory>
+#include <iosfwd>
 #endif
 
 #include "opencv2/core/gpumat.hpp"
@@ -1990,6 +1991,105 @@ private:
     std::auto_ptr<Impl> impl_;
 };
 
+
+////////////////////////////////// Video Decoding //////////////////////////////////////////
+
+namespace detail
+{
+    class FrameQueue;
+    class VideoParser;
+}
+
+class CV_EXPORTS VideoReader_GPU
+{
+public:
+    enum Codec
+    {
+        MPEG1 = 0,
+        MPEG2,
+        MPEG4,
+        VC1,
+        H264,
+        JPEG,
+        H264_SVC,
+        H264_MVC,
+
+        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)
+        Uncompressed_NV12   = (('N'<<24)|('V'<<16)|('1'<<8)|('2')),   // Y,UV  (4:2:0)
+        Uncompressed_YUYV   = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V')),   // YUYV/YUY2 (4:2:2)
+        Uncompressed_UYVY   = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y')),   // UYVY (4:2:2)
+    };
+
+    enum ChromaFormat
+    {
+        Monochrome=0,
+        YUV420,
+        YUV422,
+        YUV444,
+    };
+
+    struct FormatInfo
+    {
+        Codec codec;
+        ChromaFormat chromaFormat;
+        int width;
+        int height;
+    };
+
+    class VideoSource;
+
+    VideoReader_GPU();
+    explicit VideoReader_GPU(const std::string& filename);
+    explicit VideoReader_GPU(const cv::Ptr<VideoSource>& source);
+
+    ~VideoReader_GPU();
+
+    void open(const std::string& filename);
+    void open(const cv::Ptr<VideoSource>& source);
+    bool isOpened() const;
+
+    void close();
+
+    bool read(GpuMat& image);
+
+    FormatInfo format() const;
+    void dumpFormat(std::ostream& st);
+
+    class VideoSource
+    {
+    public:
+        VideoSource() : frameQueue_(0), videoParser_(0) {}
+        virtual ~VideoSource() {}
+
+        virtual FormatInfo format() const = 0;
+        virtual void start() = 0;
+        virtual void stop() = 0;
+        virtual bool isStarted() const = 0;
+        virtual bool hasError() const = 0;
+
+        void setFrameQueue(detail::FrameQueue* frameQueue) { frameQueue_ = frameQueue; }
+        void setVideoParser(detail::VideoParser* videoParser) { videoParser_ = videoParser; }
+
+    protected:
+        bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream = false);
+
+    private:
+        VideoSource(const VideoSource&);
+        VideoSource& operator =(const VideoSource&);
+
+        detail::FrameQueue* frameQueue_;
+        detail::VideoParser* videoParser_;
+    };
+
+private:
+    VideoReader_GPU(const VideoReader_GPU&);
+    VideoReader_GPU& operator =(const VideoReader_GPU&);
+
+    class Impl;
+    std::auto_ptr<Impl> impl_;
+};
+
 } // namespace gpu
 
 } // namespace cv

modules/gpu/src/cu_safe_call.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "cu_safe_call.h"
+
+#ifdef HAVE_CUDA
+
+namespace
+{    
+    #define error_entry(entry)  { entry, #entry }
+
+    struct ErrorEntry
+    {
+        int code;
+        std::string str;
+    }; 
+
+    class ErrorEntryComparer
+    {
+    public:
+        inline ErrorEntryComparer(int code) : code_(code) {}
+
+        inline bool operator()(const ErrorEntry& e) const { return e.code == code_; }
+
+    private:
+        int code_;
+    };
+
+    std::string getErrorString(int code, const ErrorEntry* errors, size_t n)
+    {
+        size_t idx = std::find_if(errors, errors + n, ErrorEntryComparer(code)) - errors;
+
+        const std::string& msg = (idx != n) ? errors[idx].str : std::string("Unknown error code");
+
+        std::ostringstream ostr;
+        ostr << msg << " [Code = " << code << "]";
+
+        return ostr.str();
+    }
+
+    const ErrorEntry cu_errors [] = 
+    {
+        error_entry( CUDA_SUCCESS                              ),
+        error_entry( CUDA_ERROR_INVALID_VALUE                  ),
+        error_entry( CUDA_ERROR_OUT_OF_MEMORY                  ),
+        error_entry( CUDA_ERROR_NOT_INITIALIZED                ),
+        error_entry( CUDA_ERROR_DEINITIALIZED                  ),
+        error_entry( CUDA_ERROR_PROFILER_DISABLED              ),
+        error_entry( CUDA_ERROR_PROFILER_NOT_INITIALIZED       ),
+        error_entry( CUDA_ERROR_PROFILER_ALREADY_STARTED       ),
+        error_entry( CUDA_ERROR_PROFILER_ALREADY_STOPPED       ),  
+        error_entry( CUDA_ERROR_NO_DEVICE                      ),
+        error_entry( CUDA_ERROR_INVALID_DEVICE                 ),
+        error_entry( CUDA_ERROR_INVALID_IMAGE                  ),
+        error_entry( CUDA_ERROR_INVALID_CONTEXT                ),
+        error_entry( CUDA_ERROR_CONTEXT_ALREADY_CURRENT        ),
+        error_entry( CUDA_ERROR_MAP_FAILED                     ),
+        error_entry( CUDA_ERROR_UNMAP_FAILED                   ),
+        error_entry( CUDA_ERROR_ARRAY_IS_MAPPED                ),
+        error_entry( CUDA_ERROR_ALREADY_MAPPED                 ),
+        error_entry( CUDA_ERROR_NO_BINARY_FOR_GPU              ),
+        error_entry( CUDA_ERROR_ALREADY_ACQUIRED               ),
+        error_entry( CUDA_ERROR_NOT_MAPPED                     ),
+        error_entry( CUDA_ERROR_NOT_MAPPED_AS_ARRAY            ),
+        error_entry( CUDA_ERROR_NOT_MAPPED_AS_POINTER          ),
+        error_entry( CUDA_ERROR_ECC_UNCORRECTABLE              ),
+        error_entry( CUDA_ERROR_UNSUPPORTED_LIMIT              ),
+        error_entry( CUDA_ERROR_CONTEXT_ALREADY_IN_USE         ),
+        error_entry( CUDA_ERROR_INVALID_SOURCE                 ),
+        error_entry( CUDA_ERROR_FILE_NOT_FOUND                 ),
+        error_entry( CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND ),
+        error_entry( CUDA_ERROR_SHARED_OBJECT_INIT_FAILED      ),
+        error_entry( CUDA_ERROR_OPERATING_SYSTEM               ),
+        error_entry( CUDA_ERROR_INVALID_HANDLE                 ),
+        error_entry( CUDA_ERROR_NOT_FOUND                      ),
+        error_entry( CUDA_ERROR_NOT_READY                      ),
+        error_entry( CUDA_ERROR_LAUNCH_FAILED                  ),
+        error_entry( CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES        ),
+        error_entry( CUDA_ERROR_LAUNCH_TIMEOUT                 ),
+        error_entry( CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING  ),
+        error_entry( CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED    ),
+        error_entry( CUDA_ERROR_PEER_ACCESS_NOT_ENABLED        ),
+        error_entry( CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE         ),
+        error_entry( CUDA_ERROR_CONTEXT_IS_DESTROYED           ),
+        error_entry( CUDA_ERROR_ASSERT                         ),
+        error_entry( CUDA_ERROR_TOO_MANY_PEERS                 ),
+        error_entry( CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED ),
+        error_entry( CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED     ),
+        error_entry( CUDA_ERROR_UNKNOWN                        )
+    };
+
+    const size_t cu_errors_num = sizeof(cu_errors) / sizeof(cu_errors[0]);
+}
+
+std::string cv::gpu::detail::cuGetErrString(CUresult res)
+{
+    return getErrorString(res, cu_errors, cu_errors_num);
+}
+
+#endif // HAVE_CUDA

modules/gpu/src/cu_safe_call.h

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef __CU_SAFE_CALL_H__
+#define __CU_SAFE_CALL_H__
+
+#include "precomp.hpp"
+
+#ifdef HAVE_CUDA
+
+namespace cv { namespace gpu { 
+    namespace detail
+    {
+        std::string cuGetErrString(CUresult res);
+
+        inline void cuSafeCall_impl(CUresult res, const char* file, int line)
+        {
+            if (res != CUDA_SUCCESS)
+                cv::error( cv::Exception(CV_GpuApiCallError, cuGetErrString(res), "unknown function", file, line) );
+        }
+    }
+}}
+
+#define cuSafeCall( op ) cv::gpu::detail::cuSafeCall_impl( (op), __FILE__, __LINE__ )
+
+#endif // HAVE_CUDA
+
+#endif // __CU_SAFE_CALL_H__

modules/gpu/src/cuda/NV12ToARGB.cu

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+/*
+ * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
+ *
+ * Please refer to the NVIDIA end user license agreement (EULA) associated
+ * with this source code for terms and conditions that govern your use of
+ * this software. Any use, reproduction, disclosure, or distribution of
+ * this software and related documentation outside the terms of the EULA
+ * is strictly prohibited.
+ *
+ */
+ 
+/*
+    NV12ToARGB color space conversion CUDA kernel
+
+    This sample uses CUDA to perform a simple NV12 (YUV 4:2:0 planar) 
+    source and converts to output in ARGB format
+*/
+
+#include "opencv2/gpu/device/common.hpp"
+
+namespace cv { namespace gpu { namespace device {
+    namespace video_decoding
+    {
+        __constant__ uint constAlpha = ((uint)0xff << 24);
+
+        __constant__ float constHueColorSpaceMat[9];
+
+        void loadHueCSC(float hueCSC[9])
+        {
+            cudaSafeCall( cudaMemcpyToSymbol(constHueColorSpaceMat, hueCSC, 9 * sizeof(float)) );
+        }
+
+        __device__ void YUV2RGB(const uint* yuvi, float* red, float* green, float* blue)
+        {
+            float luma, chromaCb, chromaCr;
+
+            // Prepare for hue adjustment
+            luma     = (float)yuvi[0];
+            chromaCb = (float)((int)yuvi[1] - 512.0f);
+            chromaCr = (float)((int)yuvi[2] - 512.0f);
+
+           // Convert YUV To RGB with hue adjustment
+           *red   = (luma     * constHueColorSpaceMat[0]) + 
+                    (chromaCb * constHueColorSpaceMat[1]) + 
+                    (chromaCr * constHueColorSpaceMat[2]);
+
+           *green = (luma     * constHueColorSpaceMat[3]) + 
+                    (chromaCb * constHueColorSpaceMat[4]) + 
+                    (chromaCr * constHueColorSpaceMat[5]);
+
+           *blue  = (luma     * constHueColorSpaceMat[6]) + 
+                    (chromaCb * constHueColorSpaceMat[7]) + 
+                    (chromaCr * constHueColorSpaceMat[8]);
+        }
+
+        __device__ uint RGBAPACK_10bit(float red, float green, float blue, uint alpha)
+        {
+            uint ARGBpixel = 0;
+
+            // Clamp final 10 bit results
+            red   = ::fmin(::fmax(red,   0.0f), 1023.f);
+            green = ::fmin(::fmax(green, 0.0f), 1023.f);
+            blue  = ::fmin(::fmax(blue,  0.0f), 1023.f);
+
+            // Convert to 8 bit unsigned integers per color component
+            ARGBpixel = (((uint)blue  >> 2) | 
+                        (((uint)green >> 2) << 8)  | 
+                        (((uint)red   >> 2) << 16) | 
+                        (uint)alpha);
+
+            return ARGBpixel;
+        }
+
+        // CUDA kernel for outputing the final ARGB output from NV12
+
+        #define COLOR_COMPONENT_BIT_SIZE 10
+        #define COLOR_COMPONENT_MASK     0x3FF
+
+        __global__ void NV12ToARGB(uchar* srcImage, size_t nSourcePitch, 
+                                   uint* dstImage, size_t nDestPitch,  
+                                   uint width, uint height)
+        {
+            // Pad borders with duplicate pixels, and we multiply by 2 because we process 2 pixels per thread
+            const int x = blockIdx.x * (blockDim.x << 1) + (threadIdx.x << 1);
+            const int y = blockIdx.y *  blockDim.y       +  threadIdx.y;
+
+            if (x >= width || y >= height)
+                return; 
+
+            // Read 2 Luma components at a time, so we don't waste processing since CbCr are decimated this way.
+            // if we move to texture we could read 4 luminance values
+
+            uint yuv101010Pel[2];
+
+            yuv101010Pel[0] = (srcImage[y * nSourcePitch + x    ]) << 2;
+            yuv101010Pel[1] = (srcImage[y * nSourcePitch + x + 1]) << 2;
+
+            const size_t chromaOffset = nSourcePitch * height;
+
+            const int y_chroma = y >> 1;
+
+            if (y & 1)  // odd scanline ?
+            {                
+                uint chromaCb = srcImage[chromaOffset + y_chroma * nSourcePitch + x    ];
+                uint chromaCr = srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1];
+
+                if (y_chroma < ((height >> 1) - 1)) // interpolate chroma vertically
+                {
+                    chromaCb = (chromaCb + srcImage[chromaOffset + (y_chroma + 1) * nSourcePitch + x    ] + 1) >> 1;
+                    chromaCr = (chromaCr + srcImage[chromaOffset + (y_chroma + 1) * nSourcePitch + x + 1] + 1) >> 1;
+                }
+                
+                yuv101010Pel[0] |= (chromaCb << ( COLOR_COMPONENT_BIT_SIZE       + 2));
+                yuv101010Pel[0] |= (chromaCr << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
+
+                yuv101010Pel[1] |= (chromaCb << ( COLOR_COMPONENT_BIT_SIZE       + 2));
+                yuv101010Pel[1] |= (chromaCr << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
+            }
+            else
+            {
+                yuv101010Pel[0] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x    ] << ( COLOR_COMPONENT_BIT_SIZE       + 2));
+                yuv101010Pel[0] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1] << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
+
+                yuv101010Pel[1] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x    ] << ( COLOR_COMPONENT_BIT_SIZE       + 2));
+                yuv101010Pel[1] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1] << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
+            }
+            
+            // this steps performs the color conversion
+            uint yuvi[6];
+            float red[2], green[2], blue[2];
+           
+            yuvi[0] =  (yuv101010Pel[0] &   COLOR_COMPONENT_MASK    );	
+            yuvi[1] = ((yuv101010Pel[0] >>  COLOR_COMPONENT_BIT_SIZE)       & COLOR_COMPONENT_MASK); 
+            yuvi[2] = ((yuv101010Pel[0] >> (COLOR_COMPONENT_BIT_SIZE << 1)) & COLOR_COMPONENT_MASK);
+
+            yuvi[3] =  (yuv101010Pel[1] &   COLOR_COMPONENT_MASK    );	
+            yuvi[4] = ((yuv101010Pel[1] >>  COLOR_COMPONENT_BIT_SIZE)       & COLOR_COMPONENT_MASK); 
+            yuvi[5] = ((yuv101010Pel[1] >> (COLOR_COMPONENT_BIT_SIZE << 1)) & COLOR_COMPONENT_MASK);
+
+            // YUV to RGB Transformation conversion
+            YUV2RGB(&yuvi[0], &red[0], &green[0], &blue[0]);
+            YUV2RGB(&yuvi[3], &red[1], &green[1], &blue[1]);
+
+            // Clamp the results to RGBA
+           
+            const size_t dstImagePitch = nDestPitch >> 2;
+
+            dstImage[y * dstImagePitch + x     ] = RGBAPACK_10bit(red[0], green[0], blue[0], constAlpha);
+            dstImage[y * dstImagePitch + x + 1 ] = RGBAPACK_10bit(red[1], green[1], blue[1], constAlpha);
+        }
+
+        void NV12ToARGB_gpu(const PtrStepb decodedFrame, DevMem2D_<uint> interopFrame, cudaStream_t stream)
+        {
+            dim3 block(32, 8);
+            dim3 grid(divUp(interopFrame.cols, 2 * block.x), divUp(interopFrame.rows, block.y)); 
+
+            NV12ToARGB<<<grid, block, 0, stream>>>(decodedFrame.data, decodedFrame.step, interopFrame.data, interopFrame.step, 
+                interopFrame.cols, interopFrame.rows);
+
+            cudaSafeCall( cudaGetLastError() );
+
+            if (stream == 0)
+                cudaSafeCall( cudaDeviceSynchronize() );
+        }
+    }
+}}}

modules/gpu/src/cuvid_video_source.cpp

+#include "cuvid_video_source.h"
+#include "cu_safe_call.h"
+
+#if defined(HAVE_CUDA) && !defined(__APPLE__)
+
+cv::gpu::detail::CuvidVideoSource::CuvidVideoSource(const std::string& fname)
+{
+    CUVIDSOURCEPARAMS params;
+    std::memset(&params, 0, sizeof(CUVIDSOURCEPARAMS));
+
+    // Fill parameter struct
+    params.pUserData = this;                        // will be passed to data handlers
+    params.pfnVideoDataHandler = HandleVideoData;   // our local video-handler callback
+    params.pfnAudioDataHandler = 0;
+
+    // now create the actual source
+    CUresult res = cuvidCreateVideoSource(&videoSource_, fname.c_str(), &params);
+    if (res == CUDA_ERROR_INVALID_SOURCE)
+        throw std::runtime_error("Unsupported video source");
+    cuSafeCall( res );
+
+    CUVIDEOFORMAT vidfmt;
+    cuSafeCall( cuvidGetSourceVideoFormat(videoSource_, &vidfmt, 0) );
+
+    format_.codec = static_cast<VideoReader_GPU::Codec>(vidfmt.codec);
+    format_.chromaFormat = static_cast<VideoReader_GPU::ChromaFormat>(vidfmt.chroma_format);
+    format_.width = vidfmt.coded_width;
+    format_.height = vidfmt.coded_height;
+}
+
+cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::detail::CuvidVideoSource::format() const
+{ 
+    return format_;
+}
+
+void cv::gpu::detail::CuvidVideoSource::start()
+{
+    cuSafeCall( cuvidSetVideoSourceState(videoSource_, cudaVideoState_Started) );
+}
+
+void cv::gpu::detail::CuvidVideoSource::stop()
+{
+    cuSafeCall( cuvidSetVideoSourceState(videoSource_, cudaVideoState_Stopped) );
+}
+
+bool cv::gpu::detail::CuvidVideoSource::isStarted() const
+{
+    return (cuvidGetVideoSourceState(videoSource_) == cudaVideoState_Started);
+}
+
+bool cv::gpu::detail::CuvidVideoSource::hasError() const
+{
+    return (cuvidGetVideoSourceState(videoSource_) == cudaVideoState_Error);
+}
+
+int CUDAAPI cv::gpu::detail::CuvidVideoSource::HandleVideoData(void* userData, CUVIDSOURCEDATAPACKET* packet)
+{
+    CuvidVideoSource* thiz = static_cast<CuvidVideoSource*>(userData);
+
+    return thiz->parseVideoData(packet->payload, packet->payload_size, (packet->flags & CUVID_PKT_ENDOFSTREAM) != 0);
+}
+
+#endif // defined(HAVE_CUDA) && !defined(__APPLE__)

modules/gpu/src/cuvid_video_source.h

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef __CUVUD_VIDEO_SOURCE_H__
+#define __CUVUD_VIDEO_SOURCE_H__
+
+#include "precomp.hpp"
+
+#if defined(HAVE_CUDA) && !defined(__APPLE__)
+
+namespace cv { namespace gpu 
+{
+    namespace detail
+    {
+        class CuvidVideoSource : public VideoReader_GPU::VideoSource
+        {
+        public:
+            explicit CuvidVideoSource(const std::string& fname);
+            ~CuvidVideoSource() { cuvidDestroyVideoSource(videoSource_); }
+
+            VideoReader_GPU::FormatInfo format() const;
+            void start();
+            void stop();
+            bool isStarted() const;
+            bool hasError() const;
+
+        private:
+            CuvidVideoSource(const CuvidVideoSource&);
+            CuvidVideoSource& operator =(const CuvidVideoSource&);
+
+            // Callback for handling packages of demuxed video data.
+            //
+            // Parameters:
+            //      pUserData - Pointer to user data. We must pass a pointer to a 
+            //          VideoSourceData struct here, that contains a valid CUvideoparser
+            //          and FrameQueue.
+            //      pPacket - video-source data packet. 
+            //
+            // NOTE: called from a different thread that doesn't not have a cuda context
+            //
+            static int CUDAAPI HandleVideoData(void* pUserData, CUVIDSOURCEDATAPACKET* pPacket);
+
+            CUvideosource videoSource_;
+            VideoReader_GPU::FormatInfo format_;
+        };
+    }
+}}
+
+#endif // defined(HAVE_CUDA) && !defined(__APPLE__)
+
+#endif // __CUVUD_VIDEO_SOURCE_H__

modules/gpu/src/ffmpeg_video_source.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "ffmpeg_video_source.h"
+
+#ifdef HAVE_CUDA
+
+#ifdef HAVE_FFMPEG
+    #ifdef NEW_FFMPEG
+        #include "cap_ffmpeg_impl_v2.hpp"
+    #else
+        #include "cap_ffmpeg_impl.hpp"
+    #endif
+#else
+    #include "cap_ffmpeg_api.hpp"
+#endif
+
+namespace
+{
+    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;
+
+    bool init_MediaStream_FFMPEG()
+    {
+        static bool initialized = 0;
+
+        if (!initialized)
+        {
+            #if defined WIN32 || defined _WIN32
+                const char* module_name = "opencv_ffmpeg"
+                #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;
+    }
+}
+
+cv::gpu::detail::FFmpegVideoSource::FFmpegVideoSource(const std::string& fname) :
+    stream_(0)
+{
+    CV_Assert( init_MediaStream_FFMPEG() );
+
+    int codec;
+    int chroma_format;
+    int width;
+    int height;
+
+    stream_ = create_InputMediaStream_FFMPEG_p(fname.c_str(), &codec, &chroma_format, &width, &height);
+    if (!stream_)
+        CV_Error(CV_StsUnsupportedFormat, "Unsupported video source");
+
+    format_.codec = static_cast<VideoReader_GPU::Codec>(codec);
+    format_.chromaFormat = static_cast<VideoReader_GPU::ChromaFormat>(chroma_format);
+    format_.width = width;
+    format_.height = height;
+}
+
+cv::gpu::detail::FFmpegVideoSource::~FFmpegVideoSource()
+{
+    release_InputMediaStream_FFMPEG_p(stream_);
+}
+
+cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::detail::FFmpegVideoSource::format() const
+{
+    return format_;
+}
+
+void cv::gpu::detail::FFmpegVideoSource::start()
+{
+    stop_ = false;
+    hasError_ = false;
+    thread_.reset(new Thread(readLoop, this));
+}
+
+void cv::gpu::detail::FFmpegVideoSource::stop()
+{
+    stop_ = true;
+    thread_->wait();
+    thread_.reset();
+}
+
+bool cv::gpu::detail::FFmpegVideoSource::isStarted() const
+{
+    return !stop_;
+}
+
+bool cv::gpu::detail::FFmpegVideoSource::hasError() const
+{
+    return hasError_;
+}
+
+void cv::gpu::detail::FFmpegVideoSource::readLoop(void* userData)
+{
+    FFmpegVideoSource* thiz = static_cast<FFmpegVideoSource*>(userData);
+
+    for (;;)
+    {
+        unsigned char* data;
+        int size;
+        int endOfFile;
+
+        if (!read_InputMediaStream_FFMPEG_p(thiz->stream_, &data, &size, &endOfFile))
+        {
+            thiz->hasError_ = !endOfFile;
+            break;
+        }
+
+        if (!thiz->parseVideoData(data, size))
+        {
+            thiz->hasError_ = true;
+            break;
+        }
+
+        if (thiz->stop_)
+            break;
+    }
+
+    thiz->parseVideoData(0, 0, true);
+}
+
+#endif // HAVE_CUDA

modules/gpu/src/ffmpeg_video_source.h

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef __FFMPEG_VIDEO_SOURCE_H__
+#define __FFMPEG_VIDEO_SOURCE_H__
+
+#include "precomp.hpp"
+#include "thread_wrappers.h"
+
+#ifdef HAVE_CUDA
+
+struct InputMediaStream_FFMPEG;
+
+namespace cv { namespace gpu
+{
+    namespace detail
+    {
+        class FFmpegVideoSource : public VideoReader_GPU::VideoSource
+        {
+        public:
+            FFmpegVideoSource(const std::string& fname);
+            ~FFmpegVideoSource();
+
+            VideoReader_GPU::FormatInfo format() const;
+            void start();
+            void stop();
+            bool isStarted() const;
+            bool hasError() const;
+
+        private:
+            FFmpegVideoSource(const FFmpegVideoSource&);
+            FFmpegVideoSource& operator =(const FFmpegVideoSource&);
+
+            VideoReader_GPU::FormatInfo format_;
+
+            InputMediaStream_FFMPEG* stream_;
+
+            std::auto_ptr<Thread> thread_;
+            volatile bool stop_;
+            volatile bool hasError_;
+
+            static void readLoop(void* userData);
+        };
+    }
+}}
+
+#endif // HAVE_CUDA
+
+#endif // __CUVUD_VIDEO_SOURCE_H__

modules/gpu/src/frame_queue.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "frame_queue.h"
+
+cv::gpu::detail::FrameQueue::FrameQueue() :
+    endOfDecode_(0),
+    framesInQueue_(0),
+    readPosition_(0)
+{
+    std::memset(displayQueue_, 0, sizeof(displayQueue_));
+    std::memset((void*)isFrameInUse_, 0, sizeof(isFrameInUse_));
+}
+
+bool cv::gpu::detail::FrameQueue::waitUntilFrameAvailable(int pictureIndex)
+{
+    while (isInUse(pictureIndex))
+    {
+        // Decoder is getting too far ahead from display
+        Thread::sleep(1);
+
+        if (isEndOfDecode())
+            return false;
+    }
+
+    return true;
+}
+
+void cv::gpu::detail::FrameQueue::enqueue(const CUVIDPARSERDISPINFO* picParams)
+{
+    // Mark the frame as 'in-use' so we don't re-use it for decoding until it is no longer needed
+    // for display
+    isFrameInUse_[picParams->picture_index] = true;
+
+    // Wait until we have a free entry in the display queue (should never block if we have enough entries)
+    do
+    {
+        bool isFramePlaced = false;
+
+        {
+            CriticalSection::AutoLock autoLock(criticalSection_);
+
+            if (framesInQueue_ < MaximumSize)
+            {
+                int writePosition = (readPosition_ + framesInQueue_) % MaximumSize;
+                displayQueue_[writePosition] = *picParams;
+                framesInQueue_++;
+                isFramePlaced = true;
+            }
+        }
+
+        if (isFramePlaced) // Done
+            break;
+
+        // Wait a bit
+        Thread::sleep(1);
+    } while (!isEndOfDecode());
+}
+
+bool cv::gpu::detail::FrameQueue::dequeue(CUVIDPARSERDISPINFO& displayInfo)
+{
+    CriticalSection::AutoLock autoLock(criticalSection_);
+
+    if (framesInQueue_ > 0)
+    {
+        int entry = readPosition_;
+        displayInfo = displayQueue_[entry];
+        readPosition_ = (entry + 1) % MaximumSize;
+        framesInQueue_--;
+        return true;
+    }
+
+    return false;
+}

modules/gpu/src/frame_queue.h

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef __FRAME_QUEUE_H__
+#define __FRAME_QUEUE_H__
+
+#include "precomp.hpp"
+#include "thread_wrappers.h"
+
+#ifdef HAVE_CUDA
+
+namespace cv { namespace gpu
+{
+    namespace detail
+    {
+        class FrameQueue
+        {
+        public:
+            static const int MaximumSize = 20; // MAX_FRM_CNT;
+
+            FrameQueue();
+
+            void endDecode() { endOfDecode_ = true; }
+            bool isEndOfDecode() const { return endOfDecode_ != 0;}
+
+            // Spins until frame becomes available or decoding gets canceled. 
+            // If the requested frame is available the method returns true.
+            // If decoding was interupted before the requested frame becomes
+            // available, the method returns false.
+            bool waitUntilFrameAvailable(int pictureIndex);
+
+            void enqueue(const CUVIDPARSERDISPINFO* picParams);
+
+            // Deque the next frame.
+            // Parameters:
+            //      displayInfo - New frame info gets placed into this object.
+            // Returns:
+            //      true, if a new frame was returned,
+            //      false, if the queue was empty and no new frame could be returned. 
+            bool dequeue(CUVIDPARSERDISPINFO& displayInfo);
+
+            void releaseFrame(const CUVIDPARSERDISPINFO& picParams) { isFrameInUse_[picParams.picture_index] = false; }
+
+        private:
+            FrameQueue(const FrameQueue&);
+            FrameQueue& operator =(const FrameQueue&);
+
+            bool isInUse(int pictureIndex) const { return isFrameInUse_[pictureIndex] != 0; }
+
+            CriticalSection criticalSection_;
+
+            volatile int isFrameInUse_[MaximumSize];
+            volatile int endOfDecode_;
+
+            int framesInQueue_;
+            int readPosition_;
+            CUVIDPARSERDISPINFO displayQueue_[MaximumSize];
+        };
+    }
+}}
+
+#endif // HAVE_CUDA
+
+#endif // __FRAME_QUEUE_H__

modules/gpu/src/precomp.hpp

@@ -51,6 +51,7 @@
     #include "cvconfig.h"
 #endif
 
+#include <cstring>
 #include <iostream>
 #include <limits>
 #include <vector>
@@ -60,6 +61,10 @@
 #include <iterator>
 #include <functional>
 #include <utility>
+#include <deque>
+#include <stdexcept>
+#include <memory>
+#include <string>
 
 #include "opencv2/gpu/gpu.hpp"
 #include "opencv2/imgproc/imgproc.hpp"

modules/gpu/src/thread_wrappers.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "thread_wrappers.h"
+
+#ifdef HAVE_CUDA
+
+#ifdef WIN32
+    #define NOMINMAX
+    #include <windows.h>
+#else
+    #include <pthread.h>
+    #include <unistd.h>
+#endif
+
+#ifdef WIN32
+    class cv::gpu::detail::CriticalSection::Impl
+    {
+    public:
+        Impl()
+        {
+            InitializeCriticalSection(&criticalSection_);
+        }
+
+        ~Impl()
+        {
+            DeleteCriticalSection(&criticalSection_);
+        }
+
+        void enter()
+        {
+            EnterCriticalSection(&criticalSection_);
+        }
+
+        void leave()
+        {
+            LeaveCriticalSection(&criticalSection_);
+        }
+
+    private:
+        CRITICAL_SECTION criticalSection_;
+    };
+#else
+    class cv::gpu::detail::CriticalSection::Impl
+    {
+    public:
+        Impl()
+        {
+            pthread_mutexattr_t mutex_attribute;
+            pthread_mutexattr_init(&mutex_attribute);
+            pthread_mutexattr_settype(&mutex_attribute, PTHREAD_MUTEX_RECURSIVE);
+            pthread_mutex_init(&mutex_, 0);
+            pthread_mutexattr_destroy(&mutex_attribute);
+        }
+
+        ~Impl()
+        {
+            pthread_mutex_destroy(&mutex_);
+        }
+
+        void enter()
+        {
+            pthread_mutex_lock(&mutex_);
+        }
+
+        void leave()
+        {
+            pthread_mutex_unlock(&mutex_);
+        }
+
+    private:
+        pthread_mutex_t mutex_;
+    };
+#endif
+
+cv::gpu::detail::CriticalSection::CriticalSection() :
+    impl_(new Impl)
+{
+}
+
+cv::gpu::detail::CriticalSection::~CriticalSection()
+{
+}
+
+void cv::gpu::detail::CriticalSection::enter()
+{
+    impl_->enter();
+}
+
+void cv::gpu::detail::CriticalSection::leave()
+{
+    impl_->leave();
+}
+
+#ifdef WIN32
+    namespace
+    {
+        struct UserData
+        {
+            void (*func)(void* userData);
+            void* param;
+        };
+
+        DWORD WINAPI WinThreadFunction(LPVOID lpParam)
+        {
+            UserData* userData = static_cast<UserData*>(lpParam);
+
+            userData->func(userData->param);
+
+            return 0;
+        }
+    }
+
+    class cv::gpu::detail::Thread::Impl
+    {
+    public:
+        Impl(void (*func)(void* userData), void* userData)
+        {
+            userData_.func = func;
+            userData_.param = userData;
+
+            thread_ = CreateThread(
+                NULL,                   // default security attributes
+                0,                      // use default stack size
+                WinThreadFunction,      // thread function name
+                &userData_,             // argument to thread function
+                0,                      // use default creation flags
+                &threadId_);            // returns the thread identifier
+        }
+
+        ~Impl()
+        {
+            CloseHandle(thread_);
+        }
+
+        void wait()
+        {
+            WaitForSingleObject(thread_, INFINITE);
+        }
+
+    private:
+        UserData userData_;
+        HANDLE thread_;
+        DWORD threadId_;
+    };
+#else
+    namespace
+    {
+        struct UserData
+        {
+            void (*func)(void* userData);
+            void* param;
+        };
+
+        void* PThreadFunction(void* lpParam)
+        {
+            UserData* userData = static_cast<UserData*>(lpParam);
+
+            userData->func(userData->param);
+
+            return 0;
+        }
+    }
+
+    class cv::gpu::detail::Thread::Impl
+    {
+    public:
+        Impl(void (*func)(void* userData), void* userData)
+        {
+            userData_.func = func;
+            userData_.param = userData;
+
+            pthread_create(&thread_, NULL, PThreadFunction, &userData_);
+        }
+
+        ~Impl()
+        {
+            pthread_detach(thread_);
+        }
+
+        void wait()
+        {
+            pthread_join(thread_, NULL);
+        }
+
+    private:
+        pthread_t thread_;
+        UserData userData_;
+    };
+#endif
+
+cv::gpu::detail::Thread::Thread(void (*func)(void* userData), void* userData) :
+    impl_(new Impl(func, userData))
+{
+}
+
+cv::gpu::detail::Thread::~Thread()
+{
+}
+
+void cv::gpu::detail::Thread::wait()
+{
+    impl_->wait();
+}
+
+void cv::gpu::detail::Thread::sleep(int ms)
+{
+#ifdef WIN32
+    ::Sleep(ms);
+#else
+    ::usleep(ms * 1000);
+#endif
+}
+
+#endif // HAVE_CUDA

modules/gpu/src/thread_wrappers.h

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef __THREAD_WRAPPERS_H__
+#define __THREAD_WRAPPERS_H__
+
+#include "precomp.hpp"
+
+#ifdef HAVE_CUDA
+
+namespace cv { namespace gpu 
+{
+    namespace detail
+    {
+        class CriticalSection
+        {
+        public:
+            CriticalSection();
+            ~CriticalSection();
+
+            void enter();
+            void leave();
+
+            class AutoLock
+            {
+            public:
+                explicit AutoLock(CriticalSection& criticalSection) :
+                    criticalSection_(criticalSection)
+                {
+                    criticalSection_.enter();
+                }
+
+                ~AutoLock()
+                {
+                    criticalSection_.leave();
+                }
+
+            private:
+                CriticalSection& criticalSection_;
+            };
+
+        private:
+            CriticalSection(const CriticalSection&);
+            CriticalSection& operator=(const CriticalSection&);
+
+            class Impl;
+            std::auto_ptr<Impl> impl_;
+        };
+
+        class Thread
+        {
+        public:
+            explicit Thread(void (*func)(void* userData), void* userData = 0);
+            ~Thread();
+
+            void wait();
+
+            static void sleep(int ms);
+
+        private:
+            Thread(const Thread&);
+            Thread& operator=(const Thread&);
+
+            class Impl;
+            std::auto_ptr<Impl> impl_;
+        };
+
+    }
+}}
+
+#endif // HAVE_CUDA
+
+#endif // __THREAD_WRAPPERS_H__

modules/gpu/src/video_decoder.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "video_decoder.h"
+#include "frame_queue.h"
+
+#ifdef HAVE_CUDA
+
+void cv::gpu::detail::VideoDecoder::create(const VideoReader_GPU::FormatInfo& videoFormat)
+{
+    release();
+
+    cudaVideoCodec codec = static_cast<cudaVideoCodec>(videoFormat.codec);
+    cudaVideoChromaFormat chromaFormat = static_cast<cudaVideoChromaFormat>(videoFormat.chromaFormat);
+
+    cudaVideoCreateFlags videoCreateFlags = (codec == cudaVideoCodec_JPEG || codec == cudaVideoCodec_MPEG2) ?
+                                            cudaVideoCreate_PreferCUDA :
+                                            cudaVideoCreate_PreferCUVID;
+
+    // Validate video format.  These are the currently supported formats via NVCUVID
+    CV_Assert(cudaVideoCodec_MPEG1 == codec ||
+              cudaVideoCodec_MPEG2 == codec ||
+              cudaVideoCodec_MPEG4 == codec ||
+              cudaVideoCodec_VC1   == codec ||
+              cudaVideoCodec_H264  == codec ||
+              cudaVideoCodec_JPEG  == codec ||
+              cudaVideoCodec_YUV420== codec ||
+              cudaVideoCodec_YV12  == codec ||
+              cudaVideoCodec_NV12  == codec ||
+              cudaVideoCodec_YUYV  == codec ||
+              cudaVideoCodec_UYVY  == codec );
+
+    CV_Assert(cudaVideoChromaFormat_Monochrome == chromaFormat ||
+              cudaVideoChromaFormat_420        == chromaFormat ||
+              cudaVideoChromaFormat_422        == chromaFormat ||
+              cudaVideoChromaFormat_444        == chromaFormat);
+
+    // Fill the decoder-create-info struct from the given video-format struct.
+    std::memset(&createInfo_, 0, sizeof(CUVIDDECODECREATEINFO));
+
+    // Create video decoder
+    createInfo_.CodecType           = codec;
+    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;
+
+    // No scaling
+    static const int MAX_FRAME_COUNT = 2;
+
+    createInfo_.ulTargetWidth       = createInfo_.ulWidth;
+    createInfo_.ulTargetHeight      = createInfo_.ulHeight;
+    createInfo_.ulNumOutputSurfaces = MAX_FRAME_COUNT;  // We won't simultaneously map more than 8 surfaces
+    createInfo_.ulCreationFlags     = videoCreateFlags;
+    createInfo_.vidLock = lock_;
+
+    // create the decoder
+    cuSafeCall( cuvidCreateDecoder(&decoder_, &createInfo_) );
+}
+
+void cv::gpu::detail::VideoDecoder::release()
+{
+    if (decoder_)
+    {
+        cuvidDestroyDecoder(decoder_);
+        decoder_ = 0;
+    }
+}
+
+#endif // HAVE_CUDA

modules/gpu/src/video_decoder.h

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef __VIDEO_DECODER_H__
+#define __VIDEO_DECODER_H__
+
+#include "precomp.hpp"
+#include "cu_safe_call.h"
+
+#ifdef HAVE_CUDA
+
+namespace cv { namespace gpu 
+{
+    namespace detail
+    {
+        class VideoDecoder
+        {
+        public:
+            VideoDecoder(const VideoReader_GPU::FormatInfo& videoFormat, CUvideoctxlock lock) : lock_(lock), decoder_(0)
+            {
+                create(videoFormat);
+            }
+
+            ~VideoDecoder()
+            {
+                release();
+            }
+
+            void create(const VideoReader_GPU::FormatInfo& videoFormat);
+            void release();
+
+            // Get the code-type currently used.
+            cudaVideoCodec codec() const { return createInfo_.CodecType; }
+            unsigned long maxDecodeSurfaces() const { return createInfo_.ulNumDecodeSurfaces; }
+
+            unsigned long frameWidth() const { return createInfo_.ulWidth; }
+            unsigned long frameHeight() const { return createInfo_.ulHeight; }
+
+            unsigned long targetWidth() const { return createInfo_.ulTargetWidth; }
+            unsigned long targetHeight() const { return createInfo_.ulTargetHeight; }
+
+            cudaVideoChromaFormat chromaFormat() const { return createInfo_.ChromaFormat; }
+
+            bool decodePicture(CUVIDPICPARAMS* picParams)
+            {
+                return cuvidDecodePicture(decoder_, picParams) == CUDA_SUCCESS;
+            }
+
+            cv::gpu::GpuMat mapFrame(int picIdx, CUVIDPROCPARAMS& videoProcParams)
+            {
+                CUdeviceptr ptr;
+                unsigned int pitch;
+
+                cuSafeCall( cuvidMapVideoFrame(decoder_, picIdx, &ptr, &pitch, &videoProcParams) );
+
+                return GpuMat(targetHeight() * 3 / 2, targetWidth(), CV_8UC1, (void*) ptr, pitch);
+            }
+
+            void unmapFrame(cv::gpu::GpuMat& frame)
+            {
+                cuSafeCall( cuvidUnmapVideoFrame(decoder_, (CUdeviceptr) frame.data) );
+                frame.release();
+            }
+
+        private:
+            VideoDecoder(const VideoDecoder&);
+            VideoDecoder& operator =(const VideoDecoder&);
+
+            CUvideoctxlock lock_;
+            CUVIDDECODECREATEINFO createInfo_;
+            CUvideodecoder        decoder_;
+        };
+    }
+}}
+
+#endif // HAVE_CUDA
+
+#endif // __VIDEO_DECODER_H__

modules/gpu/src/video_parser.cpp

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "video_parser.h"
+#include "cu_safe_call.h"
+
+#ifdef HAVE_CUDA
+
+cv::gpu::detail::VideoParser::VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue) :
+    videoDecoder_(videoDecoder), frameQueue_(frameQueue), unparsedPackets_(0), hasError_(false)
+{
+    CUVIDPARSERPARAMS params;
+    memset(&params, 0, sizeof(CUVIDPARSERPARAMS));
+
+    params.CodecType              = videoDecoder->codec();
+    params.ulMaxNumDecodeSurfaces = videoDecoder->maxDecodeSurfaces();
+    params.ulMaxDisplayDelay      = 1; // this flag is needed so the parser will push frames out to the decoder as quickly as it can
+    params.pUserData              = this;
+    params.pfnSequenceCallback    = HandleVideoSequence;    // Called before decoding frames and/or whenever there is a format change
+    params.pfnDecodePicture       = HandlePictureDecode;    // Called when a picture is ready to be decoded (decode order)
+    params.pfnDisplayPicture      = HandlePictureDisplay;   // Called whenever a picture is ready to be displayed (display order)
+
+    cuSafeCall( cuvidCreateVideoParser(&parser_, &params) );
+}
+
+bool cv::gpu::detail::VideoParser::parseVideoData(const unsigned char* data, size_t size, bool endOfStream)
+{
+    CUVIDSOURCEDATAPACKET packet;
+    std::memset(&packet, 0, sizeof(CUVIDSOURCEDATAPACKET));
+
+    if (endOfStream)
+        packet.flags |= CUVID_PKT_ENDOFSTREAM;
+
+    packet.payload_size = size;
+    packet.payload = data;
+
+    if (cuvidParseVideoData(parser_, &packet) != CUDA_SUCCESS)
+    {
+        hasError_ = true;
+        frameQueue_->endDecode();
+        return false;
+    }
+
+    const int maxUnparsedPackets = 15;
+
+    ++unparsedPackets_;
+    if (unparsedPackets_ > maxUnparsedPackets)
+    {
+        hasError_ = true;
+        frameQueue_->endDecode();
+        return false;
+    }
+
+    if (endOfStream)
+        frameQueue_->endDecode();
+
+    return !frameQueue_->isEndOfDecode();
+}
+
+int CUDAAPI cv::gpu::detail::VideoParser::HandleVideoSequence(void* userData, CUVIDEOFORMAT* format)
+{
+    VideoParser* thiz = static_cast<VideoParser*>(userData);
+
+    thiz->unparsedPackets_ = 0;
+
+    if (format->codec         != thiz->videoDecoder_->codec()       ||
+        format->coded_width   != thiz->videoDecoder_->frameWidth()  ||
+        format->coded_height  != thiz->videoDecoder_->frameHeight() ||
+        format->chroma_format != thiz->videoDecoder_->chromaFormat())
+    {
+        VideoReader_GPU::FormatInfo newFormat;
+
+        newFormat.codec = static_cast<VideoReader_GPU::Codec>(format->codec);
+        newFormat.chromaFormat = static_cast<VideoReader_GPU::ChromaFormat>(format->chroma_format);
+        newFormat.width = format->coded_width;
+        newFormat.height = format->coded_height;
+
+        try
+        {
+            thiz->videoDecoder_->create(newFormat);
+        }
+        catch (const cv::Exception&)
+        {
+            thiz->hasError_ = true;
+            return false;
+        }
+    }
+
+    return true;
+}
+
+int CUDAAPI cv::gpu::detail::VideoParser::HandlePictureDecode(void* userData, CUVIDPICPARAMS* picParams)
+{
+    VideoParser* thiz = static_cast<VideoParser*>(userData);
+
+    thiz->unparsedPackets_ = 0;
+
+    bool isFrameAvailable = thiz->frameQueue_->waitUntilFrameAvailable(picParams->CurrPicIdx);
+
+    if (!isFrameAvailable)
+        return false;
+
+    if (!thiz->videoDecoder_->decodePicture(picParams))
+    {
+        thiz->hasError_ = true;
+        return false;
+    }
+
+    return true;
+}
+
+int CUDAAPI cv::gpu::detail::VideoParser::HandlePictureDisplay(void* userData, CUVIDPARSERDISPINFO* picParams)
+{
+    VideoParser* thiz = static_cast<VideoParser*>(userData);
+
+    thiz->unparsedPackets_ = 0;
+
+    thiz->frameQueue_->enqueue(picParams);
+
+    return true;
+}
+
+#endif // HAVE_CUDA

modules/gpu/src/video_parser.h

+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                          License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef __VIDEO_PARSER_H__
+#define __VIDEO_PARSER_H__
+
+#include "precomp.hpp"
+
+#include "frame_queue.h"
+#include "video_decoder.h"
+
+#ifdef HAVE_CUDA
+
+namespace cv { namespace gpu 
+{
+    namespace detail
+    {
+        class VideoParser
+        {
+        public:
+            VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue);
+
+            ~VideoParser()
+            {
+                cuvidDestroyVideoParser(parser_);
+            }
+
+            bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream);
+
+            bool hasError() const { return hasError_; }
+
+        private:
+            VideoParser(const VideoParser&);
+            VideoParser& operator =(const VideoParser&);
+
+            VideoDecoder* videoDecoder_;
+            FrameQueue* frameQueue_;
+            CUvideoparser parser_;
+            int unparsedPackets_;
+            volatile bool hasError_;
+
+            // Called when the decoder encounters a video format change (or initial sequence header)
+            // This particular implementation of the callback returns 0 in case the video format changes
+            // to something different than the original format. Returning 0 causes a stop of the app.
+            static int CUDAAPI HandleVideoSequence(void* pUserData, CUVIDEOFORMAT* pFormat);
+
+            // Called by the video parser to decode a single picture
+            // Since the parser will deliver data as fast as it can, we need to make sure that the picture
+            // index we're attempting to use for decode is no longer used for display
+            static int CUDAAPI HandlePictureDecode(void* pUserData, CUVIDPICPARAMS* pPicParams);
+
+            // Called by the video parser to display a video frame (in the case of field pictures, there may be
+            // 2 decode calls per 1 display call, since two fields make up one frame)
+            static int CUDAAPI HandlePictureDisplay(void* pUserData, CUVIDPARSERDISPINFO* pPicParams);
+        };
+    }
+}}
+
+#endif // HAVE_CUDA
+
+#endif // __VIDEO_PARSER_H__

modules/gpu/src/video_writer.cpp

@@ -577,7 +577,7 @@ void cv::gpu::VideoWriter_GPU::Impl::write(const cv::gpu::GpuMat& frame, bool la
         switch (surfaceFormat_)
         {
         case UYVY: // UYVY (4:2:2)
-        case YUY2:	// YUY2 (4:2:2)
+        case YUY2: // YUY2 (4:2:2)
             copyUYVYorYUY2Frame(frameSize_, frame, videoFrame_);
             break;
 
@@ -662,33 +662,33 @@ namespace
 
         if (!initialized)
         {
-        #if defined WIN32 || defined _WIN32
-            const char* module_name = "opencv_ffmpeg"
-            #if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__)
-                "_64"
+            #if defined WIN32 || defined _WIN32
+                const char* module_name = "opencv_ffmpeg"
+                #if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__)
+                    "_64"
+                #endif
+                    ".dll";
+
+                static HMODULE cvFFOpenCV = LoadLibrary(module_name);
+
+                if (cvFFOpenCV)
+                {
+                    create_OutputMediaStream_FFMPEG_p =
+                        (Create_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "create_OutputMediaStream_FFMPEG");
+                    release_OutputMediaStream_FFMPEG_p =
+                        (Release_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "release_OutputMediaStream_FFMPEG");
+                    write_OutputMediaStream_FFMPEG_p =
+                        (Write_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "write_OutputMediaStream_FFMPEG");
+
+                    initialized = create_OutputMediaStream_FFMPEG_p != 0 && release_OutputMediaStream_FFMPEG_p != 0 && write_OutputMediaStream_FFMPEG_p != 0;
+                }
+            #elif defined HAVE_FFMPEG
+                create_OutputMediaStream_FFMPEG_p = create_OutputMediaStream_FFMPEG;
+                release_OutputMediaStream_FFMPEG_p = release_OutputMediaStream_FFMPEG;
+                write_OutputMediaStream_FFMPEG_p = write_OutputMediaStream_FFMPEG;
+
+                initialized = true;
             #endif
-                ".dll";
-
-            static HMODULE cvFFOpenCV = LoadLibrary(module_name);
-
-            if (cvFFOpenCV)
-            {
-                create_OutputMediaStream_FFMPEG_p =
-                    (Create_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "create_OutputMediaStream_FFMPEG");
-                release_OutputMediaStream_FFMPEG_p =
-                    (Release_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "release_OutputMediaStream_FFMPEG");
-                write_OutputMediaStream_FFMPEG_p =
-                    (Write_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "write_OutputMediaStream_FFMPEG");
-
-                initialized = create_OutputMediaStream_FFMPEG_p != 0 && release_OutputMediaStream_FFMPEG_p != 0 && write_OutputMediaStream_FFMPEG_p != 0;
-            }
-        #elif defined HAVE_FFMPEG
-            create_OutputMediaStream_FFMPEG_p = create_OutputMediaStream_FFMPEG;
-            release_OutputMediaStream_FFMPEG_p = release_OutputMediaStream_FFMPEG;
-            write_OutputMediaStream_FFMPEG_p = write_OutputMediaStream_FFMPEG;
-
-            initialized = true;
-        #endif
         }
 
         return initialized;

modules/gpu/test/test_video.cpp

@@ -387,6 +387,8 @@ INSTANTIATE_TEST_CASE_P(GPU_Video, FarnebackOpticalFlow, testing::Combine(
 /////////////////////////////////////////////////////////////////////////////////////////////////
 // VideoWriter
 
+#ifdef WIN32
+
 PARAM_TEST_CASE(VideoWriter, cv::gpu::DeviceInfo, std::string)
 {
     cv::gpu::DeviceInfo devInfo;
@@ -440,10 +442,58 @@ TEST_P(VideoWriter, Regression)
 
     reader.open(outputFile);
     ASSERT_TRUE( reader.isOpened() );
+
+    for (int i = 0; i < 5; ++i)
+    {
+        reader >> frame;
+        ASSERT_FALSE( frame.empty() );
+    }
 }
 
 INSTANTIATE_TEST_CASE_P(GPU_Video, VideoWriter, testing::Combine(
     ALL_DEVICES,
-    testing::Values("VID00003-20100701-2204.3GP", "big_buck_bunny.mpg")));
+    testing::Values(std::string("VID00003-20100701-2204.mpg"), std::string("big_buck_bunny.mpg"))));
+
+#endif // WIN32
+
+/////////////////////////////////////////////////////////////////////////////////////////////////
+// VideoReader
+
+PARAM_TEST_CASE(VideoReader, cv::gpu::DeviceInfo, std::string)
+{
+    cv::gpu::DeviceInfo devInfo;
+    std::string inputFile;
+
+    virtual void SetUp()
+    {
+        devInfo = GET_PARAM(0);
+        inputFile = GET_PARAM(1);
+
+        cv::gpu::setDevice(devInfo.deviceID());
+
+        inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + inputFile;
+    }
+};
+
+TEST_P(VideoReader, Regression)
+{
+    cv::gpu::VideoReader_GPU reader(inputFile);
+    ASSERT_TRUE( reader.isOpened() );
+
+    cv::gpu::GpuMat frame;
+
+    for (int i = 0; i < 5; ++i)
+    {
+        ASSERT_TRUE( reader.read(frame) );
+        ASSERT_FALSE( frame.empty() );
+    }
+
+    reader.close();
+    ASSERT_FALSE( reader.isOpened() );
+}
+
+INSTANTIATE_TEST_CASE_P(GPU_Video, VideoReader, testing::Combine(
+    ALL_DEVICES,
+    testing::Values(std::string("VID00003-20100701-2204.mpg"))));
 
 } // namespace

modules/highgui/src/cap_ffmpeg_api.hpp

@@ -77,6 +77,18 @@ typedef struct OutputMediaStream_FFMPEG* (*Create_OutputMediaStream_FFMPEG_Plugi
 typedef void (*Release_OutputMediaStream_FFMPEG_Plugin)(struct OutputMediaStream_FFMPEG* stream);
 typedef void (*Write_OutputMediaStream_FFMPEG_Plugin)(struct OutputMediaStream_FFMPEG* stream, unsigned char* data, int size, int keyFrame);
 
+/*
+ * For CUDA decoder
+ */
+
+OPENCV_FFMPEG_API struct InputMediaStream_FFMPEG* create_InputMediaStream_FFMPEG(const char* fileName, int* codec, int* chroma_format, int* width, int* height);
+OPENCV_FFMPEG_API void release_InputMediaStream_FFMPEG(struct InputMediaStream_FFMPEG* stream);
+OPENCV_FFMPEG_API int read_InputMediaStream_FFMPEG(struct InputMediaStream_FFMPEG* stream, unsigned char** data, int* size, int* endOfFile);
+
+typedef struct InputMediaStream_FFMPEG* (*Create_InputMediaStream_FFMPEG_Plugin)(const char* fileName, int* codec, int* chroma_format, int* width, int* height);
+typedef void (*Release_InputMediaStream_FFMPEG_Plugin)(struct InputMediaStream_FFMPEG* stream);
+typedef int (*Read_InputMediaStream_FFMPEG_Plugin)(struct InputMediaStream_FFMPEG* stream, unsigned char** data, int* size, int* endOfFile);
+
 #ifdef __cplusplus
 }
 #endif

modules/highgui/src/cap_ffmpeg_impl.hpp

@@ -1741,3 +1741,232 @@ void write_OutputMediaStream_FFMPEG(struct OutputMediaStream_FFMPEG* stream, uns
 {
     stream->write(data, size, keyFrame);
 }
+
+/*
+ * For CUDA decoder
+ */
+
+enum
+{
+    VideoCodec_MPEG1 = 0,
+    VideoCodec_MPEG2,
+    VideoCodec_MPEG4,
+    VideoCodec_VC1,
+    VideoCodec_H264,
+    VideoCodec_JPEG,
+    VideoCodec_H264_SVC,
+    VideoCodec_H264_MVC,
+
+    // Uncompressed YUV
+    VideoCodec_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V')),   // Y,U,V (4:2:0)
+    VideoCodec_YV12   = (('Y'<<24)|('V'<<16)|('1'<<8)|('2')),   // Y,V,U (4:2:0)
+    VideoCodec_NV12   = (('N'<<24)|('V'<<16)|('1'<<8)|('2')),   // Y,UV  (4:2:0)
+    VideoCodec_YUYV   = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V')),   // YUYV/YUY2 (4:2:2)
+    VideoCodec_UYVY   = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y')),   // UYVY (4:2:2)
+};
+
+enum
+{
+    VideoChromaFormat_Monochrome = 0,
+    VideoChromaFormat_YUV420,
+    VideoChromaFormat_YUV422,
+    VideoChromaFormat_YUV444,
+};
+
+struct InputMediaStream_FFMPEG
+{
+public:
+    bool open(const char* fileName, int* codec, int* chroma_format, int* width, int* height);
+    void close();
+
+    bool read(unsigned char** data, int* size, int* endOfFile);
+
+private:
+    InputMediaStream_FFMPEG(const InputMediaStream_FFMPEG&);
+    InputMediaStream_FFMPEG& operator =(const InputMediaStream_FFMPEG&);
+
+    AVFormatContext* ctx_;
+    int video_stream_id_;
+    AVPacket pkt_;
+};
+
+bool InputMediaStream_FFMPEG::open(const char* fileName, int* codec, int* chroma_format, int* width, int* height)
+{
+    int err;
+
+    ctx_ = 0;
+    video_stream_id_ = -1;
+    memset(&pkt_, 0, sizeof(AVPacket));
+
+    #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 13, 0)
+        avformat_network_init();
+    #endif
+
+    // register all codecs, demux and protocols
+    av_register_all();
+
+    av_log_set_level(AV_LOG_ERROR);
+
+    #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 6, 0)
+        err = avformat_open_input(&ctx_, fileName, 0, 0);
+    #else
+        err = av_open_input_file(&ctx_, fileName, 0, 0, 0);
+    #endif
+    if (err < 0)
+        return false;
+
+    #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 3, 0)
+        err = avformat_find_stream_info(ctx_, 0);
+    #else
+        err = av_find_stream_info(ctx_);
+    #endif
+    if (err < 0)
+        return false;
+
+    for (unsigned int i = 0; i < ctx_->nb_streams; ++i)
+    {
+        #if LIBAVFORMAT_BUILD > 4628
+            AVCodecContext *enc = ctx_->streams[i]->codec;
+        #else
+            AVCodecContext *enc = &ctx_->streams[i]->codec;
+        #endif
+
+        if (enc->codec_type == AVMEDIA_TYPE_VIDEO)
+        {
+            video_stream_id_ = static_cast<int>(i);
+
+            switch (enc->codec_id)
+            {
+            case CODEC_ID_MPEG1VIDEO:
+                *codec = ::VideoCodec_MPEG1;
+                break;
+
+            case CODEC_ID_MPEG2VIDEO:
+                *codec = ::VideoCodec_MPEG2;
+                break;
+
+            case CODEC_ID_MPEG4:
+                *codec = ::VideoCodec_MPEG4;
+                break;
+
+            case CODEC_ID_VC1:
+                *codec = ::VideoCodec_VC1;
+                break;
+
+            case CODEC_ID_H264:
+                *codec = ::VideoCodec_H264;
+                break;
+
+            default:
+                return false;
+            };
+
+            switch (enc->pix_fmt)
+            {
+            case PIX_FMT_YUV420P:
+                *chroma_format = ::VideoChromaFormat_YUV420;
+                break;
+
+            case PIX_FMT_YUV422P:
+                *chroma_format = ::VideoChromaFormat_YUV422;
+                break;
+
+            case PIX_FMT_YUV444P:
+                *chroma_format = ::VideoChromaFormat_YUV444;
+                break;
+
+            default:
+                return false;
+            }
+
+            *width = enc->coded_width;
+            *height = enc->coded_height;
+
+            break;
+        }
+    }
+
+    if (video_stream_id_ < 0)
+        return false;
+
+    av_init_packet(&pkt_);
+
+    return true;
+}
+
+void InputMediaStream_FFMPEG::close()
+{
+    if (ctx_)
+    {
+        #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 24, 2)
+            avformat_close_input(&ctx_);
+        #else
+            av_close_input_file(ctx_);
+        #endif
+    }
+
+    // free last packet if exist
+    if (pkt_.data)
+        av_free_packet(&pkt_);
+}
+
+bool InputMediaStream_FFMPEG::read(unsigned char** data, int* size, int* endOfFile)
+{
+    // free last packet if exist
+    if (pkt_.data)
+        av_free_packet(&pkt_);
+
+    // get the next frame
+    for (;;)
+    {
+        int ret = av_read_frame(ctx_, &pkt_);
+
+        if (ret == AVERROR(EAGAIN))
+            continue;
+
+        if (ret < 0)
+        {
+            if (ret == AVERROR_EOF)
+                *endOfFile = true;
+            return false;
+        }
+
+        if (pkt_.stream_index != video_stream_id_)
+        {
+            av_free_packet(&pkt_);
+            continue;
+        }
+
+        break;
+    }
+
+    *data = pkt_.data;
+    *size = pkt_.size;
+    *endOfFile = false;
+
+    return true;
+}
+
+InputMediaStream_FFMPEG* create_InputMediaStream_FFMPEG(const char* fileName, int* codec, int* chroma_format, int* width, int* height)
+{
+    InputMediaStream_FFMPEG* stream = (InputMediaStream_FFMPEG*) malloc(sizeof(InputMediaStream_FFMPEG));
+
+    if (stream && stream->open(fileName, codec, chroma_format, width, height))
+        return stream;
+
+    stream->close();
+    free(stream);
+
+    return 0;
+}
+
+void release_InputMediaStream_FFMPEG(InputMediaStream_FFMPEG* stream)
+{
+    stream->close();
+    free(stream);
+}
+
+int read_InputMediaStream_FFMPEG(InputMediaStream_FFMPEG* stream, unsigned char** data, int* size, int* endOfFile)
+{
+    return stream->read(data, size, endOfFile);
+}

modules/highgui/src/cap_ffmpeg_impl_v2.hpp

@@ -1907,3 +1907,232 @@ void write_OutputMediaStream_FFMPEG(struct OutputMediaStream_FFMPEG* stream, uns
 {
     stream->write(data, size, keyFrame);
 }
+
+/*
+ * For CUDA decoder
+ */
+
+enum
+{
+    VideoCodec_MPEG1 = 0,
+    VideoCodec_MPEG2,
+    VideoCodec_MPEG4,
+    VideoCodec_VC1,
+    VideoCodec_H264,
+    VideoCodec_JPEG,
+    VideoCodec_H264_SVC,
+    VideoCodec_H264_MVC,
+
+    // Uncompressed YUV
+    VideoCodec_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V')),   // Y,U,V (4:2:0)
+    VideoCodec_YV12   = (('Y'<<24)|('V'<<16)|('1'<<8)|('2')),   // Y,V,U (4:2:0)
+    VideoCodec_NV12   = (('N'<<24)|('V'<<16)|('1'<<8)|('2')),   // Y,UV  (4:2:0)
+    VideoCodec_YUYV   = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V')),   // YUYV/YUY2 (4:2:2)
+    VideoCodec_UYVY   = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y')),   // UYVY (4:2:2)
+};
+
+enum
+{
+    VideoChromaFormat_Monochrome = 0,
+    VideoChromaFormat_YUV420,
+    VideoChromaFormat_YUV422,
+    VideoChromaFormat_YUV444,
+};
+
+struct InputMediaStream_FFMPEG
+{
+public:
+    bool open(const char* fileName, int* codec, int* chroma_format, int* width, int* height);
+    void close();
+
+    bool read(unsigned char** data, int* size, int* endOfFile);
+
+private:
+    InputMediaStream_FFMPEG(const InputMediaStream_FFMPEG&);
+    InputMediaStream_FFMPEG& operator =(const InputMediaStream_FFMPEG&);
+
+    AVFormatContext* ctx_;
+    int video_stream_id_;
+    AVPacket pkt_;
+};
+
+bool InputMediaStream_FFMPEG::open(const char* fileName, int* codec, int* chroma_format, int* width, int* height)
+{
+    int err;
+
+    ctx_ = 0;
+    video_stream_id_ = -1;
+    memset(&pkt_, 0, sizeof(AVPacket));
+
+    #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 13, 0)
+        avformat_network_init();
+    #endif
+
+    // register all codecs, demux and protocols
+    av_register_all();
+
+    av_log_set_level(AV_LOG_ERROR);
+
+    #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 6, 0)
+        err = avformat_open_input(&ctx_, fileName, 0, 0);
+    #else
+        err = av_open_input_file(&ctx_, fileName, 0, 0, 0);
+    #endif
+    if (err < 0)
+        return false;
+
+    #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 3, 0)
+        err = avformat_find_stream_info(ctx_, 0);
+    #else
+        err = av_find_stream_info(ctx_);
+    #endif
+    if (err < 0)
+        return false;
+
+    for (unsigned int i = 0; i < ctx_->nb_streams; ++i)
+    {
+        #if LIBAVFORMAT_BUILD > 4628
+            AVCodecContext *enc = ctx_->streams[i]->codec;
+        #else
+            AVCodecContext *enc = &ctx_->streams[i]->codec;
+        #endif
+
+        if (enc->codec_type == AVMEDIA_TYPE_VIDEO)
+        {
+            video_stream_id_ = static_cast<int>(i);
+
+            switch (enc->codec_id)
+            {
+            case CODEC_ID_MPEG1VIDEO:
+                *codec = ::VideoCodec_MPEG1;
+                break;
+
+            case CODEC_ID_MPEG2VIDEO:
+                *codec = ::VideoCodec_MPEG2;
+                break;
+
+            case CODEC_ID_MPEG4:
+                *codec = ::VideoCodec_MPEG4;
+                break;
+
+            case CODEC_ID_VC1:
+                *codec = ::VideoCodec_VC1;
+                break;
+
+            case CODEC_ID_H264:
+                *codec = ::VideoCodec_H264;
+                break;
+
+            default:
+                return false;
+            };
+
+            switch (enc->pix_fmt)
+            {
+            case PIX_FMT_YUV420P:
+                *chroma_format = ::VideoChromaFormat_YUV420;
+                break;
+
+            case PIX_FMT_YUV422P:
+                *chroma_format = ::VideoChromaFormat_YUV422;
+                break;
+
+            case PIX_FMT_YUV444P:
+                *chroma_format = ::VideoChromaFormat_YUV444;
+                break;
+
+            default:
+                return false;
+            }
+
+            *width = enc->coded_width;
+            *height = enc->coded_height;
+
+            break;
+        }
+    }
+
+    if (video_stream_id_ < 0)
+        return false;
+
+    av_init_packet(&pkt_);
+
+    return true;
+}
+
+void InputMediaStream_FFMPEG::close()
+{
+    if (ctx_)
+    {
+        #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 24, 2)
+            avformat_close_input(&ctx_);
+        #else
+            av_close_input_file(ctx_);
+        #endif
+    }
+
+    // free last packet if exist
+    if (pkt_.data)
+        av_free_packet(&pkt_);
+}
+
+bool InputMediaStream_FFMPEG::read(unsigned char** data, int* size, int* endOfFile)
+{
+    // free last packet if exist
+    if (pkt_.data)
+        av_free_packet(&pkt_);
+
+    // get the next frame
+    for (;;)
+    {
+        int ret = av_read_frame(ctx_, &pkt_);
+
+        if (ret == AVERROR(EAGAIN))
+            continue;
+
+        if (ret < 0)
+        {
+            if (ret == AVERROR_EOF)
+                *endOfFile = true;
+            return false;
+        }
+
+        if (pkt_.stream_index != video_stream_id_)
+        {
+            av_free_packet(&pkt_);
+            continue;
+        }
+
+        break;
+    }
+
+    *data = pkt_.data;
+    *size = pkt_.size;
+    *endOfFile = false;
+
+    return true;
+}
+
+InputMediaStream_FFMPEG* create_InputMediaStream_FFMPEG(const char* fileName, int* codec, int* chroma_format, int* width, int* height)
+{
+    InputMediaStream_FFMPEG* stream = (InputMediaStream_FFMPEG*) malloc(sizeof(InputMediaStream_FFMPEG));
+
+    if (stream && stream->open(fileName, codec, chroma_format, width, height))
+        return stream;
+
+    stream->close();
+    free(stream);
+
+    return 0;
+}
+
+void release_InputMediaStream_FFMPEG(InputMediaStream_FFMPEG* stream)
+{
+    stream->close();
+    free(stream);
+}
+
+int read_InputMediaStream_FFMPEG(InputMediaStream_FFMPEG* stream, unsigned char** data, int* size, int* endOfFile)
+{
+    return stream->read(data, size, endOfFile);
+}

samples/gpu/brox_optical_flow.cpp

@@ -215,7 +215,7 @@ int main(int argc, const char* argv[])
             switch (key)
             {
             case 27:
-                break;
+                return 0;
 
             case 'A':
                 if (currentFrame > 0) 
@@ -243,8 +243,6 @@ int main(int argc, const char* argv[])
         cerr << "Unknow error" << endl;
         return -1;
     }
-
-    return 0;
 }
 
 template <typename T> inline T clamp (T x, T a, T b)

samples/gpu/video_reader.cpp

+#include <iostream>
+#include <string>
+#include <vector>
+#include <algorithm>
+#include <numeric>
+
+#include <opencv2/core/core.hpp>
+#include <opencv2/core/opengl_interop.hpp>
+#include <opencv2/gpu/gpu.hpp>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/contrib/contrib.hpp>
+
+int main(int argc, const char* argv[])
+{
+    if (argc != 2)
+        return -1;
+
+    const std::string fname(argv[1]);
+
+    cv::namedWindow("CPU", cv::WINDOW_NORMAL);
+    cv::namedWindow("GPU", cv::WINDOW_OPENGL);
+    cv::gpu::setGlDevice();
+
+    cv::Mat frame;
+    cv::VideoCapture reader(fname);
+
+    cv::gpu::GpuMat d_frame;
+    cv::gpu::VideoReader_GPU d_reader(fname);
+    d_reader.dumpFormat(std::cout);
+
+    cv::TickMeter tm;
+    std::vector<double> cpu_times;
+    std::vector<double> gpu_times;
+
+    for (;;)
+    {
+        tm.reset(); tm.start();
+        if (!reader.read(frame))
+            break;
+        tm.stop();
+        cpu_times.push_back(tm.getTimeMilli());
+
+        tm.reset(); tm.start();
+        if (!d_reader.read(d_frame))
+            break;
+        tm.stop();
+        gpu_times.push_back(tm.getTimeMilli());
+
+        cv::imshow("CPU", frame);
+        cv::imshow("GPU", frame);
+
+        if (cv::waitKey(3) > 0)
+            break;
+    }
+
+    if (!cpu_times.empty() && !gpu_times.empty())
+    {
+        std::cout << std::endl << "Results:" << std::endl;
+
+        std::sort(cpu_times.begin(), cpu_times.end());
+        std::sort(gpu_times.begin(), gpu_times.end());
+
+        double cpu_avg = std::accumulate(cpu_times.begin(), cpu_times.end(), 0.0) / cpu_times.size();
+        double gpu_avg = std::accumulate(gpu_times.begin(), gpu_times.end(), 0.0) / gpu_times.size();
+
+        std::cout << "CPU : Avg : " << cpu_avg << " ms FPS : " << 1000.0 / cpu_avg << std::endl;
+        std::cout << "GPU : Avg : " << gpu_avg << " ms FPS : " << 1000.0 / gpu_avg << std::endl;
+    }
+
+    return 0;
+}