if(WITH_NVCUVID)
find_cuda_helper_libs(nvcuvid)
+ if(WIN32)
+ find_cuda_helper_libs(nvcuvenc)
+ endif()
set(HAVE_NVCUVID 1)
endif()
ocv_cuda_compile(cuda_objs ${lib_cuda} ${ncv_cuda})
set(cuda_link_libs ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
-
- if(WITH_NVCUVID)
- set(cuda_link_libs ${cuda_link_libs} ${CUDA_nvcuvid_LIBRARY})
- endif()
-
- if(WIN32)
- find_cuda_helper_libs(nvcuvenc)
- set(cuda_link_libs ${cuda_link_libs} ${CUDA_nvcuvenc_LIBRARY})
- endif()
-
- if(WITH_FFMPEG)
- set(cuda_link_libs ${cuda_link_libs} ${HIGHGUI_LIBRARIES})
- endif()
else()
set(lib_cuda "")
set(cuda_objs "")
-gpu::VideoWriter_GPU
----------------------
-Video writer class.
-
-.. ocv:class:: gpu::VideoWriter_GPU
-
-The class uses H264 video codec.
-
-.. note:: Currently only Windows platform is supported.
-
-
-
-gpu::VideoWriter_GPU::VideoWriter_GPU
--------------------------------------
-Constructors.
-
-.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU()
-.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
-.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
-.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
-.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
-
- :param fileName: Name of the output video file. Only AVI file format is supported.
-
- :param frameSize: Size of the input video frames.
-
- :param fps: Framerate of the created video stream.
-
- :param params: Encoder parameters. See :ocv:struct:`gpu::VideoWriter_GPU::EncoderParams` .
-
- :param format: Surface format of input frames ( ``SF_UYVY`` , ``SF_YUY2`` , ``SF_YV12`` , ``SF_NV12`` , ``SF_IYUV`` , ``SF_BGR`` or ``SF_GRAY``). BGR or gray frames will be converted to YV12 format before encoding, frames with other formats will be used as is.
-
- :param encoderCallback: Callbacks for video encoder. See :ocv:class:`gpu::VideoWriter_GPU::EncoderCallBack` . Use it if you want to work with raw video stream.
-
-The constructors initialize video writer. FFMPEG is used to write videos. User can implement own multiplexing with :ocv:class:`gpu::VideoWriter_GPU::EncoderCallBack` .
-
-
-
-gpu::VideoWriter_GPU::open
---------------------------
-Initializes or reinitializes video writer.
-
-.. ocv:function:: void gpu::VideoWriter_GPU::open(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
-.. ocv:function:: void gpu::VideoWriter_GPU::open(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
-.. ocv:function:: void gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
-.. ocv:function:: void gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
-
-The method opens video writer. Parameters are the same as in the constructor :ocv:func:`gpu::VideoWriter_GPU::VideoWriter_GPU` . The method throws :ocv:class:`Exception` if error occurs.
-
-
-
-gpu::VideoWriter_GPU::isOpened
-------------------------------
-Returns true if video writer has been successfully initialized.
-
-.. ocv:function:: bool gpu::VideoWriter_GPU::isOpened() const
-
-
-
-gpu::VideoWriter_GPU::close
----------------------------
-Releases the video writer.
-
-.. ocv:function:: void gpu::VideoWriter_GPU::close()
-
-
-
-gpu::VideoWriter_GPU::write
----------------------------
-Writes the next video frame.
-
-.. ocv:function:: void gpu::VideoWriter_GPU::write(const cv::gpu::GpuMat& image, bool lastFrame = false)
-
- :param image: The written frame.
-
- :param lastFrame: Indicates that it is end of stream. The parameter can be ignored.
-
-The method write the specified image to video file. The image must have the same size and the same surface format as has been specified when opening the video writer.
-
-
-
-gpu::VideoWriter_GPU::EncoderParams
------------------------------------
-.. ocv:struct:: gpu::VideoWriter_GPU::EncoderParams
-
-Different parameters for CUDA video encoder. ::
-
- struct EncoderParams
- {
- int P_Interval; // NVVE_P_INTERVAL,
- int IDR_Period; // NVVE_IDR_PERIOD,
- int DynamicGOP; // NVVE_DYNAMIC_GOP,
- int RCType; // NVVE_RC_TYPE,
- int AvgBitrate; // NVVE_AVG_BITRATE,
- int PeakBitrate; // NVVE_PEAK_BITRATE,
- int QP_Level_Intra; // NVVE_QP_LEVEL_INTRA,
- int QP_Level_InterP; // NVVE_QP_LEVEL_INTER_P,
- int QP_Level_InterB; // NVVE_QP_LEVEL_INTER_B,
- int DeblockMode; // NVVE_DEBLOCK_MODE,
- int ProfileLevel; // NVVE_PROFILE_LEVEL,
- int ForceIntra; // NVVE_FORCE_INTRA,
- int ForceIDR; // NVVE_FORCE_IDR,
- int ClearStat; // NVVE_CLEAR_STAT,
- int DIMode; // NVVE_SET_DEINTERLACE,
- int Presets; // NVVE_PRESETS,
- int DisableCabac; // NVVE_DISABLE_CABAC,
- int NaluFramingType; // NVVE_CONFIGURE_NALU_FRAMING_TYPE
- int DisableSPSPPS; // NVVE_DISABLE_SPS_PPS
-
- EncoderParams();
- explicit EncoderParams(const String& configFile);
-
- void load(const String& configFile);
- void save(const String& configFile) const;
- };
-
-
-
-gpu::VideoWriter_GPU::EncoderParams::EncoderParams
---------------------------------------------------
-Constructors.
-
-.. ocv:function:: gpu::VideoWriter_GPU::EncoderParams::EncoderParams()
-.. ocv:function:: gpu::VideoWriter_GPU::EncoderParams::EncoderParams(const String& configFile)
-
- :param configFile: Config file name.
-
-Creates default parameters or reads parameters from config file.
-
-
-
-gpu::VideoWriter_GPU::EncoderParams::load
------------------------------------------
-Reads parameters from config file.
-
-.. ocv:function:: void gpu::VideoWriter_GPU::EncoderParams::load(const String& configFile)
-
- :param configFile: Config file name.
-
-
-
-gpu::VideoWriter_GPU::EncoderParams::save
------------------------------------------
-Saves parameters to config file.
-
-.. ocv:function:: void gpu::VideoWriter_GPU::EncoderParams::save(const String& configFile) const
-
- :param configFile: Config file name.
-
-
-
-gpu::VideoWriter_GPU::EncoderCallBack
--------------------------------------
-.. ocv:class:: gpu::VideoWriter_GPU::EncoderCallBack
-
-Callbacks for CUDA video encoder. ::
-
- class EncoderCallBack
- {
- public:
- enum PicType
- {
- IFRAME = 1,
- PFRAME = 2,
- BFRAME = 3
- };
-
- virtual ~EncoderCallBack() {}
-
- virtual unsigned char* acquireBitStream(int* bufferSize) = 0;
- virtual void releaseBitStream(unsigned char* data, int size) = 0;
- virtual void onBeginFrame(int frameNumber, PicType picType) = 0;
- virtual void onEndFrame(int frameNumber, PicType picType) = 0;
- };
-
-
-
-gpu::VideoWriter_GPU::EncoderCallBack::acquireBitStream
--------------------------------------------------------
-Callback function to signal the start of bitstream that is to be encoded.
-
-.. ocv:function:: virtual uchar* gpu::VideoWriter_GPU::EncoderCallBack::acquireBitStream(int* bufferSize) = 0
-
-Callback must allocate buffer for CUDA encoder and return pointer to it and it's size.
-
-
-
-gpu::VideoWriter_GPU::EncoderCallBack::releaseBitStream
--------------------------------------------------------
-Callback function to signal that the encoded bitstream is ready to be written to file.
-
-.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::releaseBitStream(unsigned char* data, int size) = 0
-
-
-
-gpu::VideoWriter_GPU::EncoderCallBack::onBeginFrame
----------------------------------------------------
-Callback function to signal that the encoding operation on the frame has started.
-
-.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::onBeginFrame(int frameNumber, PicType picType) = 0
-
- :param picType: Specify frame type (I-Frame, P-Frame or B-Frame).
-
-
-
-gpu::VideoWriter_GPU::EncoderCallBack::onEndFrame
--------------------------------------------------
-Callback function signals that the encoding operation on the frame has finished.
-
-.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::onEndFrame(int frameNumber, PicType picType) = 0
-
- :param picType: Specify frame type (I-Frame, P-Frame or B-Frame).
-
-
-
-gpu::VideoReader_GPU
---------------------
-Class for reading video from files.
-
-.. ocv:class:: gpu::VideoReader_GPU
-
-.. note:: Currently only Windows and Linux platforms are supported.
-
-
-
-gpu::VideoReader_GPU::Codec
----------------------------
-
-Video codecs supported by :ocv:class:`gpu::VideoReader_GPU` .
-
-.. ocv:enum:: gpu::VideoReader_GPU::Codec
-
- .. ocv:emember:: MPEG1 = 0
- .. ocv:emember:: MPEG2
- .. ocv:emember:: MPEG4
- .. ocv:emember:: VC1
- .. ocv:emember:: H264
- .. ocv:emember:: JPEG
- .. ocv:emember:: H264_SVC
- .. ocv:emember:: H264_MVC
-
- .. ocv:emember:: Uncompressed_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V'))
-
- Y,U,V (4:2:0)
-
- .. ocv:emember:: Uncompressed_YV12 = (('Y'<<24)|('V'<<16)|('1'<<8)|('2'))
-
- Y,V,U (4:2:0)
-
- .. ocv:emember:: Uncompressed_NV12 = (('N'<<24)|('V'<<16)|('1'<<8)|('2'))
-
- Y,UV (4:2:0)
-
- .. ocv:emember:: Uncompressed_YUYV = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V'))
-
- YUYV/YUY2 (4:2:2)
-
- .. ocv:emember:: Uncompressed_UYVY = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y'))
-
- UYVY (4:2:2)
-
-
-gpu::VideoReader_GPU::ChromaFormat
-----------------------------------
-
-Chroma formats supported by :ocv:class:`gpu::VideoReader_GPU` .
-
-.. ocv:enum:: gpu::VideoReader_GPU::ChromaFormat
-
- .. ocv:emember:: Monochrome = 0
- .. ocv:emember:: YUV420
- .. ocv:emember:: YUV422
- .. ocv:emember:: YUV444
-
-
-gpu::VideoReader_GPU::FormatInfo
---------------------------------
-.. ocv:struct:: gpu::VideoReader_GPU::FormatInfo
-
-Struct providing information about video file format. ::
-
- struct FormatInfo
- {
- Codec codec;
- ChromaFormat chromaFormat;
- int width;
- int height;
- };
-
-
-gpu::VideoReader_GPU::VideoReader_GPU
--------------------------------------
-Constructors.
-
-.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU()
-.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU(const String& filename)
-.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU(const cv::Ptr<VideoSource>& source)
-
- :param filename: Name of the input video file.
-
- :param source: Video file parser implemented by user.
-
-The constructors initialize video reader. FFMPEG is used to read videos. User can implement own demultiplexing with :ocv:class:`gpu::VideoReader_GPU::VideoSource` .
-
-
-
-gpu::VideoReader_GPU::open
---------------------------
-Initializes or reinitializes video reader.
-
-.. ocv:function:: void gpu::VideoReader_GPU::open(const String& filename)
-.. ocv:function:: void gpu::VideoReader_GPU::open(const cv::Ptr<VideoSource>& source)
-
-The method opens video reader. Parameters are the same as in the constructor :ocv:func:`gpu::VideoReader_GPU::VideoReader_GPU` . The method throws :ocv:class:`Exception` if error occurs.
-
-
-
-gpu::VideoReader_GPU::isOpened
-------------------------------
-Returns true if video reader has been successfully initialized.
-
-.. ocv:function:: bool gpu::VideoReader_GPU::isOpened() const
-
-
-
-gpu::VideoReader_GPU::close
----------------------------
-Releases the video reader.
-
-.. ocv:function:: void gpu::VideoReader_GPU::close()
-
-
-
-gpu::VideoReader_GPU::read
---------------------------
-Grabs, decodes and returns the next video frame.
-
-.. ocv:function:: bool gpu::VideoReader_GPU::read(GpuMat& image)
-
-If no frames has been grabbed (there are no more frames in video file), the methods return ``false`` . The method throws :ocv:class:`Exception` if error occurs.
-
-
-
-gpu::VideoReader_GPU::format
-----------------------------
-Returns information about video file format.
-
-.. ocv:function:: FormatInfo gpu::VideoReader_GPU::format() const
-
-The method throws :ocv:class:`Exception` if video reader wasn't initialized.
-
-
-
-gpu::VideoReader_GPU::dumpFormat
---------------------------------
-Dump information about video file format to specified stream.
-
-.. ocv:function:: void gpu::VideoReader_GPU::dumpFormat(std::ostream& st)
-
- :param st: Output stream.
-
-The method throws :ocv:class:`Exception` if video reader wasn't initialized.
-
-
-
-gpu::VideoReader_GPU::VideoSource
------------------------------------
-.. ocv:class:: gpu::VideoReader_GPU::VideoSource
-
-Interface for video demultiplexing. ::
-
- class VideoSource
- {
- public:
- VideoSource();
- 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;
-
- protected:
- bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream = false);
- };
-
-User can implement own demultiplexing by implementing this interface.
-
-
-
-gpu::VideoReader_GPU::VideoSource::format
------------------------------------------
-Returns information about video file format.
-
-.. ocv:function:: virtual FormatInfo gpu::VideoReader_GPU::VideoSource::format() const = 0
-
-
-
-gpu::VideoReader_GPU::VideoSource::start
-----------------------------------------
-Starts processing.
-
-.. ocv:function:: virtual void gpu::VideoReader_GPU::VideoSource::start() = 0
-
-Implementation must create own thread with video processing and call periodic :ocv:func:`gpu::VideoReader_GPU::VideoSource::parseVideoData` .
-
-
-
-gpu::VideoReader_GPU::VideoSource::stop
----------------------------------------
-Stops processing.
-
-.. ocv:function:: virtual void gpu::VideoReader_GPU::VideoSource::stop() = 0
-
-
-
-gpu::VideoReader_GPU::VideoSource::isStarted
---------------------------------------------
-Returns ``true`` if processing was successfully started.
-
-.. ocv:function:: virtual bool gpu::VideoReader_GPU::VideoSource::isStarted() const = 0
-
-
-
-gpu::VideoReader_GPU::VideoSource::hasError
--------------------------------------------
-Returns ``true`` if error occured during processing.
-
-.. ocv:function:: virtual bool gpu::VideoReader_GPU::VideoSource::hasError() const = 0
-
-
-
-gpu::VideoReader_GPU::VideoSource::parseVideoData
--------------------------------------------------
-Parse next video frame. Implementation must call this method after new frame was grabbed.
-
-.. ocv:function:: bool gpu::VideoReader_GPU::VideoSource::parseVideoData(const uchar* data, size_t size, bool endOfStream = false)
-
- :param data: Pointer to frame data. Can be ``NULL`` if ``endOfStream`` if ``true`` .
-
- :param size: Size in bytes of current frame.
-
- :param endOfStream: Indicates that it is end of stream.
-
-
-
.. [Brox2004] T. Brox, A. Bruhn, N. Papenberg, J. Weickert. *High accuracy optical flow estimation based on a theory for warping*. ECCV 2004.
.. [FGD2003] Liyuan Li, Weimin Huang, Irene Y.H. Gu, and Qi Tian. *Foreground Object Detection from Videos Containing Complex Background*. ACM MM2003 9p, 2003.
.. [MOG2001] P. KadewTraKuPong and R. Bowden. *An improved adaptive background mixture model for real-time tracking with shadow detection*. Proc. 2nd European Workshop on Advanced Video-Based Surveillance Systems, 2001
GpuMat buf_;
};
-////////////////////////////////// Video Encoding //////////////////////////////////
-
-// Works only under Windows
-// Supports olny H264 video codec and AVI files
-class CV_EXPORTS VideoWriter_GPU
-{
-public:
- struct EncoderParams;
-
- // Callbacks for video encoder, use it if you want to work with raw video stream
- class EncoderCallBack;
-
- enum SurfaceFormat
- {
- SF_UYVY = 0,
- SF_YUY2,
- SF_YV12,
- SF_NV12,
- SF_IYUV,
- SF_BGR,
- SF_GRAY = SF_BGR
- };
-
- VideoWriter_GPU();
- VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
- VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
- VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
- VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
- ~VideoWriter_GPU();
-
- // all methods throws cv::Exception if error occurs
- void open(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
- void open(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
- void open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
- void open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
-
- bool isOpened() const;
- void close();
-
- void write(const cv::gpu::GpuMat& image, bool lastFrame = false);
-
- struct CV_EXPORTS EncoderParams
- {
- int P_Interval; // NVVE_P_INTERVAL,
- int IDR_Period; // NVVE_IDR_PERIOD,
- int DynamicGOP; // NVVE_DYNAMIC_GOP,
- int RCType; // NVVE_RC_TYPE,
- int AvgBitrate; // NVVE_AVG_BITRATE,
- int PeakBitrate; // NVVE_PEAK_BITRATE,
- int QP_Level_Intra; // NVVE_QP_LEVEL_INTRA,
- int QP_Level_InterP; // NVVE_QP_LEVEL_INTER_P,
- int QP_Level_InterB; // NVVE_QP_LEVEL_INTER_B,
- int DeblockMode; // NVVE_DEBLOCK_MODE,
- int ProfileLevel; // NVVE_PROFILE_LEVEL,
- int ForceIntra; // NVVE_FORCE_INTRA,
- int ForceIDR; // NVVE_FORCE_IDR,
- int ClearStat; // NVVE_CLEAR_STAT,
- int DIMode; // NVVE_SET_DEINTERLACE,
- int Presets; // NVVE_PRESETS,
- int DisableCabac; // NVVE_DISABLE_CABAC,
- int NaluFramingType; // NVVE_CONFIGURE_NALU_FRAMING_TYPE
- int DisableSPSPPS; // NVVE_DISABLE_SPS_PPS
-
- EncoderParams();
- explicit EncoderParams(const String& configFile);
-
- void load(const String& configFile);
- void save(const String& configFile) const;
- };
-
- EncoderParams getParams() const;
-
- class CV_EXPORTS EncoderCallBack
- {
- public:
- enum PicType
- {
- IFRAME = 1,
- PFRAME = 2,
- BFRAME = 3
- };
-
- virtual ~EncoderCallBack() {}
-
- // callback function to signal the start of bitstream that is to be encoded
- // must return pointer to buffer
- virtual uchar* acquireBitStream(int* bufferSize) = 0;
-
- // callback function to signal that the encoded bitstream is ready to be written to file
- virtual void releaseBitStream(unsigned char* data, int size) = 0;
-
- // callback function to signal that the encoding operation on the frame has started
- virtual void onBeginFrame(int frameNumber, PicType picType) = 0;
-
- // callback function signals that the encoding operation on the frame has finished
- virtual void onEndFrame(int frameNumber, PicType picType) = 0;
- };
-
-private:
- VideoWriter_GPU(const VideoWriter_GPU&);
- VideoWriter_GPU& operator=(const VideoWriter_GPU&);
-
- class Impl;
- 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 String& filename);
- explicit VideoReader_GPU(const cv::Ptr<VideoSource>& source);
-
- ~VideoReader_GPU();
-
- void open(const 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 CV_EXPORTS 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 uchar* 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_;
-};
-
//! removes points (CV_32FC2, single row matrix) with zero mask value
CV_EXPORTS void compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask);
}
#endif
-
-//////////////////////////////////////////////////////
-// VideoReader
-
-#if defined(HAVE_NVCUVID) && BUILD_WITH_VIDEO_INPUT_SUPPORT
-
-PERF_TEST_P(Video, DISABLED_Video_VideoReader, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
-{
- declare.time(20);
-
- const string inputFile = perf::TestBase::getDataPath(GetParam());
-
- if (PERF_RUN_GPU())
- {
- cv::gpu::VideoReader_GPU d_reader(inputFile);
- ASSERT_TRUE( d_reader.isOpened() );
-
- cv::gpu::GpuMat frame;
-
- TEST_CYCLE_N(10) d_reader.read(frame);
-
- GPU_SANITY_CHECK(frame);
- }
- else
- {
- cv::VideoCapture reader(inputFile);
- ASSERT_TRUE( reader.isOpened() );
-
- cv::Mat frame;
-
- TEST_CYCLE_N(10) reader >> frame;
-
- CPU_SANITY_CHECK(frame);
- }
-}
-
-#endif
-
-//////////////////////////////////////////////////////
-// VideoWriter
-
-#if defined(HAVE_NVCUVID) && defined(WIN32)
-
-PERF_TEST_P(Video, DISABLED_Video_VideoWriter, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
-{
- declare.time(30);
-
- const string inputFile = perf::TestBase::getDataPath(GetParam());
- const string outputFile = cv::tempfile(".avi");
-
- const double FPS = 25.0;
-
- cv::VideoCapture reader(inputFile);
- ASSERT_TRUE( reader.isOpened() );
-
- cv::Mat frame;
-
- if (PERF_RUN_GPU())
- {
- cv::gpu::VideoWriter_GPU d_writer;
-
- cv::gpu::GpuMat d_frame;
-
- for (int i = 0; i < 10; ++i)
- {
- reader >> frame;
- ASSERT_FALSE(frame.empty());
-
- d_frame.upload(frame);
-
- if (!d_writer.isOpened())
- d_writer.open(outputFile, frame.size(), FPS);
-
- startTimer(); next();
- d_writer.write(d_frame);
- stopTimer();
- }
- }
- else
- {
- cv::VideoWriter writer;
-
- for (int i = 0; i < 10; ++i)
- {
- reader >> frame;
- ASSERT_FALSE(frame.empty());
-
- if (!writer.isOpened())
- writer.open(outputFile, CV_FOURCC('X', 'V', 'I', 'D'), FPS, frame.size());
-
- startTimer(); next();
- writer.write(frame);
- stopTimer();
- }
- }
-
- SANITY_CHECK(frame);
-}
-
-#endif
+++ /dev/null
-/*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*/
-
-/*
- * 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
- */
-
-#if !defined CUDA_DISABLER
-
-#include "opencv2/core/cuda/common.hpp"
-
-namespace cv { namespace gpu { namespace cudev {
- 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, PtrStepSz<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() );
- }
- }
-}}}
-
-#endif /* CUDA_DISABLER */
+++ /dev/null
-/*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*/
-
-#if !defined CUDA_DISABLER
-
-#include "opencv2/core/cuda/common.hpp"
-#include "opencv2/core/cuda/vec_traits.hpp"
-
-namespace cv { namespace gpu { namespace cudev
-{
- namespace video_encoding
- {
- __device__ __forceinline__ void rgbtoy(const uchar b, const uchar g, const uchar r, uchar& y)
- {
- y = static_cast<uchar>(((int)(30 * r) + (int)(59 * g) + (int)(11 * b)) / 100);
- }
-
- __device__ __forceinline__ void rgbtoyuv(const uchar b, const uchar g, const uchar r, uchar& y, uchar& u, uchar& v)
- {
- rgbtoy(b, g, r, y);
- u = static_cast<uchar>(((int)(-17 * r) - (int)(33 * g) + (int)(50 * b) + 12800) / 100);
- v = static_cast<uchar>(((int)(50 * r) - (int)(42 * g) - (int)(8 * b) + 12800) / 100);
- }
-
- __global__ void Gray_to_YV12(const PtrStepSzb src, PtrStepb dst)
- {
- const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
- const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
-
- if (x + 1 >= src.cols || y + 1 >= src.rows)
- return;
-
- // get pointers to the data
- const size_t planeSize = src.rows * dst.step;
- PtrStepb y_plane(dst.data, dst.step);
- PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
- PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
-
- uchar pix;
- uchar y_val, u_val, v_val;
-
- pix = src(y, x);
- rgbtoy(pix, pix, pix, y_val);
- y_plane(y, x) = y_val;
-
- pix = src(y, x + 1);
- rgbtoy(pix, pix, pix, y_val);
- y_plane(y, x + 1) = y_val;
-
- pix = src(y + 1, x);
- rgbtoy(pix, pix, pix, y_val);
- y_plane(y + 1, x) = y_val;
-
- pix = src(y + 1, x + 1);
- rgbtoyuv(pix, pix, pix, y_val, u_val, v_val);
- y_plane(y + 1, x + 1) = y_val;
- u_plane(y / 2, x / 2) = u_val;
- v_plane(y / 2, x / 2) = v_val;
- }
-
- template <typename T>
- __global__ void BGR_to_YV12(const PtrStepSz<T> src, PtrStepb dst)
- {
- const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
- const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
-
- if (x + 1 >= src.cols || y + 1 >= src.rows)
- return;
-
- // get pointers to the data
- const size_t planeSize = src.rows * dst.step;
- PtrStepb y_plane(dst.data, dst.step);
- PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
- PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
-
- T pix;
- uchar y_val, u_val, v_val;
-
- pix = src(y, x);
- rgbtoy(pix.z, pix.y, pix.x, y_val);
- y_plane(y, x) = y_val;
-
- pix = src(y, x + 1);
- rgbtoy(pix.z, pix.y, pix.x, y_val);
- y_plane(y, x + 1) = y_val;
-
- pix = src(y + 1, x);
- rgbtoy(pix.z, pix.y, pix.x, y_val);
- y_plane(y + 1, x) = y_val;
-
- pix = src(y + 1, x + 1);
- rgbtoyuv(pix.z, pix.y, pix.x, y_val, u_val, v_val);
- y_plane(y + 1, x + 1) = y_val;
- u_plane(y / 2, x / 2) = u_val;
- v_plane(y / 2, x / 2) = v_val;
- }
-
- void Gray_to_YV12_caller(const PtrStepSzb src, PtrStepb dst)
- {
- dim3 block(32, 8);
- dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
-
- Gray_to_YV12<<<grid, block>>>(src, dst);
- cudaSafeCall( cudaGetLastError() );
-
- cudaSafeCall( cudaDeviceSynchronize() );
- }
- template <int cn>
- void BGR_to_YV12_caller(const PtrStepSzb src, PtrStepb dst)
- {
- typedef typename TypeVec<uchar, cn>::vec_type src_t;
-
- dim3 block(32, 8);
- dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
-
- BGR_to_YV12<<<grid, block>>>(static_cast< PtrStepSz<src_t> >(src), dst);
- cudaSafeCall( cudaGetLastError() );
-
- cudaSafeCall( cudaDeviceSynchronize() );
- }
-
- void YV12_gpu(const PtrStepSzb src, int cn, PtrStepSzb dst)
- {
- typedef void (*func_t)(const PtrStepSzb src, PtrStepb dst);
-
- static const func_t funcs[] =
- {
- 0, Gray_to_YV12_caller, 0, BGR_to_YV12_caller<3>, BGR_to_YV12_caller<4>
- };
-
- funcs[cn](src, dst);
- }
- }
-}}}
-
-#endif /* CUDA_DISABLER */
#include <cublas.h>
#endif
- #ifdef HAVE_NVCUVID
- #include <nvcuvid.h>
-
- #ifdef WIN32
- #include <windows.h>
- #undef small
- #undef min
- #undef max
- #undef abs
-
- #include <NVEncoderAPI.h>
- #endif
- #endif
-
#include "internal_shared.hpp"
#include "opencv2/core/stream_accessor.hpp"
+++ /dev/null
-/*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"
-
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
-
-#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
+++ /dev/null
-/*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"
-
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
-
-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__
--- /dev/null
+if(ANDROID OR IOS)
+ ocv_module_disable(gpucodec)
+endif()
+
+set(the_description "GPU-accelerated Video Encoding/Decoding")
+
+ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations)
+
+ocv_add_module(gpucodec opencv_highgui)
+
+ocv_module_include_directories()
+ocv_glob_module_sources()
+
+set(extra_libs ${HIGHGUI_LIBRARIES})
+
+if(HAVE_NVCUVID)
+ list(APPEND extra_libs ${CUDA_nvcuvid_LIBRARY})
+
+ if(WIN32)
+ list(APPEND extra_libs ${CUDA_nvcuvenc_LIBRARY})
+ endif()
+endif()
+
+ocv_create_module(${extra_libs})
+
+ocv_add_precompiled_headers(${the_module})
+
+ocv_add_accuracy_tests()
+ocv_add_perf_tests()
--- /dev/null
+*************************************************
+gpucodec. GPU-accelerated Video Encoding/Decoding
+*************************************************
+
+.. toctree::
+ :maxdepth: 1
+
+ videodec
+ videoenc
--- /dev/null
+Video Decoding
+==============
+
+.. highlight:: cpp
+
+
+
+gpu::VideoReader_GPU
+--------------------
+Video reader class.
+
+.. ocv:class:: gpu::VideoReader_GPU
+
+
+
+gpu::VideoReader_GPU::Codec
+---------------------------
+
+Video codecs supported by :ocv:class:`gpu::VideoReader_GPU` .
+
+.. ocv:enum:: gpu::VideoReader_GPU::Codec
+
+ .. ocv:emember:: MPEG1 = 0
+ .. ocv:emember:: MPEG2
+ .. ocv:emember:: MPEG4
+ .. ocv:emember:: VC1
+ .. ocv:emember:: H264
+ .. ocv:emember:: JPEG
+ .. ocv:emember:: H264_SVC
+ .. ocv:emember:: H264_MVC
+
+ .. ocv:emember:: Uncompressed_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V'))
+
+ Y,U,V (4:2:0)
+
+ .. ocv:emember:: Uncompressed_YV12 = (('Y'<<24)|('V'<<16)|('1'<<8)|('2'))
+
+ Y,V,U (4:2:0)
+
+ .. ocv:emember:: Uncompressed_NV12 = (('N'<<24)|('V'<<16)|('1'<<8)|('2'))
+
+ Y,UV (4:2:0)
+
+ .. ocv:emember:: Uncompressed_YUYV = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V'))
+
+ YUYV/YUY2 (4:2:2)
+
+ .. ocv:emember:: Uncompressed_UYVY = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y'))
+
+ UYVY (4:2:2)
+
+
+gpu::VideoReader_GPU::ChromaFormat
+----------------------------------
+
+Chroma formats supported by :ocv:class:`gpu::VideoReader_GPU` .
+
+.. ocv:enum:: gpu::VideoReader_GPU::ChromaFormat
+
+ .. ocv:emember:: Monochrome = 0
+ .. ocv:emember:: YUV420
+ .. ocv:emember:: YUV422
+ .. ocv:emember:: YUV444
+
+
+gpu::VideoReader_GPU::FormatInfo
+--------------------------------
+.. ocv:struct:: gpu::VideoReader_GPU::FormatInfo
+
+Struct providing information about video file format. ::
+
+ struct FormatInfo
+ {
+ Codec codec;
+ ChromaFormat chromaFormat;
+ int width;
+ int height;
+ };
+
+
+gpu::VideoReader_GPU::VideoReader_GPU
+-------------------------------------
+Constructors.
+
+.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU()
+.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU(const String& filename)
+.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU(const cv::Ptr<VideoSource>& source)
+
+ :param filename: Name of the input video file.
+
+ :param source: Video file parser implemented by user.
+
+The constructors initialize video reader. FFMPEG is used to read videos. User can implement own demultiplexing with :ocv:class:`gpu::VideoReader_GPU::VideoSource` .
+
+
+
+gpu::VideoReader_GPU::open
+--------------------------
+Initializes or reinitializes video reader.
+
+.. ocv:function:: void gpu::VideoReader_GPU::open(const String& filename)
+.. ocv:function:: void gpu::VideoReader_GPU::open(const cv::Ptr<VideoSource>& source)
+
+The method opens video reader. Parameters are the same as in the constructor :ocv:func:`gpu::VideoReader_GPU::VideoReader_GPU` . The method throws :ocv:class:`Exception` if error occurs.
+
+
+
+gpu::VideoReader_GPU::isOpened
+------------------------------
+Returns true if video reader has been successfully initialized.
+
+.. ocv:function:: bool gpu::VideoReader_GPU::isOpened() const
+
+
+
+gpu::VideoReader_GPU::close
+---------------------------
+Releases the video reader.
+
+.. ocv:function:: void gpu::VideoReader_GPU::close()
+
+
+
+gpu::VideoReader_GPU::read
+--------------------------
+Grabs, decodes and returns the next video frame.
+
+.. ocv:function:: bool gpu::VideoReader_GPU::read(GpuMat& image)
+
+If no frames has been grabbed (there are no more frames in video file), the methods return ``false`` . The method throws :ocv:class:`Exception` if error occurs.
+
+
+
+gpu::VideoReader_GPU::format
+----------------------------
+Returns information about video file format.
+
+.. ocv:function:: FormatInfo gpu::VideoReader_GPU::format() const
+
+The method throws :ocv:class:`Exception` if video reader wasn't initialized.
+
+
+
+gpu::VideoReader_GPU::dumpFormat
+--------------------------------
+Dump information about video file format to specified stream.
+
+.. ocv:function:: void gpu::VideoReader_GPU::dumpFormat(std::ostream& st)
+
+ :param st: Output stream.
+
+The method throws :ocv:class:`Exception` if video reader wasn't initialized.
+
+
+
+gpu::VideoReader_GPU::VideoSource
+-----------------------------------
+.. ocv:class:: gpu::VideoReader_GPU::VideoSource
+
+Interface for video demultiplexing. ::
+
+ class VideoSource
+ {
+ public:
+ VideoSource();
+ 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;
+
+ protected:
+ bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream = false);
+ };
+
+User can implement own demultiplexing by implementing this interface.
+
+
+
+gpu::VideoReader_GPU::VideoSource::format
+-----------------------------------------
+Returns information about video file format.
+
+.. ocv:function:: virtual FormatInfo gpu::VideoReader_GPU::VideoSource::format() const = 0
+
+
+
+gpu::VideoReader_GPU::VideoSource::start
+----------------------------------------
+Starts processing.
+
+.. ocv:function:: virtual void gpu::VideoReader_GPU::VideoSource::start() = 0
+
+Implementation must create own thread with video processing and call periodic :ocv:func:`gpu::VideoReader_GPU::VideoSource::parseVideoData` .
+
+
+
+gpu::VideoReader_GPU::VideoSource::stop
+---------------------------------------
+Stops processing.
+
+.. ocv:function:: virtual void gpu::VideoReader_GPU::VideoSource::stop() = 0
+
+
+
+gpu::VideoReader_GPU::VideoSource::isStarted
+--------------------------------------------
+Returns ``true`` if processing was successfully started.
+
+.. ocv:function:: virtual bool gpu::VideoReader_GPU::VideoSource::isStarted() const = 0
+
+
+
+gpu::VideoReader_GPU::VideoSource::hasError
+-------------------------------------------
+Returns ``true`` if error occured during processing.
+
+.. ocv:function:: virtual bool gpu::VideoReader_GPU::VideoSource::hasError() const = 0
+
+
+
+gpu::VideoReader_GPU::VideoSource::parseVideoData
+-------------------------------------------------
+Parse next video frame. Implementation must call this method after new frame was grabbed.
+
+.. ocv:function:: bool gpu::VideoReader_GPU::VideoSource::parseVideoData(const uchar* data, size_t size, bool endOfStream = false)
+
+ :param data: Pointer to frame data. Can be ``NULL`` if ``endOfStream`` if ``true`` .
+
+ :param size: Size in bytes of current frame.
+
+ :param endOfStream: Indicates that it is end of stream.
--- /dev/null
+Video Encoding
+==============
+
+.. highlight:: cpp
+
+
+
+gpu::VideoWriter_GPU
+---------------------
+Video writer class.
+
+.. ocv:class:: gpu::VideoWriter_GPU
+
+The class uses H264 video codec.
+
+.. note:: Currently only Windows platform is supported.
+
+
+
+gpu::VideoWriter_GPU::VideoWriter_GPU
+-------------------------------------
+Constructors.
+
+.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU()
+.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
+.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
+.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
+.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
+
+ :param fileName: Name of the output video file. Only AVI file format is supported.
+
+ :param frameSize: Size of the input video frames.
+
+ :param fps: Framerate of the created video stream.
+
+ :param params: Encoder parameters. See :ocv:struct:`gpu::VideoWriter_GPU::EncoderParams` .
+
+ :param format: Surface format of input frames ( ``SF_UYVY`` , ``SF_YUY2`` , ``SF_YV12`` , ``SF_NV12`` , ``SF_IYUV`` , ``SF_BGR`` or ``SF_GRAY``). BGR or gray frames will be converted to YV12 format before encoding, frames with other formats will be used as is.
+
+ :param encoderCallback: Callbacks for video encoder. See :ocv:class:`gpu::VideoWriter_GPU::EncoderCallBack` . Use it if you want to work with raw video stream.
+
+The constructors initialize video writer. FFMPEG is used to write videos. User can implement own multiplexing with :ocv:class:`gpu::VideoWriter_GPU::EncoderCallBack` .
+
+
+
+gpu::VideoWriter_GPU::open
+--------------------------
+Initializes or reinitializes video writer.
+
+.. ocv:function:: void gpu::VideoWriter_GPU::open(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
+.. ocv:function:: void gpu::VideoWriter_GPU::open(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
+.. ocv:function:: void gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
+.. ocv:function:: void gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
+
+The method opens video writer. Parameters are the same as in the constructor :ocv:func:`gpu::VideoWriter_GPU::VideoWriter_GPU` . The method throws :ocv:class:`Exception` if error occurs.
+
+
+
+gpu::VideoWriter_GPU::isOpened
+------------------------------
+Returns true if video writer has been successfully initialized.
+
+.. ocv:function:: bool gpu::VideoWriter_GPU::isOpened() const
+
+
+
+gpu::VideoWriter_GPU::close
+---------------------------
+Releases the video writer.
+
+.. ocv:function:: void gpu::VideoWriter_GPU::close()
+
+
+
+gpu::VideoWriter_GPU::write
+---------------------------
+Writes the next video frame.
+
+.. ocv:function:: void gpu::VideoWriter_GPU::write(const cv::gpu::GpuMat& image, bool lastFrame = false)
+
+ :param image: The written frame.
+
+ :param lastFrame: Indicates that it is end of stream. The parameter can be ignored.
+
+The method write the specified image to video file. The image must have the same size and the same surface format as has been specified when opening the video writer.
+
+
+
+gpu::VideoWriter_GPU::EncoderParams
+-----------------------------------
+.. ocv:struct:: gpu::VideoWriter_GPU::EncoderParams
+
+Different parameters for CUDA video encoder. ::
+
+ struct EncoderParams
+ {
+ int P_Interval; // NVVE_P_INTERVAL,
+ int IDR_Period; // NVVE_IDR_PERIOD,
+ int DynamicGOP; // NVVE_DYNAMIC_GOP,
+ int RCType; // NVVE_RC_TYPE,
+ int AvgBitrate; // NVVE_AVG_BITRATE,
+ int PeakBitrate; // NVVE_PEAK_BITRATE,
+ int QP_Level_Intra; // NVVE_QP_LEVEL_INTRA,
+ int QP_Level_InterP; // NVVE_QP_LEVEL_INTER_P,
+ int QP_Level_InterB; // NVVE_QP_LEVEL_INTER_B,
+ int DeblockMode; // NVVE_DEBLOCK_MODE,
+ int ProfileLevel; // NVVE_PROFILE_LEVEL,
+ int ForceIntra; // NVVE_FORCE_INTRA,
+ int ForceIDR; // NVVE_FORCE_IDR,
+ int ClearStat; // NVVE_CLEAR_STAT,
+ int DIMode; // NVVE_SET_DEINTERLACE,
+ int Presets; // NVVE_PRESETS,
+ int DisableCabac; // NVVE_DISABLE_CABAC,
+ int NaluFramingType; // NVVE_CONFIGURE_NALU_FRAMING_TYPE
+ int DisableSPSPPS; // NVVE_DISABLE_SPS_PPS
+
+ EncoderParams();
+ explicit EncoderParams(const String& configFile);
+
+ void load(const String& configFile);
+ void save(const String& configFile) const;
+ };
+
+
+
+gpu::VideoWriter_GPU::EncoderParams::EncoderParams
+--------------------------------------------------
+Constructors.
+
+.. ocv:function:: gpu::VideoWriter_GPU::EncoderParams::EncoderParams()
+.. ocv:function:: gpu::VideoWriter_GPU::EncoderParams::EncoderParams(const String& configFile)
+
+ :param configFile: Config file name.
+
+Creates default parameters or reads parameters from config file.
+
+
+
+gpu::VideoWriter_GPU::EncoderParams::load
+-----------------------------------------
+Reads parameters from config file.
+
+.. ocv:function:: void gpu::VideoWriter_GPU::EncoderParams::load(const String& configFile)
+
+ :param configFile: Config file name.
+
+
+
+gpu::VideoWriter_GPU::EncoderParams::save
+-----------------------------------------
+Saves parameters to config file.
+
+.. ocv:function:: void gpu::VideoWriter_GPU::EncoderParams::save(const String& configFile) const
+
+ :param configFile: Config file name.
+
+
+
+gpu::VideoWriter_GPU::EncoderCallBack
+-------------------------------------
+.. ocv:class:: gpu::VideoWriter_GPU::EncoderCallBack
+
+Callbacks for CUDA video encoder. ::
+
+ class EncoderCallBack
+ {
+ public:
+ enum PicType
+ {
+ IFRAME = 1,
+ PFRAME = 2,
+ BFRAME = 3
+ };
+
+ virtual ~EncoderCallBack() {}
+
+ virtual unsigned char* acquireBitStream(int* bufferSize) = 0;
+ virtual void releaseBitStream(unsigned char* data, int size) = 0;
+ virtual void onBeginFrame(int frameNumber, PicType picType) = 0;
+ virtual void onEndFrame(int frameNumber, PicType picType) = 0;
+ };
+
+
+
+gpu::VideoWriter_GPU::EncoderCallBack::acquireBitStream
+-------------------------------------------------------
+Callback function to signal the start of bitstream that is to be encoded.
+
+.. ocv:function:: virtual uchar* gpu::VideoWriter_GPU::EncoderCallBack::acquireBitStream(int* bufferSize) = 0
+
+Callback must allocate buffer for CUDA encoder and return pointer to it and it's size.
+
+
+
+gpu::VideoWriter_GPU::EncoderCallBack::releaseBitStream
+-------------------------------------------------------
+Callback function to signal that the encoded bitstream is ready to be written to file.
+
+.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::releaseBitStream(unsigned char* data, int size) = 0
+
+
+
+gpu::VideoWriter_GPU::EncoderCallBack::onBeginFrame
+---------------------------------------------------
+Callback function to signal that the encoding operation on the frame has started.
+
+.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::onBeginFrame(int frameNumber, PicType picType) = 0
+
+ :param picType: Specify frame type (I-Frame, P-Frame or B-Frame).
+
+
+
+gpu::VideoWriter_GPU::EncoderCallBack::onEndFrame
+-------------------------------------------------
+Callback function signals that the encoding operation on the frame has finished.
+
+.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::onEndFrame(int frameNumber, PicType picType) = 0
+
+ :param picType: Specify frame type (I-Frame, P-Frame or B-Frame).
--- /dev/null
+/*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 __OPENCV_GPUCODEC_HPP__
+#define __OPENCV_GPUCODEC_HPP__
+
+#ifndef __cplusplus
+# error gpucodec.hpp header must be compiled as C++
+#endif
+
+#include <iosfwd>
+
+#include "opencv2/core/gpumat.hpp"
+
+namespace cv { namespace gpu {
+
+////////////////////////////////// Video Encoding //////////////////////////////////
+
+// Works only under Windows
+// Supports olny H264 video codec and AVI files
+class CV_EXPORTS VideoWriter_GPU
+{
+public:
+ struct EncoderParams;
+
+ // Callbacks for video encoder, use it if you want to work with raw video stream
+ class EncoderCallBack;
+
+ enum SurfaceFormat
+ {
+ SF_UYVY = 0,
+ SF_YUY2,
+ SF_YV12,
+ SF_NV12,
+ SF_IYUV,
+ SF_BGR,
+ SF_GRAY = SF_BGR
+ };
+
+ VideoWriter_GPU();
+ VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
+ VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
+ VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
+ VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
+ ~VideoWriter_GPU();
+
+ // all methods throws cv::Exception if error occurs
+ void open(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
+ void open(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
+ void open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
+ void open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
+
+ bool isOpened() const;
+ void close();
+
+ void write(const cv::gpu::GpuMat& image, bool lastFrame = false);
+
+ struct CV_EXPORTS EncoderParams
+ {
+ int P_Interval; // NVVE_P_INTERVAL,
+ int IDR_Period; // NVVE_IDR_PERIOD,
+ int DynamicGOP; // NVVE_DYNAMIC_GOP,
+ int RCType; // NVVE_RC_TYPE,
+ int AvgBitrate; // NVVE_AVG_BITRATE,
+ int PeakBitrate; // NVVE_PEAK_BITRATE,
+ int QP_Level_Intra; // NVVE_QP_LEVEL_INTRA,
+ int QP_Level_InterP; // NVVE_QP_LEVEL_INTER_P,
+ int QP_Level_InterB; // NVVE_QP_LEVEL_INTER_B,
+ int DeblockMode; // NVVE_DEBLOCK_MODE,
+ int ProfileLevel; // NVVE_PROFILE_LEVEL,
+ int ForceIntra; // NVVE_FORCE_INTRA,
+ int ForceIDR; // NVVE_FORCE_IDR,
+ int ClearStat; // NVVE_CLEAR_STAT,
+ int DIMode; // NVVE_SET_DEINTERLACE,
+ int Presets; // NVVE_PRESETS,
+ int DisableCabac; // NVVE_DISABLE_CABAC,
+ int NaluFramingType; // NVVE_CONFIGURE_NALU_FRAMING_TYPE
+ int DisableSPSPPS; // NVVE_DISABLE_SPS_PPS
+
+ EncoderParams();
+ explicit EncoderParams(const String& configFile);
+
+ void load(const String& configFile);
+ void save(const String& configFile) const;
+ };
+
+ EncoderParams getParams() const;
+
+ class CV_EXPORTS EncoderCallBack
+ {
+ public:
+ enum PicType
+ {
+ IFRAME = 1,
+ PFRAME = 2,
+ BFRAME = 3
+ };
+
+ virtual ~EncoderCallBack() {}
+
+ // callback function to signal the start of bitstream that is to be encoded
+ // must return pointer to buffer
+ virtual uchar* acquireBitStream(int* bufferSize) = 0;
+
+ // callback function to signal that the encoded bitstream is ready to be written to file
+ virtual void releaseBitStream(unsigned char* data, int size) = 0;
+
+ // callback function to signal that the encoding operation on the frame has started
+ virtual void onBeginFrame(int frameNumber, PicType picType) = 0;
+
+ // callback function signals that the encoding operation on the frame has finished
+ virtual void onEndFrame(int frameNumber, PicType picType) = 0;
+ };
+
+ class Impl;
+
+private:
+ cv::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 String& filename);
+ explicit VideoReader_GPU(const cv::Ptr<VideoSource>& source);
+
+ ~VideoReader_GPU();
+
+ void open(const 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 CV_EXPORTS 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 uchar* data, size_t size, bool endOfStream = false);
+
+ private:
+ VideoSource(const VideoSource&);
+ VideoSource& operator =(const VideoSource&);
+
+ detail::FrameQueue* frameQueue_;
+ detail::VideoParser* videoParser_;
+ };
+
+ class Impl;
+
+private:
+ cv::Ptr<Impl> impl_;
+};
+
+}} // namespace cv { namespace gpu {
+
+namespace cv {
+
+template <> CV_EXPORTS void Ptr<cv::gpu::VideoWriter_GPU::Impl>::delete_obj();
+template <> CV_EXPORTS void Ptr<cv::gpu::VideoReader_GPU::Impl>::delete_obj();
+
+}
+
+#endif /* __OPENCV_GPUCODEC_HPP__ */
--- /dev/null
+/*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 "perf_precomp.hpp"
+
+using namespace perf;
+
+CV_PERF_TEST_MAIN(gpucodec, printCudaInfo())
--- /dev/null
+/*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 "perf_precomp.hpp"
--- /dev/null
+/*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*/
+
+#ifdef __GNUC__
+# pragma GCC diagnostic ignored "-Wmissing-declarations"
+# if defined __clang__ || defined __APPLE__
+# pragma GCC diagnostic ignored "-Wmissing-prototypes"
+# pragma GCC diagnostic ignored "-Wextra"
+# endif
+#endif
+
+#ifndef __OPENCV_PERF_PRECOMP_HPP__
+#define __OPENCV_PERF_PRECOMP_HPP__
+
+#include "opencv2/ts.hpp"
+#include "opencv2/ts/gpu_perf.hpp"
+
+#include "opencv2/gpucodec.hpp"
+#include "opencv2/highgui.hpp"
+
+#ifdef GTEST_CREATE_SHARED_LIBRARY
+#error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined
+#endif
+
+#endif
--- /dev/null
+/*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 "perf_precomp.hpp"
+#include "opencv2/highgui/highgui_c.h"
+
+using namespace std;
+using namespace testing;
+using namespace perf;
+
+#if defined(HAVE_XINE) || \
+ defined(HAVE_GSTREAMER) || \
+ defined(HAVE_QUICKTIME) || \
+ defined(HAVE_AVFOUNDATION) || \
+ defined(HAVE_FFMPEG) || \
+ defined(WIN32) /* assume that we have ffmpeg */
+
+# define BUILD_WITH_VIDEO_INPUT_SUPPORT 1
+#else
+# define BUILD_WITH_VIDEO_INPUT_SUPPORT 0
+#endif
+
+DEF_PARAM_TEST_1(FileName, string);
+
+//////////////////////////////////////////////////////
+// VideoReader
+
+#if defined(HAVE_NVCUVID) && BUILD_WITH_VIDEO_INPUT_SUPPORT
+
+PERF_TEST_P(FileName, VideoReader, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
+{
+ declare.time(20);
+
+ const string inputFile = perf::TestBase::getDataPath(GetParam());
+
+ if (PERF_RUN_GPU())
+ {
+ cv::gpu::VideoReader_GPU d_reader(inputFile);
+ ASSERT_TRUE( d_reader.isOpened() );
+
+ cv::gpu::GpuMat frame;
+
+ TEST_CYCLE_N(10) d_reader.read(frame);
+
+ GPU_SANITY_CHECK(frame);
+ }
+ else
+ {
+ cv::VideoCapture reader(inputFile);
+ ASSERT_TRUE( reader.isOpened() );
+
+ cv::Mat frame;
+
+ TEST_CYCLE_N(10) reader >> frame;
+
+ CPU_SANITY_CHECK(frame);
+ }
+}
+
+#endif
+
+//////////////////////////////////////////////////////
+// VideoWriter
+
+#if defined(HAVE_NVCUVID) && defined(WIN32)
+
+PERF_TEST_P(FileName, VideoWriter, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
+{
+ declare.time(30);
+
+ const string inputFile = perf::TestBase::getDataPath(GetParam());
+ const string outputFile = cv::tempfile(".avi");
+
+ const double FPS = 25.0;
+
+ cv::VideoCapture reader(inputFile);
+ ASSERT_TRUE( reader.isOpened() );
+
+ cv::Mat frame;
+
+ if (PERF_RUN_GPU())
+ {
+ cv::gpu::VideoWriter_GPU d_writer;
+
+ cv::gpu::GpuMat d_frame;
+
+ for (int i = 0; i < 10; ++i)
+ {
+ reader >> frame;
+ ASSERT_FALSE(frame.empty());
+
+ d_frame.upload(frame);
+
+ if (!d_writer.isOpened())
+ d_writer.open(outputFile, frame.size(), FPS);
+
+ startTimer(); next();
+ d_writer.write(d_frame);
+ stopTimer();
+ }
+ }
+ else
+ {
+ cv::VideoWriter writer;
+
+ for (int i = 0; i < 10; ++i)
+ {
+ reader >> frame;
+ ASSERT_FALSE(frame.empty());
+
+ if (!writer.isOpened())
+ writer.open(outputFile, CV_FOURCC('X', 'V', 'I', 'D'), FPS, frame.size());
+
+ startTimer(); next();
+ writer.write(frame);
+ stopTimer();
+ }
+ }
+
+ SANITY_CHECK(frame);
+}
+
+#endif
--- /dev/null
+/*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*/
+
+/*
+ * 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/core/cuda/common.hpp"
+
+namespace cv { namespace gpu { namespace cudev
+{
+ __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 RGBA_pack_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 NV12_to_RGB(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 ] = RGBA_pack_10bit(red[0], green[0], blue[0], ((uint)0xff << 24));
+ dstImage[y * dstImagePitch + x + 1 ] = RGBA_pack_10bit(red[1], green[1], blue[1], ((uint)0xff << 24));
+ }
+
+ void NV12_to_RGB(const PtrStepb decodedFrame, PtrStepSz<uint> interopFrame, cudaStream_t stream)
+ {
+ dim3 block(32, 8);
+ dim3 grid(divUp(interopFrame.cols, 2 * block.x), divUp(interopFrame.rows, block.y));
+
+ NV12_to_RGB<<<grid, block, 0, stream>>>(decodedFrame.data, decodedFrame.step, interopFrame.data, interopFrame.step,
+ interopFrame.cols, interopFrame.rows);
+
+ cudaSafeCall( cudaGetLastError() );
+
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+}}}
--- /dev/null
+/*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 "opencv2/core/cuda/common.hpp"
+#include "opencv2/core/cuda/vec_traits.hpp"
+
+namespace cv { namespace gpu { namespace cudev
+{
+ __device__ __forceinline__ void rgb_to_y(const uchar b, const uchar g, const uchar r, uchar& y)
+ {
+ y = static_cast<uchar>(((int)(30 * r) + (int)(59 * g) + (int)(11 * b)) / 100);
+ }
+
+ __device__ __forceinline__ void rgb_to_yuv(const uchar b, const uchar g, const uchar r, uchar& y, uchar& u, uchar& v)
+ {
+ rgb_to_y(b, g, r, y);
+ u = static_cast<uchar>(((int)(-17 * r) - (int)(33 * g) + (int)(50 * b) + 12800) / 100);
+ v = static_cast<uchar>(((int)(50 * r) - (int)(42 * g) - (int)(8 * b) + 12800) / 100);
+ }
+
+ __global__ void Gray_to_YV12(const PtrStepSzb src, PtrStepb dst)
+ {
+ const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
+ const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
+
+ if (x + 1 >= src.cols || y + 1 >= src.rows)
+ return;
+
+ // get pointers to the data
+ const size_t planeSize = src.rows * dst.step;
+ PtrStepb y_plane(dst.data, dst.step);
+ PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
+ PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
+
+ uchar pix;
+ uchar y_val, u_val, v_val;
+
+ pix = src(y, x);
+ rgb_to_y(pix, pix, pix, y_val);
+ y_plane(y, x) = y_val;
+
+ pix = src(y, x + 1);
+ rgb_to_y(pix, pix, pix, y_val);
+ y_plane(y, x + 1) = y_val;
+
+ pix = src(y + 1, x);
+ rgb_to_y(pix, pix, pix, y_val);
+ y_plane(y + 1, x) = y_val;
+
+ pix = src(y + 1, x + 1);
+ rgb_to_yuv(pix, pix, pix, y_val, u_val, v_val);
+ y_plane(y + 1, x + 1) = y_val;
+ u_plane(y / 2, x / 2) = u_val;
+ v_plane(y / 2, x / 2) = v_val;
+ }
+
+ template <typename T>
+ __global__ void RGB_to_YV12(const PtrStepSz<T> src, PtrStepb dst)
+ {
+ const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
+ const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
+
+ if (x + 1 >= src.cols || y + 1 >= src.rows)
+ return;
+
+ // get pointers to the data
+ const size_t planeSize = src.rows * dst.step;
+ PtrStepb y_plane(dst.data, dst.step);
+ PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
+ PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
+
+ T pix;
+ uchar y_val, u_val, v_val;
+
+ pix = src(y, x);
+ rgb_to_y(pix.z, pix.y, pix.x, y_val);
+ y_plane(y, x) = y_val;
+
+ pix = src(y, x + 1);
+ rgb_to_y(pix.z, pix.y, pix.x, y_val);
+ y_plane(y, x + 1) = y_val;
+
+ pix = src(y + 1, x);
+ rgb_to_y(pix.z, pix.y, pix.x, y_val);
+ y_plane(y + 1, x) = y_val;
+
+ pix = src(y + 1, x + 1);
+ rgb_to_yuv(pix.z, pix.y, pix.x, y_val, u_val, v_val);
+ y_plane(y + 1, x + 1) = y_val;
+ u_plane(y / 2, x / 2) = u_val;
+ v_plane(y / 2, x / 2) = v_val;
+ }
+
+ void Gray_to_YV12_caller(const PtrStepSzb src, PtrStepb dst, cudaStream_t stream)
+ {
+ dim3 block(32, 8);
+ dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
+
+ Gray_to_YV12<<<grid, block, 0, stream>>>(src, dst);
+ cudaSafeCall( cudaGetLastError() );
+
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+ template <int cn>
+ void RGB_to_YV12_caller(const PtrStepSzb src, PtrStepb dst, cudaStream_t stream)
+ {
+ typedef typename TypeVec<uchar, cn>::vec_type src_t;
+
+ dim3 block(32, 8);
+ dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
+
+ RGB_to_YV12<<<grid, block, 0, stream>>>(static_cast< PtrStepSz<src_t> >(src), dst);
+ cudaSafeCall( cudaGetLastError() );
+
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+
+ void RGB_to_YV12(const PtrStepSzb src, int cn, PtrStepSzb dst, cudaStream_t stream)
+ {
+ typedef void (*func_t)(const PtrStepSzb src, PtrStepb dst, cudaStream_t stream);
+
+ static const func_t funcs[] =
+ {
+ 0, Gray_to_YV12_caller, 0, RGB_to_YV12_caller<3>, RGB_to_YV12_caller<4>
+ };
+
+ funcs[cn](src, dst, stream);
+ }
+}}}
//
//M*/
-#include "cuvid_video_source.h"
+#include "precomp.hpp"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#ifdef HAVE_NVCUVID
cv::gpu::detail::CuvidVideoSource::CuvidVideoSource(const String& fname)
{
format_.height = vidfmt.coded_height;
}
+cv::gpu::detail::CuvidVideoSource::~CuvidVideoSource()
+{
+ cuvidDestroyVideoSource(videoSource_);
+}
+
cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::detail::CuvidVideoSource::format() const
{
return format_;
return thiz->parseVideoData(packet->payload, packet->payload_size, (packet->flags & CUVID_PKT_ENDOFSTREAM) != 0);
}
-#endif // defined(HAVE_CUDA) && !defined(__APPLE__)
+#endif // HAVE_NVCUVID
#ifndef __CUVUD_VIDEO_SOURCE_H__
#define __CUVUD_VIDEO_SOURCE_H__
-#include "precomp.hpp"
+#include "opencv2/core/gpu_private.hpp"
+#include "opencv2/gpucodec.hpp"
+#include "thread.h"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#include <nvcuvid.h>
-namespace cv { namespace gpu
+namespace cv { namespace gpu { namespace detail
{
- namespace detail
- {
- class CuvidVideoSource : public VideoReader_GPU::VideoSource
- {
- public:
- explicit CuvidVideoSource(const String& fname);
- ~CuvidVideoSource() { cuvidDestroyVideoSource(videoSource_); }
- VideoReader_GPU::FormatInfo format() const;
- void start();
- void stop();
- bool isStarted() const;
- bool hasError() const;
+class CuvidVideoSource : public VideoReader_GPU::VideoSource
+{
+public:
+ explicit CuvidVideoSource(const String& fname);
+ ~CuvidVideoSource();
- private:
- CuvidVideoSource(const CuvidVideoSource&);
- CuvidVideoSource& operator =(const CuvidVideoSource&);
+ VideoReader_GPU::FormatInfo format() const;
+ void start();
+ void stop();
+ bool isStarted() const;
+ bool hasError() const;
- // 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);
+private:
+ // 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_;
- };
- }
-}}
+ CUvideosource videoSource_;
+ VideoReader_GPU::FormatInfo format_;
+};
-#endif // defined(HAVE_CUDA) && !defined(__APPLE__)
+}}}
#endif // __CUVUD_VIDEO_SOURCE_H__
//
//M*/
-#include "ffmpeg_video_source.h"
+#include "precomp.hpp"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#ifdef HAVE_NVCUVID
-#if defined(HAVE_FFMPEG) && defined(BUILD_SHARED_LIBS)
+#if defined(HAVE_FFMPEG) && defined(BUILD_SHARED_LIBS) && !defined(WIN32)
#include "../src/cap_ffmpeg_impl.hpp"
-#else
- #include "../src/cap_ffmpeg_api.hpp"
#endif
namespace
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_;
{
stop_ = false;
hasError_ = false;
- thread_.reset(new Thread(readLoop, this));
+ thread_ = new Thread(readLoop, this);
}
void cv::gpu::detail::FFmpegVideoSource::stop()
{
stop_ = true;
thread_->wait();
- thread_.reset();
+ thread_.release();
}
bool cv::gpu::detail::FFmpegVideoSource::isStarted() const
thiz->parseVideoData(0, 0, true);
}
+template <> void cv::Ptr<InputMediaStream_FFMPEG>::delete_obj()
+{
+ if (obj) release_InputMediaStream_FFMPEG_p(obj);
+}
+
#endif // HAVE_CUDA
#ifndef __FFMPEG_VIDEO_SOURCE_H__
#define __FFMPEG_VIDEO_SOURCE_H__
-#include "precomp.hpp"
-#include "thread_wrappers.h"
-
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#include "opencv2/gpucodec.hpp"
+#include "thread.h"
struct InputMediaStream_FFMPEG;
-namespace cv { namespace gpu
+namespace cv { namespace gpu { namespace detail {
+
+class FFmpegVideoSource : public VideoReader_GPU::VideoSource
{
- namespace detail
- {
- class FFmpegVideoSource : public VideoReader_GPU::VideoSource
- {
- public:
- FFmpegVideoSource(const String& fname);
- ~FFmpegVideoSource();
+public:
+ FFmpegVideoSource(const String& fname);
- VideoReader_GPU::FormatInfo format() const;
- void start();
- void stop();
- bool isStarted() const;
- bool hasError() const;
+ VideoReader_GPU::FormatInfo format() const;
+ void start();
+ void stop();
+ bool isStarted() const;
+ bool hasError() const;
- private:
- FFmpegVideoSource(const FFmpegVideoSource&);
- FFmpegVideoSource& operator =(const FFmpegVideoSource&);
+private:
+ VideoReader_GPU::FormatInfo format_;
- VideoReader_GPU::FormatInfo format_;
+ cv::Ptr<InputMediaStream_FFMPEG> stream_;
- InputMediaStream_FFMPEG* stream_;
+ cv::Ptr<Thread> thread_;
+ volatile bool stop_;
+ volatile bool hasError_;
- std::auto_ptr<Thread> thread_;
- volatile bool stop_;
- volatile bool hasError_;
+ static void readLoop(void* userData);
+};
- static void readLoop(void* userData);
- };
- }
-}}
+}}}
-#endif // HAVE_CUDA
+namespace cv {
+ template <> void Ptr<InputMediaStream_FFMPEG>::delete_obj();
+}
-#endif // __CUVUD_VIDEO_SOURCE_H__
+#endif // __FFMPEG_VIDEO_SOURCE_H__
//
//M*/
-#include "frame_queue.h"
+#include "precomp.hpp"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#ifdef HAVE_NVCUVID
cv::gpu::detail::FrameQueue::FrameQueue() :
endOfDecode_(0),
bool isFramePlaced = false;
{
- CriticalSection::AutoLock autoLock(criticalSection_);
+ AutoLock autoLock(mtx_);
if (framesInQueue_ < MaximumSize)
{
bool cv::gpu::detail::FrameQueue::dequeue(CUVIDPARSERDISPINFO& displayInfo)
{
- CriticalSection::AutoLock autoLock(criticalSection_);
+ AutoLock autoLock(mtx_);
if (framesInQueue_ > 0)
{
return false;
}
-#endif // HAVE_CUDA
+#endif // HAVE_NVCUVID
#ifndef __FRAME_QUEUE_H__
#define __FRAME_QUEUE_H__
-#include "precomp.hpp"
-#include "thread_wrappers.h"
+#include "opencv2/core/utility.hpp"
+#include "opencv2/core/gpu_private.hpp"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#include <nvcuvid.h>
-namespace cv { namespace gpu
+namespace cv { namespace gpu { namespace detail
{
- namespace detail
- {
- class FrameQueue
- {
- public:
- static const int MaximumSize = 20; // MAX_FRM_CNT;
- FrameQueue();
+class FrameQueue
+{
+public:
+ static const int MaximumSize = 20; // MAX_FRM_CNT;
- void endDecode() { endOfDecode_ = true; }
- bool isEndOfDecode() const { return endOfDecode_ != 0;}
+ FrameQueue();
- // 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 endDecode() { endOfDecode_ = true; }
+ bool isEndOfDecode() const { return endOfDecode_ != 0;}
- void enqueue(const CUVIDPARSERDISPINFO* picParams);
+ // 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);
- // 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 enqueue(const CUVIDPARSERDISPINFO* picParams);
- void releaseFrame(const CUVIDPARSERDISPINFO& picParams) { isFrameInUse_[picParams.picture_index] = false; }
+ // 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);
- private:
- FrameQueue(const FrameQueue&);
- FrameQueue& operator =(const FrameQueue&);
+ void releaseFrame(const CUVIDPARSERDISPINFO& picParams) { isFrameInUse_[picParams.picture_index] = false; }
- bool isInUse(int pictureIndex) const { return isFrameInUse_[pictureIndex] != 0; }
+private:
+ bool isInUse(int pictureIndex) const { return isFrameInUse_[pictureIndex] != 0; }
- CriticalSection criticalSection_;
+ Mutex mtx_;
- volatile int isFrameInUse_[MaximumSize];
- volatile int endOfDecode_;
+ volatile int isFrameInUse_[MaximumSize];
+ volatile int endOfDecode_;
- int framesInQueue_;
- int readPosition_;
- CUVIDPARSERDISPINFO displayQueue_[MaximumSize];
- };
- }
-}}
+ int framesInQueue_;
+ int readPosition_;
+ CUVIDPARSERDISPINFO displayQueue_[MaximumSize];
+};
-#endif // HAVE_CUDA
+}}}
#endif // __FRAME_QUEUE_H__
--- /dev/null
+/*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 "precomp.hpp"
--- /dev/null
+/*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 __OPENCV_PRECOMP_H__
+#define __OPENCV_PRECOMP_H__
+
+#include <cstdlib>
+#include <cstring>
+#include <deque>
+#include <utility>
+#include <stdexcept>
+#include <iostream>
+
+#include "opencv2/gpucodec.hpp"
+
+#include "opencv2/core/gpu_private.hpp"
+
+#ifdef HAVE_NVCUVID
+ #include <nvcuvid.h>
+
+ #ifdef WIN32
+ #define NOMINMAX
+ #include <windows.h>
+ #include <NVEncoderAPI.h>
+ #else
+ #include <pthread.h>
+ #include <unistd.h>
+ #endif
+
+ #include "thread.h"
+ #include "ffmpeg_video_source.h"
+ #include "cuvid_video_source.h"
+ #include "frame_queue.h"
+ #include "video_decoder.h"
+ #include "video_parser.h"
+
+ #include "../src/cap_ffmpeg_api.hpp"
+#endif
+
+#endif /* __OPENCV_PRECOMP_H__ */
--- /dev/null
+/*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 "precomp.hpp"
+
+#ifdef HAVE_NVCUVID
+
+using namespace cv::gpu::detail;
+
+#ifdef WIN32
+
+namespace
+{
+ struct UserData
+ {
+ Thread::Func func;
+ 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(Thread::Func func, 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
+ {
+ Thread::Func func;
+ 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(Thread::Func func, 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(Func func, void* userData) :
+ impl_(new Impl(func, userData))
+{
+}
+
+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
+}
+
+template <> void cv::Ptr<cv::gpu::detail::Thread::Impl>::delete_obj()
+{
+ if (obj) delete obj;
+}
+
+#endif // HAVE_NVCUVID
#ifndef __THREAD_WRAPPERS_H__
#define __THREAD_WRAPPERS_H__
-#include "precomp.hpp"
+#include "opencv2/core.hpp"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+namespace cv { namespace gpu { namespace detail {
-namespace cv { namespace gpu
+class Thread
{
- namespace detail
- {
- class CriticalSection
- {
- public:
- CriticalSection();
- ~CriticalSection();
+public:
+ typedef void (*Func)(void* userData);
- void enter();
- void leave();
+ explicit Thread(Func func, void* userData = 0);
- class AutoLock
- {
- public:
- explicit AutoLock(CriticalSection& criticalSection) :
- criticalSection_(criticalSection)
- {
- criticalSection_.enter();
- }
+ void wait();
- ~AutoLock()
- {
- criticalSection_.leave();
- }
+ static void sleep(int ms);
- private:
- CriticalSection& criticalSection_;
- };
+ class Impl;
- private:
- CriticalSection(const CriticalSection&);
- CriticalSection& operator=(const CriticalSection&);
+private:
+ cv::Ptr<Impl> impl_;
+};
- 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
+namespace cv {
+ template <> void Ptr<cv::gpu::detail::Thread::Impl>::delete_obj();
+}
#endif // __THREAD_WRAPPERS_H__
//
//M*/
-#include "video_decoder.h"
-#include "frame_queue.h"
+#include "precomp.hpp"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#ifdef HAVE_NVCUVID
void cv::gpu::detail::VideoDecoder::create(const VideoReader_GPU::FormatInfo& videoFormat)
{
}
}
-#endif // HAVE_CUDA
+#endif // HAVE_NVCUVID
--- /dev/null
+/*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 "opencv2/core/gpu_private.hpp"
+#include "opencv2/gpucodec.hpp"
+
+#include <nvcuvid.h>
+
+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:
+ CUvideoctxlock lock_;
+ CUVIDDECODECREATEINFO createInfo_;
+ CUvideodecoder decoder_;
+};
+
+}}}
+
+#endif // __VIDEO_DECODER_H__
//
//M*/
-#include "video_parser.h"
+#include "precomp.hpp"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#ifdef HAVE_NVCUVID
cv::gpu::detail::VideoParser::VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue) :
videoDecoder_(videoDecoder), frameQueue_(frameQueue), unparsedPackets_(0), hasError_(false)
return true;
}
-#endif // HAVE_CUDA
+#endif // HAVE_NVCUVID
#ifndef __VIDEO_PARSER_H__
#define __VIDEO_PARSER_H__
-#include "precomp.hpp"
-
+#include "opencv2/core/gpu_private.hpp"
+#include "opencv2/gpucodec.hpp"
#include "frame_queue.h"
#include "video_decoder.h"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#include <nvcuvid.h>
-namespace cv { namespace gpu
+namespace cv { namespace gpu { namespace detail
{
- namespace detail
- {
- class VideoParser
- {
- public:
- VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue);
- ~VideoParser()
- {
- cuvidDestroyVideoParser(parser_);
- }
+class VideoParser
+{
+public:
+ VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue);
- bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream);
+ ~VideoParser()
+ {
+ cuvidDestroyVideoParser(parser_);
+ }
- bool hasError() const { return hasError_; }
+ bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream);
- private:
- VideoParser(const VideoParser&);
- VideoParser& operator =(const VideoParser&);
+ bool hasError() const { return hasError_; }
- VideoDecoder* videoDecoder_;
- FrameQueue* frameQueue_;
- CUvideoparser parser_;
- int unparsedPackets_;
- volatile bool hasError_;
+private:
+ 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 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 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);
- };
- }
-}}
+ // 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__
#include "precomp.hpp"
-#if !defined(HAVE_CUDA) || defined(CUDA_DISABLER) || !defined(HAVE_NVCUVID)
+#ifndef HAVE_NVCUVID
class cv::gpu::VideoReader_GPU::Impl
{
bool cv::gpu::VideoReader_GPU::VideoSource::parseVideoData(const unsigned char*, size_t, bool) { throw_no_cuda(); return false; }
void cv::gpu::VideoReader_GPU::dumpFormat(std::ostream&) { throw_no_cuda(); }
-#else // HAVE_CUDA
-
-#include "frame_queue.h"
-#include "video_decoder.h"
-#include "video_parser.h"
-
-#include "cuvid_video_source.h"
-#include "ffmpeg_video_source.h"
+#else // HAVE_NVCUVID
class cv::gpu::VideoReader_GPU::Impl
{
cv::gpu::VideoReader_GPU::FormatInfo format() const { return videoSource_->format(); }
private:
- Impl(const Impl&);
- Impl& operator =(const Impl&);
-
cv::Ptr<cv::gpu::VideoReader_GPU::VideoSource> videoSource_;
- std::auto_ptr<cv::gpu::detail::FrameQueue> frameQueue_;
- std::auto_ptr<cv::gpu::detail::VideoDecoder> videoDecoder_;
- std::auto_ptr<cv::gpu::detail::VideoParser> videoParser_;
+ cv::Ptr<cv::gpu::detail::FrameQueue> frameQueue_;
+ cv::Ptr<cv::gpu::detail::VideoDecoder> videoDecoder_;
+ cv::Ptr<cv::gpu::detail::VideoParser> videoParser_;
CUvideoctxlock lock_;
cuSafeCall( cuCtxGetCurrent(&ctx) );
cuSafeCall( cuvidCtxLockCreate(&lock_, ctx) );
- frameQueue_.reset(new detail::FrameQueue);
- videoDecoder_.reset(new detail::VideoDecoder(videoSource_->format(), lock_));
- videoParser_.reset(new detail::VideoParser(videoDecoder_.get(), frameQueue_.get()));
+ frameQueue_ = new detail::FrameQueue;
+ videoDecoder_ = new detail::VideoDecoder(videoSource_->format(), lock_);
+ videoParser_ = new detail::VideoParser(videoDecoder_, frameQueue_);
- videoSource_->setFrameQueue(frameQueue_.get());
- videoSource_->setVideoParser(videoParser_.get());
+ videoSource_->setFrameQueue(frameQueue_);
+ videoSource_->setVideoParser(videoParser_);
videoSource_->start();
}
videoSource_->stop();
}
-namespace cv { namespace gpu { namespace cudev {
- namespace video_decoding
- {
- void loadHueCSC(float hueCSC[9]);
- void NV12ToARGB_gpu(const PtrStepb decodedFrame, PtrStepSz<unsigned int> interopFrame, cudaStream_t stream = 0);
- }
+namespace cv { namespace gpu { namespace cudev
+{
+ void loadHueCSC(float hueCSC[9]);
+ void NV12_to_RGB(const PtrStepb decodedFrame, PtrStepSz<uint> interopFrame, cudaStream_t stream = 0);
}}}
namespace
void cudaPostProcessFrame(const cv::gpu::GpuMat& decodedFrame, cv::gpu::GpuMat& interopFrame, int width, int height)
{
- using namespace cv::gpu::cudev::video_decoding;
+ using namespace cv::gpu::cudev;
static bool updateCSC = true;
static float hueColorSpaceMat[9];
loadHueCSC(hueColorSpaceMat);
- NV12ToARGB_gpu(decodedFrame, interopFrame);
+ NV12_to_RGB(decodedFrame, interopFrame);
}
}
{
CV_Assert( !source.empty() );
close();
- impl_.reset(new Impl(source));
+ impl_ = new Impl(source);
}
bool cv::gpu::VideoReader_GPU::isOpened() const
{
- return impl_.get() != 0;
+ return !impl_.empty();
}
void cv::gpu::VideoReader_GPU::close()
{
- impl_.reset();
+ impl_.release();
}
bool cv::gpu::VideoReader_GPU::read(GpuMat& image)
st << "Chroma Format : " << chromas[_format.chromaFormat] << std::endl;
}
-#endif // HAVE_CUDA
+#endif // HAVE_NVCUVID
+
+template <> void cv::Ptr<cv::gpu::VideoReader_GPU::Impl>::delete_obj()
+{
+ if (obj) delete obj;
+}
#include "precomp.hpp"
-#if !defined(HAVE_CUDA) || defined(CUDA_DISABLER) || !defined(HAVE_NVCUVID) || !defined(WIN32)
+#if !defined(HAVE_NVCUVID) || !defined(WIN32)
class cv::gpu::VideoWriter_GPU::Impl
{
#else // !defined HAVE_CUDA || !defined WIN32
-#ifdef HAVE_FFMPEG
- #include "../src/cap_ffmpeg_impl.hpp"
-#else
- #include "../src/cap_ffmpeg_api.hpp"
-#endif
-
-
///////////////////////////////////////////////////////////////////////////
// VideoWriter_GPU::Impl
err = NVGetHWEncodeCaps();
if (err)
- CV_Error(CV_GpuNotSupported, "No CUDA capability present");
+ CV_Error(cv::Error::GpuNotSupported, "No CUDA capability present");
// Create the Encoder API Interface
err = NVCreateEncoder(&encoder_);
};
err = NVSetCodec(encoder_, codecs_id[codec_]);
if (err)
- CV_Error(CV_StsNotImplemented, "Codec format is not supported");
+ CV_Error(cv::Error::StsNotImplemented, "Codec format is not supported");
// Set default params
CV_Assert( err == 0 );
}
-namespace cv { namespace gpu { namespace cudev
-{
- namespace video_encoding
- {
- void YV12_gpu(const PtrStepSzb src, int cn, PtrStepSzb dst);
- }
-}}}
-
namespace
{
// UYVY/YUY2 are both 4:2:2 formats (16bpc)
}
}
+namespace cv { namespace gpu { namespace cudev
+{
+ void RGB_to_YV12(const PtrStepSzb src, int cn, PtrStepSzb dst, cudaStream_t stream = 0);
+}}}
+
void cv::gpu::VideoWriter_GPU::Impl::write(const cv::gpu::GpuMat& frame, bool lastFrame)
{
if (inputFormat_ == SF_BGR)
CV_Assert( res == CUDA_SUCCESS );
if (inputFormat_ == SF_BGR)
- cv::gpu::cudev::video_encoding::YV12_gpu(frame, frame.channels(), videoFrame_);
+ cv::gpu::cudev::RGB_to_YV12(frame, frame.channels(), videoFrame_);
else
{
switch (surfaceFormat_)
void EncoderCallBackFFMPEG::onBeginFrame(int frameNumber, PicType picType)
{
+ (void) frameNumber;
isKeyFrame_ = picType == IFRAME;
}
void EncoderCallBackFFMPEG::onEndFrame(int frameNumber, PicType picType)
{
+ (void) frameNumber;
+ (void) picType;
}
///////////////////////////////////////////////////////////////////////////
void cv::gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format)
{
close();
- impl_.reset(new Impl(encoderCallback, frameSize, fps, format));
+ impl_ = new Impl(encoderCallback, frameSize, fps, format);
}
void cv::gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format)
{
close();
- impl_.reset(new Impl(encoderCallback, frameSize, fps, params, format));
+ impl_ = new Impl(encoderCallback, frameSize, fps, params, format);
}
bool cv::gpu::VideoWriter_GPU::isOpened() const
{
- return impl_.get() != 0;
+ return !impl_.empty();
}
void cv::gpu::VideoWriter_GPU::close()
{
- impl_.reset();
+ impl_.release();
}
void cv::gpu::VideoWriter_GPU::write(const cv::gpu::GpuMat& image, bool lastFrame)
}
#endif // !defined HAVE_CUDA || !defined WIN32
+
+template <> void cv::Ptr<cv::gpu::VideoWriter_GPU::Impl>::delete_obj()
+{
+ if (obj) delete obj;
+}
--- /dev/null
+/*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 "test_precomp.hpp"
+
+CV_TEST_MAIN("gpu")
--- /dev/null
+/*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 "test_precomp.hpp"
--- /dev/null
+/*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*/
+
+#ifdef __GNUC__
+# pragma GCC diagnostic ignored "-Wmissing-declarations"
+# if defined __clang__ || defined __APPLE__
+# pragma GCC diagnostic ignored "-Wmissing-prototypes"
+# pragma GCC diagnostic ignored "-Wextra"
+# endif
+#endif
+
+#ifndef __OPENCV_TEST_PRECOMP_HPP__
+#define __OPENCV_TEST_PRECOMP_HPP__
+
+#include "opencv2/ts.hpp"
+#include "opencv2/ts/gpu_test.hpp"
+
+#include "opencv2/gpucodec.hpp"
+#include "opencv2/highgui.hpp"
+
+#endif
#include "test_precomp.hpp"
-#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#ifdef HAVE_NVCUVID
+
+PARAM_TEST_CASE(Video, cv::gpu::DeviceInfo, std::string)
+{
+};
//////////////////////////////////////////////////////
// VideoReader
-PARAM_TEST_CASE(VideoReader, cv::gpu::DeviceInfo, std::string)
+GPU_TEST_P(Video, Reader)
{
- cv::gpu::DeviceInfo devInfo;
- std::string inputFile;
+ cv::gpu::setDevice(GET_PARAM(0).deviceID());
- virtual void SetUp()
- {
- devInfo = GET_PARAM(0);
- inputFile = GET_PARAM(1);
-
- cv::gpu::setDevice(devInfo.deviceID());
+ const std::string inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + GET_PARAM(1);
- inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + inputFile;
- }
-};
-
-GPU_TEST_P(VideoReader, Regression)
-{
cv::gpu::VideoReader_GPU reader(inputFile);
ASSERT_TRUE(reader.isOpened());
ASSERT_FALSE(reader.isOpened());
}
-INSTANTIATE_TEST_CASE_P(GPU_Video, VideoReader, testing::Combine(
- ALL_DEVICES,
- testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi"))));
-
//////////////////////////////////////////////////////
// VideoWriter
#ifdef WIN32
-PARAM_TEST_CASE(VideoWriter, cv::gpu::DeviceInfo, std::string)
+GPU_TEST_P(Video, Writer)
{
- cv::gpu::DeviceInfo devInfo;
- std::string inputFile;
-
- virtual void SetUp()
- {
- devInfo = GET_PARAM(0);
- inputFile = GET_PARAM(1);
+ cv::gpu::setDevice(GET_PARAM(0).deviceID());
- cv::gpu::setDevice(devInfo.deviceID());
-
- inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + std::string("video/") + inputFile;
- }
-};
+ const std::string inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + GET_PARAM(1);
-GPU_TEST_P(VideoWriter, Regression)
-{
std::string outputFile = cv::tempfile(".avi");
const double FPS = 25.0;
}
}
-INSTANTIATE_TEST_CASE_P(GPU_Video, VideoWriter, testing::Combine(
+#endif // WIN32
+
+INSTANTIATE_TEST_CASE_P(GPU_Codec, Video, testing::Combine(
ALL_DEVICES,
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi"))));
-#endif // WIN32
-
-#endif // defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
+#endif // HAVE_NVCUVID
set(the_description "Super Resolution")
ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 -Wundef)
-ocv_define_module(superres opencv_imgproc opencv_video OPTIONAL opencv_gpu opencv_highgui)
+ocv_define_module(superres opencv_imgproc opencv_video OPTIONAL opencv_gpu opencv_highgui opencv_gpucodec)
//////////////////////////////////////////////////////
// VideoFrameSource_GPU
-#ifndef HAVE_OPENCV_GPU
+#ifndef HAVE_OPENCV_GPUCODEC
Ptr<FrameSource> cv::superres::createFrameSource_Video_GPU(const String& fileName)
{
return Ptr<FrameSource>();
}
-#else // HAVE_OPENCV_GPU
+#else // HAVE_OPENCV_GPUCODEC
namespace
{
return new VideoFrameSource(fileName);
}
-#endif // HAVE_OPENCV_GPU
+#endif // HAVE_OPENCV_GPUCODEC
# include "opencv2/core/gpu_private.hpp"
#endif
+#ifdef HAVE_OPENCV_GPUCODEC
+# include "opencv2/gpucodec.hpp"
+#endif
+
#ifdef HAVE_OPENCV_HIGHGUI
#include "opencv2/highgui.hpp"
#endif
SET(OPENCV_GPU_SAMPLES_REQUIRED_DEPS opencv_core opencv_flann opencv_imgproc opencv_highgui
opencv_ml opencv_video opencv_objdetect opencv_features2d
opencv_calib3d opencv_legacy opencv_contrib opencv_gpu
- opencv_nonfree opencv_softcascade opencv_superres)
+ opencv_nonfree opencv_softcascade opencv_superres
+ opencv_gpucodec)
ocv_check_dependencies(${OPENCV_GPU_SAMPLES_REQUIRED_DEPS})
#include <algorithm>
#include <numeric>
-#include <opencv2/core/core.hpp>
+#include <opencv2/core.hpp>
#include <opencv2/core/opengl.hpp>
-#include <opencv2/gpu/gpu.hpp>
-#include <opencv2/highgui/highgui.hpp>
-#include <opencv2/contrib/contrib.hpp>
+#include <opencv2/gpucodec.hpp>
+#include <opencv2/highgui.hpp>
+#include <opencv2/contrib.hpp>
int main(int argc, const char* argv[])
{
#include <vector>
#include <numeric>
-#include "opencv2/core/core.hpp"
-#include "opencv2/gpu/gpu.hpp"
-#include "opencv2/highgui/highgui.hpp"
-#include "opencv2/contrib/contrib.hpp"
+#include "opencv2/core.hpp"
+#include "opencv2/gpucodec.hpp"
+#include "opencv2/highgui.hpp"
+#include "opencv2/contrib.hpp"
int main(int argc, const char* argv[])
{
projects / platform / upstream / opencv.git / commitdiff