diff --git a/3rdparty/ffmpeg/opencv_ffmpeg.dll b/3rdparty/ffmpeg/opencv_ffmpeg.dll index f4101c027..7aefa5ff7 100644 Binary files a/3rdparty/ffmpeg/opencv_ffmpeg.dll and b/3rdparty/ffmpeg/opencv_ffmpeg.dll differ diff --git a/3rdparty/ffmpeg/opencv_ffmpeg_64.dll b/3rdparty/ffmpeg/opencv_ffmpeg_64.dll index 604f581a1..604f28112 100644 Binary files a/3rdparty/ffmpeg/opencv_ffmpeg_64.dll and b/3rdparty/ffmpeg/opencv_ffmpeg_64.dll differ diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 03399a0c7..ad30beb26 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -46,6 +46,7 @@ #ifndef SKIP_INCLUDES #include #include +#include #endif #include "opencv2/core/gpumat.hpp" @@ -1990,6 +1991,105 @@ private: std::auto_ptr impl_; }; + +////////////////////////////////// Video Decoding ////////////////////////////////////////// + +namespace detail +{ + class FrameQueue; + class VideoParser; +} + +class CV_EXPORTS VideoReader_GPU +{ +public: + enum Codec + { + MPEG1 = 0, + MPEG2, + MPEG4, + VC1, + H264, + JPEG, + H264_SVC, + H264_MVC, + + Uncompressed_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V')), // Y,U,V (4:2:0) + Uncompressed_YV12 = (('Y'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,V,U (4:2:0) + Uncompressed_NV12 = (('N'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,UV (4:2:0) + Uncompressed_YUYV = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V')), // YUYV/YUY2 (4:2:2) + Uncompressed_UYVY = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y')), // UYVY (4:2:2) + }; + + enum ChromaFormat + { + Monochrome=0, + YUV420, + YUV422, + YUV444, + }; + + struct FormatInfo + { + Codec codec; + ChromaFormat chromaFormat; + int width; + int height; + }; + + class VideoSource; + + VideoReader_GPU(); + explicit VideoReader_GPU(const std::string& filename); + explicit VideoReader_GPU(const cv::Ptr& source); + + ~VideoReader_GPU(); + + void open(const std::string& filename); + void open(const cv::Ptr& source); + bool isOpened() const; + + void close(); + + bool read(GpuMat& image); + + FormatInfo format() const; + void dumpFormat(std::ostream& st); + + class VideoSource + { + public: + VideoSource() : frameQueue_(0), videoParser_(0) {} + virtual ~VideoSource() {} + + virtual FormatInfo format() const = 0; + virtual void start() = 0; + virtual void stop() = 0; + virtual bool isStarted() const = 0; + virtual bool hasError() const = 0; + + void setFrameQueue(detail::FrameQueue* frameQueue) { frameQueue_ = frameQueue; } + void setVideoParser(detail::VideoParser* videoParser) { videoParser_ = videoParser; } + + protected: + bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream = false); + + private: + VideoSource(const VideoSource&); + VideoSource& operator =(const VideoSource&); + + detail::FrameQueue* frameQueue_; + detail::VideoParser* videoParser_; + }; + +private: + VideoReader_GPU(const VideoReader_GPU&); + VideoReader_GPU& operator =(const VideoReader_GPU&); + + class Impl; + std::auto_ptr impl_; +}; + } // namespace gpu } // namespace cv diff --git a/modules/gpu/src/cu_safe_call.cpp b/modules/gpu/src/cu_safe_call.cpp new file mode 100644 index 000000000..0fb658697 --- /dev/null +++ b/modules/gpu/src/cu_safe_call.cpp @@ -0,0 +1,139 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "cu_safe_call.h" + +#ifdef HAVE_CUDA + +namespace +{ + #define error_entry(entry) { entry, #entry } + + struct ErrorEntry + { + int code; + std::string str; + }; + + class ErrorEntryComparer + { + public: + inline ErrorEntryComparer(int code) : code_(code) {} + + inline bool operator()(const ErrorEntry& e) const { return e.code == code_; } + + private: + int code_; + }; + + std::string getErrorString(int code, const ErrorEntry* errors, size_t n) + { + size_t idx = std::find_if(errors, errors + n, ErrorEntryComparer(code)) - errors; + + const std::string& msg = (idx != n) ? errors[idx].str : std::string("Unknown error code"); + + std::ostringstream ostr; + ostr << msg << " [Code = " << code << "]"; + + return ostr.str(); + } + + const ErrorEntry cu_errors [] = + { + error_entry( CUDA_SUCCESS ), + error_entry( CUDA_ERROR_INVALID_VALUE ), + error_entry( CUDA_ERROR_OUT_OF_MEMORY ), + error_entry( CUDA_ERROR_NOT_INITIALIZED ), + error_entry( CUDA_ERROR_DEINITIALIZED ), + error_entry( CUDA_ERROR_PROFILER_DISABLED ), + error_entry( CUDA_ERROR_PROFILER_NOT_INITIALIZED ), + error_entry( CUDA_ERROR_PROFILER_ALREADY_STARTED ), + error_entry( CUDA_ERROR_PROFILER_ALREADY_STOPPED ), + error_entry( CUDA_ERROR_NO_DEVICE ), + error_entry( CUDA_ERROR_INVALID_DEVICE ), + error_entry( CUDA_ERROR_INVALID_IMAGE ), + error_entry( CUDA_ERROR_INVALID_CONTEXT ), + error_entry( CUDA_ERROR_CONTEXT_ALREADY_CURRENT ), + error_entry( CUDA_ERROR_MAP_FAILED ), + error_entry( CUDA_ERROR_UNMAP_FAILED ), + error_entry( CUDA_ERROR_ARRAY_IS_MAPPED ), + error_entry( CUDA_ERROR_ALREADY_MAPPED ), + error_entry( CUDA_ERROR_NO_BINARY_FOR_GPU ), + error_entry( CUDA_ERROR_ALREADY_ACQUIRED ), + error_entry( CUDA_ERROR_NOT_MAPPED ), + error_entry( CUDA_ERROR_NOT_MAPPED_AS_ARRAY ), + error_entry( CUDA_ERROR_NOT_MAPPED_AS_POINTER ), + error_entry( CUDA_ERROR_ECC_UNCORRECTABLE ), + error_entry( CUDA_ERROR_UNSUPPORTED_LIMIT ), + error_entry( CUDA_ERROR_CONTEXT_ALREADY_IN_USE ), + error_entry( CUDA_ERROR_INVALID_SOURCE ), + error_entry( CUDA_ERROR_FILE_NOT_FOUND ), + error_entry( CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND ), + error_entry( CUDA_ERROR_SHARED_OBJECT_INIT_FAILED ), + error_entry( CUDA_ERROR_OPERATING_SYSTEM ), + error_entry( CUDA_ERROR_INVALID_HANDLE ), + error_entry( CUDA_ERROR_NOT_FOUND ), + error_entry( CUDA_ERROR_NOT_READY ), + error_entry( CUDA_ERROR_LAUNCH_FAILED ), + error_entry( CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES ), + error_entry( CUDA_ERROR_LAUNCH_TIMEOUT ), + error_entry( CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING ), + error_entry( CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED ), + error_entry( CUDA_ERROR_PEER_ACCESS_NOT_ENABLED ), + error_entry( CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE ), + error_entry( CUDA_ERROR_CONTEXT_IS_DESTROYED ), + error_entry( CUDA_ERROR_ASSERT ), + error_entry( CUDA_ERROR_TOO_MANY_PEERS ), + error_entry( CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED ), + error_entry( CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED ), + error_entry( CUDA_ERROR_UNKNOWN ) + }; + + const size_t cu_errors_num = sizeof(cu_errors) / sizeof(cu_errors[0]); +} + +std::string cv::gpu::detail::cuGetErrString(CUresult res) +{ + return getErrorString(res, cu_errors, cu_errors_num); +} + +#endif // HAVE_CUDA diff --git a/modules/gpu/src/cu_safe_call.h b/modules/gpu/src/cu_safe_call.h new file mode 100644 index 000000000..7e9313d34 --- /dev/null +++ b/modules/gpu/src/cu_safe_call.h @@ -0,0 +1,67 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __CU_SAFE_CALL_H__ +#define __CU_SAFE_CALL_H__ + +#include "precomp.hpp" + +#ifdef HAVE_CUDA + +namespace cv { namespace gpu { + namespace detail + { + std::string cuGetErrString(CUresult res); + + inline void cuSafeCall_impl(CUresult res, const char* file, int line) + { + if (res != CUDA_SUCCESS) + cv::error( cv::Exception(CV_GpuApiCallError, cuGetErrString(res), "unknown function", file, line) ); + } + } +}} + +#define cuSafeCall( op ) cv::gpu::detail::cuSafeCall_impl( (op), __FILE__, __LINE__ ) + +#endif // HAVE_CUDA + +#endif // __CU_SAFE_CALL_H__ diff --git a/modules/gpu/src/cuda/NV12ToARGB.cu b/modules/gpu/src/cuda/NV12ToARGB.cu new file mode 100644 index 000000000..8089fca0d --- /dev/null +++ b/modules/gpu/src/cuda/NV12ToARGB.cu @@ -0,0 +1,208 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +/* + * Copyright 1993-2010 NVIDIA Corporation. All rights reserved. + * + * Please refer to the NVIDIA end user license agreement (EULA) associated + * with this source code for terms and conditions that govern your use of + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +/* + NV12ToARGB color space conversion CUDA kernel + + This sample uses CUDA to perform a simple NV12 (YUV 4:2:0 planar) + source and converts to output in ARGB format +*/ + +#include "opencv2/gpu/device/common.hpp" + +namespace cv { namespace gpu { namespace device { + namespace video_decoding + { + __constant__ uint constAlpha = ((uint)0xff << 24); + + __constant__ float constHueColorSpaceMat[9]; + + void loadHueCSC(float hueCSC[9]) + { + cudaSafeCall( cudaMemcpyToSymbol(constHueColorSpaceMat, hueCSC, 9 * sizeof(float)) ); + } + + __device__ void YUV2RGB(const uint* yuvi, float* red, float* green, float* blue) + { + float luma, chromaCb, chromaCr; + + // Prepare for hue adjustment + luma = (float)yuvi[0]; + chromaCb = (float)((int)yuvi[1] - 512.0f); + chromaCr = (float)((int)yuvi[2] - 512.0f); + + // Convert YUV To RGB with hue adjustment + *red = (luma * constHueColorSpaceMat[0]) + + (chromaCb * constHueColorSpaceMat[1]) + + (chromaCr * constHueColorSpaceMat[2]); + + *green = (luma * constHueColorSpaceMat[3]) + + (chromaCb * constHueColorSpaceMat[4]) + + (chromaCr * constHueColorSpaceMat[5]); + + *blue = (luma * constHueColorSpaceMat[6]) + + (chromaCb * constHueColorSpaceMat[7]) + + (chromaCr * constHueColorSpaceMat[8]); + } + + __device__ uint RGBAPACK_10bit(float red, float green, float blue, uint alpha) + { + uint ARGBpixel = 0; + + // Clamp final 10 bit results + red = ::fmin(::fmax(red, 0.0f), 1023.f); + green = ::fmin(::fmax(green, 0.0f), 1023.f); + blue = ::fmin(::fmax(blue, 0.0f), 1023.f); + + // Convert to 8 bit unsigned integers per color component + ARGBpixel = (((uint)blue >> 2) | + (((uint)green >> 2) << 8) | + (((uint)red >> 2) << 16) | + (uint)alpha); + + return ARGBpixel; + } + + // CUDA kernel for outputing the final ARGB output from NV12 + + #define COLOR_COMPONENT_BIT_SIZE 10 + #define COLOR_COMPONENT_MASK 0x3FF + + __global__ void NV12ToARGB(uchar* srcImage, size_t nSourcePitch, + uint* dstImage, size_t nDestPitch, + uint width, uint height) + { + // Pad borders with duplicate pixels, and we multiply by 2 because we process 2 pixels per thread + const int x = blockIdx.x * (blockDim.x << 1) + (threadIdx.x << 1); + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x >= width || y >= height) + return; + + // Read 2 Luma components at a time, so we don't waste processing since CbCr are decimated this way. + // if we move to texture we could read 4 luminance values + + uint yuv101010Pel[2]; + + yuv101010Pel[0] = (srcImage[y * nSourcePitch + x ]) << 2; + yuv101010Pel[1] = (srcImage[y * nSourcePitch + x + 1]) << 2; + + const size_t chromaOffset = nSourcePitch * height; + + const int y_chroma = y >> 1; + + if (y & 1) // odd scanline ? + { + uint chromaCb = srcImage[chromaOffset + y_chroma * nSourcePitch + x ]; + uint chromaCr = srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1]; + + if (y_chroma < ((height >> 1) - 1)) // interpolate chroma vertically + { + chromaCb = (chromaCb + srcImage[chromaOffset + (y_chroma + 1) * nSourcePitch + x ] + 1) >> 1; + chromaCr = (chromaCr + srcImage[chromaOffset + (y_chroma + 1) * nSourcePitch + x + 1] + 1) >> 1; + } + + yuv101010Pel[0] |= (chromaCb << ( COLOR_COMPONENT_BIT_SIZE + 2)); + yuv101010Pel[0] |= (chromaCr << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2)); + + yuv101010Pel[1] |= (chromaCb << ( COLOR_COMPONENT_BIT_SIZE + 2)); + yuv101010Pel[1] |= (chromaCr << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2)); + } + else + { + yuv101010Pel[0] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x ] << ( COLOR_COMPONENT_BIT_SIZE + 2)); + yuv101010Pel[0] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1] << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2)); + + yuv101010Pel[1] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x ] << ( COLOR_COMPONENT_BIT_SIZE + 2)); + yuv101010Pel[1] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1] << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2)); + } + + // this steps performs the color conversion + uint yuvi[6]; + float red[2], green[2], blue[2]; + + yuvi[0] = (yuv101010Pel[0] & COLOR_COMPONENT_MASK ); + yuvi[1] = ((yuv101010Pel[0] >> COLOR_COMPONENT_BIT_SIZE) & COLOR_COMPONENT_MASK); + yuvi[2] = ((yuv101010Pel[0] >> (COLOR_COMPONENT_BIT_SIZE << 1)) & COLOR_COMPONENT_MASK); + + yuvi[3] = (yuv101010Pel[1] & COLOR_COMPONENT_MASK ); + yuvi[4] = ((yuv101010Pel[1] >> COLOR_COMPONENT_BIT_SIZE) & COLOR_COMPONENT_MASK); + yuvi[5] = ((yuv101010Pel[1] >> (COLOR_COMPONENT_BIT_SIZE << 1)) & COLOR_COMPONENT_MASK); + + // YUV to RGB Transformation conversion + YUV2RGB(&yuvi[0], &red[0], &green[0], &blue[0]); + YUV2RGB(&yuvi[3], &red[1], &green[1], &blue[1]); + + // Clamp the results to RGBA + + const size_t dstImagePitch = nDestPitch >> 2; + + dstImage[y * dstImagePitch + x ] = RGBAPACK_10bit(red[0], green[0], blue[0], constAlpha); + dstImage[y * dstImagePitch + x + 1 ] = RGBAPACK_10bit(red[1], green[1], blue[1], constAlpha); + } + + void NV12ToARGB_gpu(const PtrStepb decodedFrame, DevMem2D_ interopFrame, cudaStream_t stream) + { + dim3 block(32, 8); + dim3 grid(divUp(interopFrame.cols, 2 * block.x), divUp(interopFrame.rows, block.y)); + + NV12ToARGB<<>>(decodedFrame.data, decodedFrame.step, interopFrame.data, interopFrame.step, + interopFrame.cols, interopFrame.rows); + + cudaSafeCall( cudaGetLastError() ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } + } +}}} diff --git a/modules/gpu/src/cuvid_video_source.cpp b/modules/gpu/src/cuvid_video_source.cpp new file mode 100644 index 000000000..9099744d3 --- /dev/null +++ b/modules/gpu/src/cuvid_video_source.cpp @@ -0,0 +1,63 @@ +#include "cuvid_video_source.h" +#include "cu_safe_call.h" + +#if defined(HAVE_CUDA) && !defined(__APPLE__) + +cv::gpu::detail::CuvidVideoSource::CuvidVideoSource(const std::string& fname) +{ + CUVIDSOURCEPARAMS params; + std::memset(¶ms, 0, sizeof(CUVIDSOURCEPARAMS)); + + // Fill parameter struct + params.pUserData = this; // will be passed to data handlers + params.pfnVideoDataHandler = HandleVideoData; // our local video-handler callback + params.pfnAudioDataHandler = 0; + + // now create the actual source + CUresult res = cuvidCreateVideoSource(&videoSource_, fname.c_str(), ¶ms); + if (res == CUDA_ERROR_INVALID_SOURCE) + throw std::runtime_error("Unsupported video source"); + cuSafeCall( res ); + + CUVIDEOFORMAT vidfmt; + cuSafeCall( cuvidGetSourceVideoFormat(videoSource_, &vidfmt, 0) ); + + format_.codec = static_cast(vidfmt.codec); + format_.chromaFormat = static_cast(vidfmt.chroma_format); + format_.width = vidfmt.coded_width; + format_.height = vidfmt.coded_height; +} + +cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::detail::CuvidVideoSource::format() const +{ + return format_; +} + +void cv::gpu::detail::CuvidVideoSource::start() +{ + cuSafeCall( cuvidSetVideoSourceState(videoSource_, cudaVideoState_Started) ); +} + +void cv::gpu::detail::CuvidVideoSource::stop() +{ + cuSafeCall( cuvidSetVideoSourceState(videoSource_, cudaVideoState_Stopped) ); +} + +bool cv::gpu::detail::CuvidVideoSource::isStarted() const +{ + return (cuvidGetVideoSourceState(videoSource_) == cudaVideoState_Started); +} + +bool cv::gpu::detail::CuvidVideoSource::hasError() const +{ + return (cuvidGetVideoSourceState(videoSource_) == cudaVideoState_Error); +} + +int CUDAAPI cv::gpu::detail::CuvidVideoSource::HandleVideoData(void* userData, CUVIDSOURCEDATAPACKET* packet) +{ + CuvidVideoSource* thiz = static_cast(userData); + + return thiz->parseVideoData(packet->payload, packet->payload_size, (packet->flags & CUVID_PKT_ENDOFSTREAM) != 0); +} + +#endif // defined(HAVE_CUDA) && !defined(__APPLE__) diff --git a/modules/gpu/src/cuvid_video_source.h b/modules/gpu/src/cuvid_video_source.h new file mode 100644 index 000000000..95830b7e5 --- /dev/null +++ b/modules/gpu/src/cuvid_video_source.h @@ -0,0 +1,90 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __CUVUD_VIDEO_SOURCE_H__ +#define __CUVUD_VIDEO_SOURCE_H__ + +#include "precomp.hpp" + +#if defined(HAVE_CUDA) && !defined(__APPLE__) + +namespace cv { namespace gpu +{ + namespace detail + { + class CuvidVideoSource : public VideoReader_GPU::VideoSource + { + public: + explicit CuvidVideoSource(const std::string& fname); + ~CuvidVideoSource() { cuvidDestroyVideoSource(videoSource_); } + + VideoReader_GPU::FormatInfo format() const; + void start(); + void stop(); + bool isStarted() const; + bool hasError() const; + + private: + CuvidVideoSource(const CuvidVideoSource&); + CuvidVideoSource& operator =(const CuvidVideoSource&); + + // Callback for handling packages of demuxed video data. + // + // Parameters: + // pUserData - Pointer to user data. We must pass a pointer to a + // VideoSourceData struct here, that contains a valid CUvideoparser + // and FrameQueue. + // pPacket - video-source data packet. + // + // NOTE: called from a different thread that doesn't not have a cuda context + // + static int CUDAAPI HandleVideoData(void* pUserData, CUVIDSOURCEDATAPACKET* pPacket); + + CUvideosource videoSource_; + VideoReader_GPU::FormatInfo format_; + }; + } +}} + +#endif // defined(HAVE_CUDA) && !defined(__APPLE__) + +#endif // __CUVUD_VIDEO_SOURCE_H__ diff --git a/modules/gpu/src/ffmpeg_video_source.cpp b/modules/gpu/src/ffmpeg_video_source.cpp new file mode 100644 index 000000000..4d1c6a51f --- /dev/null +++ b/modules/gpu/src/ffmpeg_video_source.cpp @@ -0,0 +1,185 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "ffmpeg_video_source.h" + +#ifdef HAVE_CUDA + +#ifdef HAVE_FFMPEG + #ifdef NEW_FFMPEG + #include "cap_ffmpeg_impl_v2.hpp" + #else + #include "cap_ffmpeg_impl.hpp" + #endif +#else + #include "cap_ffmpeg_api.hpp" +#endif + +namespace +{ + Create_InputMediaStream_FFMPEG_Plugin create_InputMediaStream_FFMPEG_p = 0; + Release_InputMediaStream_FFMPEG_Plugin release_InputMediaStream_FFMPEG_p = 0; + Read_InputMediaStream_FFMPEG_Plugin read_InputMediaStream_FFMPEG_p = 0; + + bool init_MediaStream_FFMPEG() + { + static bool initialized = 0; + + if (!initialized) + { + #if defined WIN32 || defined _WIN32 + const char* module_name = "opencv_ffmpeg" + #if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__) + "_64" + #endif + ".dll"; + + static HMODULE cvFFOpenCV = LoadLibrary(module_name); + + if (cvFFOpenCV) + { + create_InputMediaStream_FFMPEG_p = + (Create_InputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "create_InputMediaStream_FFMPEG"); + release_InputMediaStream_FFMPEG_p = + (Release_InputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "release_InputMediaStream_FFMPEG"); + read_InputMediaStream_FFMPEG_p = + (Read_InputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "read_InputMediaStream_FFMPEG"); + + initialized = create_InputMediaStream_FFMPEG_p != 0 && release_InputMediaStream_FFMPEG_p != 0 && read_InputMediaStream_FFMPEG_p != 0; + } + #elif defined HAVE_FFMPEG + create_InputMediaStream_FFMPEG_p = create_InputMediaStream_FFMPEG; + release_InputMediaStream_FFMPEG_p = release_InputMediaStream_FFMPEG; + read_InputMediaStream_FFMPEG_p = read_InputMediaStream_FFMPEG; + + initialized = true; + #endif + } + + return initialized; + } +} + +cv::gpu::detail::FFmpegVideoSource::FFmpegVideoSource(const std::string& fname) : + stream_(0) +{ + CV_Assert( init_MediaStream_FFMPEG() ); + + int codec; + int chroma_format; + int width; + int height; + + stream_ = create_InputMediaStream_FFMPEG_p(fname.c_str(), &codec, &chroma_format, &width, &height); + if (!stream_) + CV_Error(CV_StsUnsupportedFormat, "Unsupported video source"); + + format_.codec = static_cast(codec); + format_.chromaFormat = static_cast(chroma_format); + format_.width = width; + format_.height = height; +} + +cv::gpu::detail::FFmpegVideoSource::~FFmpegVideoSource() +{ + release_InputMediaStream_FFMPEG_p(stream_); +} + +cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::detail::FFmpegVideoSource::format() const +{ + return format_; +} + +void cv::gpu::detail::FFmpegVideoSource::start() +{ + stop_ = false; + hasError_ = false; + thread_.reset(new Thread(readLoop, this)); +} + +void cv::gpu::detail::FFmpegVideoSource::stop() +{ + stop_ = true; + thread_->wait(); + thread_.reset(); +} + +bool cv::gpu::detail::FFmpegVideoSource::isStarted() const +{ + return !stop_; +} + +bool cv::gpu::detail::FFmpegVideoSource::hasError() const +{ + return hasError_; +} + +void cv::gpu::detail::FFmpegVideoSource::readLoop(void* userData) +{ + FFmpegVideoSource* thiz = static_cast(userData); + + for (;;) + { + unsigned char* data; + int size; + int endOfFile; + + if (!read_InputMediaStream_FFMPEG_p(thiz->stream_, &data, &size, &endOfFile)) + { + thiz->hasError_ = !endOfFile; + break; + } + + if (!thiz->parseVideoData(data, size)) + { + thiz->hasError_ = true; + break; + } + + if (thiz->stop_) + break; + } + + thiz->parseVideoData(0, 0, true); +} + +#endif // HAVE_CUDA diff --git a/modules/gpu/src/ffmpeg_video_source.h b/modules/gpu/src/ffmpeg_video_source.h new file mode 100644 index 000000000..dd267fcf9 --- /dev/null +++ b/modules/gpu/src/ffmpeg_video_source.h @@ -0,0 +1,88 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __FFMPEG_VIDEO_SOURCE_H__ +#define __FFMPEG_VIDEO_SOURCE_H__ + +#include "precomp.hpp" +#include "thread_wrappers.h" + +#ifdef HAVE_CUDA + +struct InputMediaStream_FFMPEG; + +namespace cv { namespace gpu +{ + namespace detail + { + class FFmpegVideoSource : public VideoReader_GPU::VideoSource + { + public: + FFmpegVideoSource(const std::string& fname); + ~FFmpegVideoSource(); + + VideoReader_GPU::FormatInfo format() const; + void start(); + void stop(); + bool isStarted() const; + bool hasError() const; + + private: + FFmpegVideoSource(const FFmpegVideoSource&); + FFmpegVideoSource& operator =(const FFmpegVideoSource&); + + VideoReader_GPU::FormatInfo format_; + + InputMediaStream_FFMPEG* stream_; + + std::auto_ptr thread_; + volatile bool stop_; + volatile bool hasError_; + + static void readLoop(void* userData); + }; + } +}} + +#endif // HAVE_CUDA + +#endif // __CUVUD_VIDEO_SOURCE_H__ diff --git a/modules/gpu/src/frame_queue.cpp b/modules/gpu/src/frame_queue.cpp new file mode 100644 index 000000000..4ba310658 --- /dev/null +++ b/modules/gpu/src/frame_queue.cpp @@ -0,0 +1,113 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "frame_queue.h" + +cv::gpu::detail::FrameQueue::FrameQueue() : + endOfDecode_(0), + framesInQueue_(0), + readPosition_(0) +{ + std::memset(displayQueue_, 0, sizeof(displayQueue_)); + std::memset((void*)isFrameInUse_, 0, sizeof(isFrameInUse_)); +} + +bool cv::gpu::detail::FrameQueue::waitUntilFrameAvailable(int pictureIndex) +{ + while (isInUse(pictureIndex)) + { + // Decoder is getting too far ahead from display + Thread::sleep(1); + + if (isEndOfDecode()) + return false; + } + + return true; +} + +void cv::gpu::detail::FrameQueue::enqueue(const CUVIDPARSERDISPINFO* picParams) +{ + // Mark the frame as 'in-use' so we don't re-use it for decoding until it is no longer needed + // for display + isFrameInUse_[picParams->picture_index] = true; + + // Wait until we have a free entry in the display queue (should never block if we have enough entries) + do + { + bool isFramePlaced = false; + + { + CriticalSection::AutoLock autoLock(criticalSection_); + + if (framesInQueue_ < MaximumSize) + { + int writePosition = (readPosition_ + framesInQueue_) % MaximumSize; + displayQueue_[writePosition] = *picParams; + framesInQueue_++; + isFramePlaced = true; + } + } + + if (isFramePlaced) // Done + break; + + // Wait a bit + Thread::sleep(1); + } while (!isEndOfDecode()); +} + +bool cv::gpu::detail::FrameQueue::dequeue(CUVIDPARSERDISPINFO& displayInfo) +{ + CriticalSection::AutoLock autoLock(criticalSection_); + + if (framesInQueue_ > 0) + { + int entry = readPosition_; + displayInfo = displayQueue_[entry]; + readPosition_ = (entry + 1) % MaximumSize; + framesInQueue_--; + return true; + } + + return false; +} diff --git a/modules/gpu/src/frame_queue.h b/modules/gpu/src/frame_queue.h new file mode 100644 index 000000000..040060cde --- /dev/null +++ b/modules/gpu/src/frame_queue.h @@ -0,0 +1,103 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __FRAME_QUEUE_H__ +#define __FRAME_QUEUE_H__ + +#include "precomp.hpp" +#include "thread_wrappers.h" + +#ifdef HAVE_CUDA + +namespace cv { namespace gpu +{ + namespace detail + { + class FrameQueue + { + public: + static const int MaximumSize = 20; // MAX_FRM_CNT; + + FrameQueue(); + + void endDecode() { endOfDecode_ = true; } + bool isEndOfDecode() const { return endOfDecode_ != 0;} + + // Spins until frame becomes available or decoding gets canceled. + // If the requested frame is available the method returns true. + // If decoding was interupted before the requested frame becomes + // available, the method returns false. + bool waitUntilFrameAvailable(int pictureIndex); + + void enqueue(const CUVIDPARSERDISPINFO* picParams); + + // Deque the next frame. + // Parameters: + // displayInfo - New frame info gets placed into this object. + // Returns: + // true, if a new frame was returned, + // false, if the queue was empty and no new frame could be returned. + bool dequeue(CUVIDPARSERDISPINFO& displayInfo); + + void releaseFrame(const CUVIDPARSERDISPINFO& picParams) { isFrameInUse_[picParams.picture_index] = false; } + + private: + FrameQueue(const FrameQueue&); + FrameQueue& operator =(const FrameQueue&); + + bool isInUse(int pictureIndex) const { return isFrameInUse_[pictureIndex] != 0; } + + CriticalSection criticalSection_; + + volatile int isFrameInUse_[MaximumSize]; + volatile int endOfDecode_; + + int framesInQueue_; + int readPosition_; + CUVIDPARSERDISPINFO displayQueue_[MaximumSize]; + }; + } +}} + +#endif // HAVE_CUDA + +#endif // __FRAME_QUEUE_H__ diff --git a/modules/gpu/src/precomp.hpp b/modules/gpu/src/precomp.hpp index b95ea613a..4c66f9c2c 100644 --- a/modules/gpu/src/precomp.hpp +++ b/modules/gpu/src/precomp.hpp @@ -51,6 +51,7 @@ #include "cvconfig.h" #endif +#include #include #include #include @@ -60,6 +61,10 @@ #include #include #include +#include +#include +#include +#include #include "opencv2/gpu/gpu.hpp" #include "opencv2/imgproc/imgproc.hpp" diff --git a/modules/gpu/src/thread_wrappers.cpp b/modules/gpu/src/thread_wrappers.cpp new file mode 100644 index 000000000..5a35258a0 --- /dev/null +++ b/modules/gpu/src/thread_wrappers.cpp @@ -0,0 +1,254 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "thread_wrappers.h" + +#ifdef HAVE_CUDA + +#ifdef WIN32 + #define NOMINMAX + #include +#else + #include + #include +#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(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(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 diff --git a/modules/gpu/src/thread_wrappers.h b/modules/gpu/src/thread_wrappers.h new file mode 100644 index 000000000..e96957f6f --- /dev/null +++ b/modules/gpu/src/thread_wrappers.h @@ -0,0 +1,112 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __THREAD_WRAPPERS_H__ +#define __THREAD_WRAPPERS_H__ + +#include "precomp.hpp" + +#ifdef HAVE_CUDA + +namespace cv { namespace gpu +{ + namespace detail + { + class CriticalSection + { + public: + CriticalSection(); + ~CriticalSection(); + + void enter(); + void leave(); + + class AutoLock + { + public: + explicit AutoLock(CriticalSection& criticalSection) : + criticalSection_(criticalSection) + { + criticalSection_.enter(); + } + + ~AutoLock() + { + criticalSection_.leave(); + } + + private: + CriticalSection& criticalSection_; + }; + + private: + CriticalSection(const CriticalSection&); + CriticalSection& operator=(const CriticalSection&); + + class Impl; + std::auto_ptr impl_; + }; + + 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_; + }; + + } +}} + +#endif // HAVE_CUDA + +#endif // __THREAD_WRAPPERS_H__ diff --git a/modules/gpu/src/video_decoder.cpp b/modules/gpu/src/video_decoder.cpp new file mode 100644 index 000000000..d1a7cbe1c --- /dev/null +++ b/modules/gpu/src/video_decoder.cpp @@ -0,0 +1,116 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "video_decoder.h" +#include "frame_queue.h" + +#ifdef HAVE_CUDA + +void cv::gpu::detail::VideoDecoder::create(const VideoReader_GPU::FormatInfo& videoFormat) +{ + release(); + + cudaVideoCodec codec = static_cast(videoFormat.codec); + cudaVideoChromaFormat chromaFormat = static_cast(videoFormat.chromaFormat); + + cudaVideoCreateFlags videoCreateFlags = (codec == cudaVideoCodec_JPEG || codec == cudaVideoCodec_MPEG2) ? + cudaVideoCreate_PreferCUDA : + cudaVideoCreate_PreferCUVID; + + // Validate video format. These are the currently supported formats via NVCUVID + CV_Assert(cudaVideoCodec_MPEG1 == codec || + cudaVideoCodec_MPEG2 == codec || + cudaVideoCodec_MPEG4 == codec || + cudaVideoCodec_VC1 == codec || + cudaVideoCodec_H264 == codec || + cudaVideoCodec_JPEG == codec || + cudaVideoCodec_YUV420== codec || + cudaVideoCodec_YV12 == codec || + cudaVideoCodec_NV12 == codec || + cudaVideoCodec_YUYV == codec || + cudaVideoCodec_UYVY == codec ); + + CV_Assert(cudaVideoChromaFormat_Monochrome == chromaFormat || + cudaVideoChromaFormat_420 == chromaFormat || + cudaVideoChromaFormat_422 == chromaFormat || + cudaVideoChromaFormat_444 == chromaFormat); + + // Fill the decoder-create-info struct from the given video-format struct. + std::memset(&createInfo_, 0, sizeof(CUVIDDECODECREATEINFO)); + + // Create video decoder + createInfo_.CodecType = codec; + createInfo_.ulWidth = videoFormat.width; + createInfo_.ulHeight = videoFormat.height; + createInfo_.ulNumDecodeSurfaces = FrameQueue::MaximumSize; + + // Limit decode memory to 24MB (16M pixels at 4:2:0 = 24M bytes) + while (createInfo_.ulNumDecodeSurfaces * videoFormat.width * videoFormat.height > 16 * 1024 * 1024) + createInfo_.ulNumDecodeSurfaces--; + + createInfo_.ChromaFormat = chromaFormat; + createInfo_.OutputFormat = cudaVideoSurfaceFormat_NV12; + createInfo_.DeinterlaceMode = cudaVideoDeinterlaceMode_Adaptive; + + // No scaling + static const int MAX_FRAME_COUNT = 2; + + createInfo_.ulTargetWidth = createInfo_.ulWidth; + createInfo_.ulTargetHeight = createInfo_.ulHeight; + createInfo_.ulNumOutputSurfaces = MAX_FRAME_COUNT; // We won't simultaneously map more than 8 surfaces + createInfo_.ulCreationFlags = videoCreateFlags; + createInfo_.vidLock = lock_; + + // create the decoder + cuSafeCall( cuvidCreateDecoder(&decoder_, &createInfo_) ); +} + +void cv::gpu::detail::VideoDecoder::release() +{ + if (decoder_) + { + cuvidDestroyDecoder(decoder_); + decoder_ = 0; + } +} + +#endif // HAVE_CUDA diff --git a/modules/gpu/src/video_decoder.h b/modules/gpu/src/video_decoder.h new file mode 100644 index 000000000..335d214c5 --- /dev/null +++ b/modules/gpu/src/video_decoder.h @@ -0,0 +1,117 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __VIDEO_DECODER_H__ +#define __VIDEO_DECODER_H__ + +#include "precomp.hpp" +#include "cu_safe_call.h" + +#ifdef HAVE_CUDA + +namespace cv { namespace gpu +{ + namespace detail + { + class VideoDecoder + { + public: + VideoDecoder(const VideoReader_GPU::FormatInfo& videoFormat, CUvideoctxlock lock) : lock_(lock), decoder_(0) + { + create(videoFormat); + } + + ~VideoDecoder() + { + release(); + } + + void create(const VideoReader_GPU::FormatInfo& videoFormat); + void release(); + + // Get the code-type currently used. + cudaVideoCodec codec() const { return createInfo_.CodecType; } + unsigned long maxDecodeSurfaces() const { return createInfo_.ulNumDecodeSurfaces; } + + unsigned long frameWidth() const { return createInfo_.ulWidth; } + unsigned long frameHeight() const { return createInfo_.ulHeight; } + + unsigned long targetWidth() const { return createInfo_.ulTargetWidth; } + unsigned long targetHeight() const { return createInfo_.ulTargetHeight; } + + cudaVideoChromaFormat chromaFormat() const { return createInfo_.ChromaFormat; } + + bool decodePicture(CUVIDPICPARAMS* picParams) + { + return cuvidDecodePicture(decoder_, picParams) == CUDA_SUCCESS; + } + + cv::gpu::GpuMat mapFrame(int picIdx, CUVIDPROCPARAMS& videoProcParams) + { + CUdeviceptr ptr; + unsigned int pitch; + + cuSafeCall( cuvidMapVideoFrame(decoder_, picIdx, &ptr, &pitch, &videoProcParams) ); + + return GpuMat(targetHeight() * 3 / 2, targetWidth(), CV_8UC1, (void*) ptr, pitch); + } + + void unmapFrame(cv::gpu::GpuMat& frame) + { + cuSafeCall( cuvidUnmapVideoFrame(decoder_, (CUdeviceptr) frame.data) ); + frame.release(); + } + + private: + VideoDecoder(const VideoDecoder&); + VideoDecoder& operator =(const VideoDecoder&); + + CUvideoctxlock lock_; + CUVIDDECODECREATEINFO createInfo_; + CUvideodecoder decoder_; + }; + } +}} + +#endif // HAVE_CUDA + +#endif // __VIDEO_DECODER_H__ diff --git a/modules/gpu/src/video_parser.cpp b/modules/gpu/src/video_parser.cpp new file mode 100644 index 000000000..dabbdbcd4 --- /dev/null +++ b/modules/gpu/src/video_parser.cpp @@ -0,0 +1,162 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "video_parser.h" +#include "cu_safe_call.h" + +#ifdef HAVE_CUDA + +cv::gpu::detail::VideoParser::VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue) : + videoDecoder_(videoDecoder), frameQueue_(frameQueue), unparsedPackets_(0), hasError_(false) +{ + CUVIDPARSERPARAMS params; + memset(¶ms, 0, sizeof(CUVIDPARSERPARAMS)); + + params.CodecType = videoDecoder->codec(); + params.ulMaxNumDecodeSurfaces = videoDecoder->maxDecodeSurfaces(); + params.ulMaxDisplayDelay = 1; // this flag is needed so the parser will push frames out to the decoder as quickly as it can + params.pUserData = this; + params.pfnSequenceCallback = HandleVideoSequence; // Called before decoding frames and/or whenever there is a format change + params.pfnDecodePicture = HandlePictureDecode; // Called when a picture is ready to be decoded (decode order) + params.pfnDisplayPicture = HandlePictureDisplay; // Called whenever a picture is ready to be displayed (display order) + + cuSafeCall( cuvidCreateVideoParser(&parser_, ¶ms) ); +} + +bool cv::gpu::detail::VideoParser::parseVideoData(const unsigned char* data, size_t size, bool endOfStream) +{ + CUVIDSOURCEDATAPACKET packet; + std::memset(&packet, 0, sizeof(CUVIDSOURCEDATAPACKET)); + + if (endOfStream) + packet.flags |= CUVID_PKT_ENDOFSTREAM; + + packet.payload_size = size; + packet.payload = data; + + if (cuvidParseVideoData(parser_, &packet) != CUDA_SUCCESS) + { + hasError_ = true; + frameQueue_->endDecode(); + return false; + } + + const int maxUnparsedPackets = 15; + + ++unparsedPackets_; + if (unparsedPackets_ > maxUnparsedPackets) + { + hasError_ = true; + frameQueue_->endDecode(); + return false; + } + + if (endOfStream) + frameQueue_->endDecode(); + + return !frameQueue_->isEndOfDecode(); +} + +int CUDAAPI cv::gpu::detail::VideoParser::HandleVideoSequence(void* userData, CUVIDEOFORMAT* format) +{ + VideoParser* thiz = static_cast(userData); + + thiz->unparsedPackets_ = 0; + + if (format->codec != thiz->videoDecoder_->codec() || + format->coded_width != thiz->videoDecoder_->frameWidth() || + format->coded_height != thiz->videoDecoder_->frameHeight() || + format->chroma_format != thiz->videoDecoder_->chromaFormat()) + { + VideoReader_GPU::FormatInfo newFormat; + + newFormat.codec = static_cast(format->codec); + newFormat.chromaFormat = static_cast(format->chroma_format); + newFormat.width = format->coded_width; + newFormat.height = format->coded_height; + + try + { + thiz->videoDecoder_->create(newFormat); + } + catch (const cv::Exception&) + { + thiz->hasError_ = true; + return false; + } + } + + return true; +} + +int CUDAAPI cv::gpu::detail::VideoParser::HandlePictureDecode(void* userData, CUVIDPICPARAMS* picParams) +{ + VideoParser* thiz = static_cast(userData); + + thiz->unparsedPackets_ = 0; + + bool isFrameAvailable = thiz->frameQueue_->waitUntilFrameAvailable(picParams->CurrPicIdx); + + if (!isFrameAvailable) + return false; + + if (!thiz->videoDecoder_->decodePicture(picParams)) + { + thiz->hasError_ = true; + return false; + } + + return true; +} + +int CUDAAPI cv::gpu::detail::VideoParser::HandlePictureDisplay(void* userData, CUVIDPARSERDISPINFO* picParams) +{ + VideoParser* thiz = static_cast(userData); + + thiz->unparsedPackets_ = 0; + + thiz->frameQueue_->enqueue(picParams); + + return true; +} + +#endif // HAVE_CUDA diff --git a/modules/gpu/src/video_parser.h b/modules/gpu/src/video_parser.h new file mode 100644 index 000000000..abbbeaccb --- /dev/null +++ b/modules/gpu/src/video_parser.h @@ -0,0 +1,100 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __VIDEO_PARSER_H__ +#define __VIDEO_PARSER_H__ + +#include "precomp.hpp" + +#include "frame_queue.h" +#include "video_decoder.h" + +#ifdef HAVE_CUDA + +namespace cv { namespace gpu +{ + namespace detail + { + class VideoParser + { + public: + VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue); + + ~VideoParser() + { + cuvidDestroyVideoParser(parser_); + } + + bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream); + + bool hasError() const { return hasError_; } + + private: + VideoParser(const VideoParser&); + VideoParser& operator =(const VideoParser&); + + VideoDecoder* videoDecoder_; + FrameQueue* frameQueue_; + CUvideoparser parser_; + int unparsedPackets_; + volatile bool hasError_; + + // Called when the decoder encounters a video format change (or initial sequence header) + // This particular implementation of the callback returns 0 in case the video format changes + // to something different than the original format. Returning 0 causes a stop of the app. + static int CUDAAPI HandleVideoSequence(void* pUserData, CUVIDEOFORMAT* pFormat); + + // Called by the video parser to decode a single picture + // Since the parser will deliver data as fast as it can, we need to make sure that the picture + // index we're attempting to use for decode is no longer used for display + static int CUDAAPI HandlePictureDecode(void* pUserData, CUVIDPICPARAMS* pPicParams); + + // Called by the video parser to display a video frame (in the case of field pictures, there may be + // 2 decode calls per 1 display call, since two fields make up one frame) + static int CUDAAPI HandlePictureDisplay(void* pUserData, CUVIDPARSERDISPINFO* pPicParams); + }; + } +}} + +#endif // HAVE_CUDA + +#endif // __VIDEO_PARSER_H__ diff --git a/modules/gpu/src/video_reader.cpp b/modules/gpu/src/video_reader.cpp new file mode 100644 index 000000000..fc2843c01 --- /dev/null +++ b/modules/gpu/src/video_reader.cpp @@ -0,0 +1,397 @@ +/*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" + +#ifndef HAVE_CUDA + +cv::gpu::VideoReader_GPU::VideoReader_GPU() { throw_nogpu(); } +cv::gpu::VideoReader_GPU::VideoReader_GPU(const std::string&) { throw_nogpu(); } +cv::gpu::VideoReader_GPU::VideoReader_GPU(const cv::Ptr&) { throw_nogpu(); } +cv::gpu::VideoReader_GPU::~VideoReader_GPU() { } +void cv::gpu::VideoReader_GPU::open(const std::string&) { throw_nogpu(); } +void cv::gpu::VideoReader_GPU::open(const cv::Ptr&) { throw_nogpu(); } +bool cv::gpu::VideoReader_GPU::isOpened() const { return false; } +void cv::gpu::VideoReader_GPU::close() { } +bool cv::gpu::VideoReader_GPU::read(GpuMat&) { throw_nogpu(); return false; } +cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::VideoReader_GPU::format() const { throw_nogpu(); FormatInfo format; return format; } +bool cv::gpu::VideoReader_GPU::VideoSource::parseVideoData(const unsigned char*, size_t, bool) { throw_nogpu(); return false; } +void cv::gpu::VideoReader_GPU::dumpFormat(std::ostream&) { throw_nogpu(); } + +#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" + +#include "cu_safe_call.h" + +class cv::gpu::VideoReader_GPU::Impl +{ +public: + explicit Impl(const cv::Ptr& source); + ~Impl(); + + bool grab(cv::gpu::GpuMat& frame); + + cv::gpu::VideoReader_GPU::FormatInfo format() const { return videoSource_->format(); } + +private: + Impl(const Impl&); + Impl& operator =(const Impl&); + + cv::Ptr videoSource_; + + std::auto_ptr frameQueue_; + std::auto_ptr videoDecoder_; + std::auto_ptr videoParser_; + + CUvideoctxlock lock_; + + std::deque< std::pair > frames_; +}; + +cv::gpu::VideoReader_GPU::Impl::Impl(const cv::Ptr& source) : + videoSource_(source), + lock_(0) +{ + // init context + GpuMat temp(1, 1, CV_8UC1); + temp.release(); + + DeviceInfo devInfo; + CV_Assert( devInfo.supports(FEATURE_SET_COMPUTE_11) ); + + CUcontext ctx; + 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())); + + videoSource_->setFrameQueue(frameQueue_.get()); + videoSource_->setVideoParser(videoParser_.get()); + + videoSource_->start(); +} + +cv::gpu::VideoReader_GPU::Impl::~Impl() +{ + frameQueue_->endDecode(); + videoSource_->stop(); +} + +namespace cv { namespace gpu { namespace device { + namespace video_decoding + { + void loadHueCSC(float hueCSC[9]); + void NV12ToARGB_gpu(const PtrStepb decodedFrame, DevMem2D_ interopFrame, cudaStream_t stream = 0); + } +}}} + +namespace +{ + class VideoCtxAutoLock + { + public: + VideoCtxAutoLock(CUvideoctxlock lock) : m_lock(lock) { cuSafeCall( cuvidCtxLock(m_lock, 0) ); } + ~VideoCtxAutoLock() { cuvidCtxUnlock(m_lock, 0); } + + private: + CUvideoctxlock m_lock; + }; + + enum ColorSpace + { + ITU601 = 1, + ITU709 = 2 + }; + + void setColorSpaceMatrix(ColorSpace CSC, float hueCSC[9], float hue) + { + float hueSin = std::sin(hue); + float hueCos = std::cos(hue); + + if (CSC == ITU601) + { + //CCIR 601 + hueCSC[0] = 1.1644f; + hueCSC[1] = hueSin * 1.5960f; + hueCSC[2] = hueCos * 1.5960f; + hueCSC[3] = 1.1644f; + hueCSC[4] = (hueCos * -0.3918f) - (hueSin * 0.8130f); + hueCSC[5] = (hueSin * 0.3918f) - (hueCos * 0.8130f); + hueCSC[6] = 1.1644f; + hueCSC[7] = hueCos * 2.0172f; + hueCSC[8] = hueSin * -2.0172f; + } + else if (CSC == ITU709) + { + //CCIR 709 + hueCSC[0] = 1.0f; + hueCSC[1] = hueSin * 1.57480f; + hueCSC[2] = hueCos * 1.57480f; + hueCSC[3] = 1.0; + hueCSC[4] = (hueCos * -0.18732f) - (hueSin * 0.46812f); + hueCSC[5] = (hueSin * 0.18732f) - (hueCos * 0.46812f); + hueCSC[6] = 1.0f; + hueCSC[7] = hueCos * 1.85560f; + hueCSC[8] = hueSin * -1.85560f; + } + } + + void cudaPostProcessFrame(const cv::gpu::GpuMat& decodedFrame, cv::gpu::GpuMat& interopFrame, int width, int height) + { + using namespace cv::gpu::device::video_decoding; + + static bool updateCSC = true; + static float hueColorSpaceMat[9]; + + // Upload the Color Space Conversion Matrices + if (updateCSC) + { + const ColorSpace colorSpace = ITU601; + const float hue = 0.0f; + + // CCIR 601/709 + setColorSpaceMatrix(colorSpace, hueColorSpaceMat, hue); + + updateCSC = false; + } + + // Final Stage: NV12toARGB color space conversion + + interopFrame.create(height, width, CV_8UC4); + + loadHueCSC(hueColorSpaceMat); + + NV12ToARGB_gpu(decodedFrame, interopFrame); + } +} + +bool cv::gpu::VideoReader_GPU::Impl::grab(GpuMat& frame) +{ + if (videoSource_->hasError() || videoParser_->hasError()) + CV_Error(CV_StsUnsupportedFormat, "Unsupported video source"); + + if (!videoSource_->isStarted() || frameQueue_->isEndOfDecode()) + return false; + + if (frames_.empty()) + { + CUVIDPARSERDISPINFO displayInfo; + + for (;;) + { + if (frameQueue_->dequeue(displayInfo)) + break; + + if (videoSource_->hasError() || videoParser_->hasError()) + CV_Error(CV_StsUnsupportedFormat, "Unsupported video source"); + + if (frameQueue_->isEndOfDecode()) + return false; + + // Wait a bit + detail::Thread::sleep(1); + } + + bool isProgressive = displayInfo.progressive_frame != 0; + const int num_fields = isProgressive ? 1 : 2 + displayInfo.repeat_first_field; + + for (int active_field = 0; active_field < num_fields; ++active_field) + { + CUVIDPROCPARAMS videoProcParams; + std::memset(&videoProcParams, 0, sizeof(CUVIDPROCPARAMS)); + + videoProcParams.progressive_frame = displayInfo.progressive_frame; + videoProcParams.second_field = active_field; + videoProcParams.top_field_first = displayInfo.top_field_first; + videoProcParams.unpaired_field = (num_fields == 1); + + frames_.push_back(std::make_pair(displayInfo, videoProcParams)); + } + } + + if (frames_.empty()) + return false; + + std::pair frameInfo = frames_.front(); + frames_.pop_front(); + + { + VideoCtxAutoLock autoLock(lock_); + + // map decoded video frame to CUDA surface + cv::gpu::GpuMat decodedFrame = videoDecoder_->mapFrame(frameInfo.first.picture_index, frameInfo.second); + + // perform post processing on the CUDA surface (performs colors space conversion and post processing) + // comment this out if we inclue the line of code seen above + cudaPostProcessFrame(decodedFrame, frame, videoDecoder_->targetWidth(), videoDecoder_->targetHeight()); + + // unmap video frame + // unmapFrame() synchronizes with the VideoDecode API (ensures the frame has finished decoding) + videoDecoder_->unmapFrame(decodedFrame); + } + + // release the frame, so it can be re-used in decoder + if (frames_.empty()) + frameQueue_->releaseFrame(frameInfo.first); + + return true; +} + +//////////////////////////////////////////////////////////////////////////// + +cv::gpu::VideoReader_GPU::VideoReader_GPU() +{ +} + +cv::gpu::VideoReader_GPU::VideoReader_GPU(const std::string& filename) +{ + open(filename); +} + +cv::gpu::VideoReader_GPU::VideoReader_GPU(const cv::Ptr& source) +{ + open(source); +} + +cv::gpu::VideoReader_GPU::~VideoReader_GPU() +{ + close(); +} + +void cv::gpu::VideoReader_GPU::open(const std::string& filename) +{ + CV_Assert( !filename.empty() ); + +#ifndef __APPLE__ + try + { + cv::Ptr source(new detail::CuvidVideoSource(filename)); + open(source); + } + catch (const std::runtime_error&) +#endif + { + cv::Ptr source(new cv::gpu::detail::FFmpegVideoSource(filename)); + open(source); + } +} + +void cv::gpu::VideoReader_GPU::open(const cv::Ptr& source) +{ + CV_Assert( !source.empty() ); + close(); + impl_.reset(new Impl(source)); +} + +bool cv::gpu::VideoReader_GPU::isOpened() const +{ + return impl_.get() != 0; +} + +void cv::gpu::VideoReader_GPU::close() +{ + impl_.reset(); +} + +bool cv::gpu::VideoReader_GPU::read(GpuMat& image) +{ + if (!isOpened()) + return false; + + if (!impl_->grab(image)) + { + close(); + return false; + } + + return true; +} + +cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::VideoReader_GPU::format() const +{ + CV_Assert( isOpened() ); + return impl_->format(); +} + +bool cv::gpu::VideoReader_GPU::VideoSource::parseVideoData(const unsigned char* data, size_t size, bool endOfStream) +{ + return videoParser_->parseVideoData(data, size, endOfStream); +} + +void cv::gpu::VideoReader_GPU::dumpFormat(std::ostream& st) +{ + static const char* codecs[] = + { + "MPEG1", + "MPEG2", + "MPEG4", + "VC1", + "H264", + "JPEG", + "H264_SVC", + "H264_MVC" + }; + + static const char* chromas[] = + { + "Monochrome", + "YUV420", + "YUV422", + "YUV444" + }; + + FormatInfo format = this->format(); + + st << "Frame Size : " << format.width << "x" << format.height << std::endl; + st << "Codec : " << (format.codec <= H264_MVC ? codecs[format.codec] : "Uncompressed YUV") << std::endl; + st << "Chroma Format : " << chromas[format.chromaFormat] << std::endl; +} + +#endif // HAVE_CUDA diff --git a/modules/gpu/src/video_writer.cpp b/modules/gpu/src/video_writer.cpp index fc7f7185e..451b64fd8 100644 --- a/modules/gpu/src/video_writer.cpp +++ b/modules/gpu/src/video_writer.cpp @@ -577,7 +577,7 @@ void cv::gpu::VideoWriter_GPU::Impl::write(const cv::gpu::GpuMat& frame, bool la switch (surfaceFormat_) { case UYVY: // UYVY (4:2:2) - case YUY2: // YUY2 (4:2:2) + case YUY2: // YUY2 (4:2:2) copyUYVYorYUY2Frame(frameSize_, frame, videoFrame_); break; @@ -662,33 +662,33 @@ namespace if (!initialized) { - #if defined WIN32 || defined _WIN32 - const char* module_name = "opencv_ffmpeg" - #if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__) - "_64" + #if defined WIN32 || defined _WIN32 + const char* module_name = "opencv_ffmpeg" + #if (defined _MSC_VER && defined _M_X64) || (defined __GNUC__ && defined __x86_64__) + "_64" + #endif + ".dll"; + + static HMODULE cvFFOpenCV = LoadLibrary(module_name); + + if (cvFFOpenCV) + { + create_OutputMediaStream_FFMPEG_p = + (Create_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "create_OutputMediaStream_FFMPEG"); + release_OutputMediaStream_FFMPEG_p = + (Release_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "release_OutputMediaStream_FFMPEG"); + write_OutputMediaStream_FFMPEG_p = + (Write_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "write_OutputMediaStream_FFMPEG"); + + initialized = create_OutputMediaStream_FFMPEG_p != 0 && release_OutputMediaStream_FFMPEG_p != 0 && write_OutputMediaStream_FFMPEG_p != 0; + } + #elif defined HAVE_FFMPEG + create_OutputMediaStream_FFMPEG_p = create_OutputMediaStream_FFMPEG; + release_OutputMediaStream_FFMPEG_p = release_OutputMediaStream_FFMPEG; + write_OutputMediaStream_FFMPEG_p = write_OutputMediaStream_FFMPEG; + + initialized = true; #endif - ".dll"; - - static HMODULE cvFFOpenCV = LoadLibrary(module_name); - - if (cvFFOpenCV) - { - create_OutputMediaStream_FFMPEG_p = - (Create_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "create_OutputMediaStream_FFMPEG"); - release_OutputMediaStream_FFMPEG_p = - (Release_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "release_OutputMediaStream_FFMPEG"); - write_OutputMediaStream_FFMPEG_p = - (Write_OutputMediaStream_FFMPEG_Plugin)GetProcAddress(cvFFOpenCV, "write_OutputMediaStream_FFMPEG"); - - initialized = create_OutputMediaStream_FFMPEG_p != 0 && release_OutputMediaStream_FFMPEG_p != 0 && write_OutputMediaStream_FFMPEG_p != 0; - } - #elif defined HAVE_FFMPEG - create_OutputMediaStream_FFMPEG_p = create_OutputMediaStream_FFMPEG; - release_OutputMediaStream_FFMPEG_p = release_OutputMediaStream_FFMPEG; - write_OutputMediaStream_FFMPEG_p = write_OutputMediaStream_FFMPEG; - - initialized = true; - #endif } return initialized; diff --git a/modules/gpu/test/test_video.cpp b/modules/gpu/test/test_video.cpp index 2d59f5540..7288a393b 100644 --- a/modules/gpu/test/test_video.cpp +++ b/modules/gpu/test/test_video.cpp @@ -387,6 +387,8 @@ INSTANTIATE_TEST_CASE_P(GPU_Video, FarnebackOpticalFlow, testing::Combine( ///////////////////////////////////////////////////////////////////////////////////////////////// // VideoWriter +#ifdef WIN32 + PARAM_TEST_CASE(VideoWriter, cv::gpu::DeviceInfo, std::string) { cv::gpu::DeviceInfo devInfo; @@ -440,10 +442,58 @@ TEST_P(VideoWriter, Regression) reader.open(outputFile); ASSERT_TRUE( reader.isOpened() ); + + for (int i = 0; i < 5; ++i) + { + reader >> frame; + ASSERT_FALSE( frame.empty() ); + } } INSTANTIATE_TEST_CASE_P(GPU_Video, VideoWriter, testing::Combine( ALL_DEVICES, - testing::Values("VID00003-20100701-2204.3GP", "big_buck_bunny.mpg"))); + testing::Values(std::string("VID00003-20100701-2204.mpg"), std::string("big_buck_bunny.mpg")))); + +#endif // WIN32 + +///////////////////////////////////////////////////////////////////////////////////////////////// +// VideoReader + +PARAM_TEST_CASE(VideoReader, cv::gpu::DeviceInfo, std::string) +{ + cv::gpu::DeviceInfo devInfo; + std::string inputFile; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + inputFile = GET_PARAM(1); + + cv::gpu::setDevice(devInfo.deviceID()); + + inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + inputFile; + } +}; + +TEST_P(VideoReader, Regression) +{ + cv::gpu::VideoReader_GPU reader(inputFile); + ASSERT_TRUE( reader.isOpened() ); + + cv::gpu::GpuMat frame; + + for (int i = 0; i < 5; ++i) + { + ASSERT_TRUE( reader.read(frame) ); + ASSERT_FALSE( frame.empty() ); + } + + reader.close(); + ASSERT_FALSE( reader.isOpened() ); +} + +INSTANTIATE_TEST_CASE_P(GPU_Video, VideoReader, testing::Combine( + ALL_DEVICES, + testing::Values(std::string("VID00003-20100701-2204.mpg")))); } // namespace diff --git a/modules/highgui/src/cap_ffmpeg_api.hpp b/modules/highgui/src/cap_ffmpeg_api.hpp index 3ce80d4b8..b585214dd 100644 --- a/modules/highgui/src/cap_ffmpeg_api.hpp +++ b/modules/highgui/src/cap_ffmpeg_api.hpp @@ -77,6 +77,18 @@ typedef struct OutputMediaStream_FFMPEG* (*Create_OutputMediaStream_FFMPEG_Plugi typedef void (*Release_OutputMediaStream_FFMPEG_Plugin)(struct OutputMediaStream_FFMPEG* stream); typedef void (*Write_OutputMediaStream_FFMPEG_Plugin)(struct OutputMediaStream_FFMPEG* stream, unsigned char* data, int size, int keyFrame); +/* + * For CUDA decoder + */ + +OPENCV_FFMPEG_API struct InputMediaStream_FFMPEG* create_InputMediaStream_FFMPEG(const char* fileName, int* codec, int* chroma_format, int* width, int* height); +OPENCV_FFMPEG_API void release_InputMediaStream_FFMPEG(struct InputMediaStream_FFMPEG* stream); +OPENCV_FFMPEG_API int read_InputMediaStream_FFMPEG(struct InputMediaStream_FFMPEG* stream, unsigned char** data, int* size, int* endOfFile); + +typedef struct InputMediaStream_FFMPEG* (*Create_InputMediaStream_FFMPEG_Plugin)(const char* fileName, int* codec, int* chroma_format, int* width, int* height); +typedef void (*Release_InputMediaStream_FFMPEG_Plugin)(struct InputMediaStream_FFMPEG* stream); +typedef int (*Read_InputMediaStream_FFMPEG_Plugin)(struct InputMediaStream_FFMPEG* stream, unsigned char** data, int* size, int* endOfFile); + #ifdef __cplusplus } #endif diff --git a/modules/highgui/src/cap_ffmpeg_impl.hpp b/modules/highgui/src/cap_ffmpeg_impl.hpp index b269543aa..40c7377b8 100644 --- a/modules/highgui/src/cap_ffmpeg_impl.hpp +++ b/modules/highgui/src/cap_ffmpeg_impl.hpp @@ -1741,3 +1741,232 @@ void write_OutputMediaStream_FFMPEG(struct OutputMediaStream_FFMPEG* stream, uns { stream->write(data, size, keyFrame); } + +/* + * For CUDA decoder + */ + +enum +{ + VideoCodec_MPEG1 = 0, + VideoCodec_MPEG2, + VideoCodec_MPEG4, + VideoCodec_VC1, + VideoCodec_H264, + VideoCodec_JPEG, + VideoCodec_H264_SVC, + VideoCodec_H264_MVC, + + // Uncompressed YUV + VideoCodec_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V')), // Y,U,V (4:2:0) + VideoCodec_YV12 = (('Y'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,V,U (4:2:0) + VideoCodec_NV12 = (('N'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,UV (4:2:0) + VideoCodec_YUYV = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V')), // YUYV/YUY2 (4:2:2) + VideoCodec_UYVY = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y')), // UYVY (4:2:2) +}; + +enum +{ + VideoChromaFormat_Monochrome = 0, + VideoChromaFormat_YUV420, + VideoChromaFormat_YUV422, + VideoChromaFormat_YUV444, +}; + +struct InputMediaStream_FFMPEG +{ +public: + bool open(const char* fileName, int* codec, int* chroma_format, int* width, int* height); + void close(); + + bool read(unsigned char** data, int* size, int* endOfFile); + +private: + InputMediaStream_FFMPEG(const InputMediaStream_FFMPEG&); + InputMediaStream_FFMPEG& operator =(const InputMediaStream_FFMPEG&); + + AVFormatContext* ctx_; + int video_stream_id_; + AVPacket pkt_; +}; + +bool InputMediaStream_FFMPEG::open(const char* fileName, int* codec, int* chroma_format, int* width, int* height) +{ + int err; + + ctx_ = 0; + video_stream_id_ = -1; + memset(&pkt_, 0, sizeof(AVPacket)); + + #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 13, 0) + avformat_network_init(); + #endif + + // register all codecs, demux and protocols + av_register_all(); + + av_log_set_level(AV_LOG_ERROR); + + #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 6, 0) + err = avformat_open_input(&ctx_, fileName, 0, 0); + #else + err = av_open_input_file(&ctx_, fileName, 0, 0, 0); + #endif + if (err < 0) + return false; + + #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 3, 0) + err = avformat_find_stream_info(ctx_, 0); + #else + err = av_find_stream_info(ctx_); + #endif + if (err < 0) + return false; + + for (unsigned int i = 0; i < ctx_->nb_streams; ++i) + { + #if LIBAVFORMAT_BUILD > 4628 + AVCodecContext *enc = ctx_->streams[i]->codec; + #else + AVCodecContext *enc = &ctx_->streams[i]->codec; + #endif + + if (enc->codec_type == AVMEDIA_TYPE_VIDEO) + { + video_stream_id_ = static_cast(i); + + switch (enc->codec_id) + { + case CODEC_ID_MPEG1VIDEO: + *codec = ::VideoCodec_MPEG1; + break; + + case CODEC_ID_MPEG2VIDEO: + *codec = ::VideoCodec_MPEG2; + break; + + case CODEC_ID_MPEG4: + *codec = ::VideoCodec_MPEG4; + break; + + case CODEC_ID_VC1: + *codec = ::VideoCodec_VC1; + break; + + case CODEC_ID_H264: + *codec = ::VideoCodec_H264; + break; + + default: + return false; + }; + + switch (enc->pix_fmt) + { + case PIX_FMT_YUV420P: + *chroma_format = ::VideoChromaFormat_YUV420; + break; + + case PIX_FMT_YUV422P: + *chroma_format = ::VideoChromaFormat_YUV422; + break; + + case PIX_FMT_YUV444P: + *chroma_format = ::VideoChromaFormat_YUV444; + break; + + default: + return false; + } + + *width = enc->coded_width; + *height = enc->coded_height; + + break; + } + } + + if (video_stream_id_ < 0) + return false; + + av_init_packet(&pkt_); + + return true; +} + +void InputMediaStream_FFMPEG::close() +{ + if (ctx_) + { + #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 24, 2) + avformat_close_input(&ctx_); + #else + av_close_input_file(ctx_); + #endif + } + + // free last packet if exist + if (pkt_.data) + av_free_packet(&pkt_); +} + +bool InputMediaStream_FFMPEG::read(unsigned char** data, int* size, int* endOfFile) +{ + // free last packet if exist + if (pkt_.data) + av_free_packet(&pkt_); + + // get the next frame + for (;;) + { + int ret = av_read_frame(ctx_, &pkt_); + + if (ret == AVERROR(EAGAIN)) + continue; + + if (ret < 0) + { + if (ret == AVERROR_EOF) + *endOfFile = true; + return false; + } + + if (pkt_.stream_index != video_stream_id_) + { + av_free_packet(&pkt_); + continue; + } + + break; + } + + *data = pkt_.data; + *size = pkt_.size; + *endOfFile = false; + + return true; +} + +InputMediaStream_FFMPEG* create_InputMediaStream_FFMPEG(const char* fileName, int* codec, int* chroma_format, int* width, int* height) +{ + InputMediaStream_FFMPEG* stream = (InputMediaStream_FFMPEG*) malloc(sizeof(InputMediaStream_FFMPEG)); + + if (stream && stream->open(fileName, codec, chroma_format, width, height)) + return stream; + + stream->close(); + free(stream); + + return 0; +} + +void release_InputMediaStream_FFMPEG(InputMediaStream_FFMPEG* stream) +{ + stream->close(); + free(stream); +} + +int read_InputMediaStream_FFMPEG(InputMediaStream_FFMPEG* stream, unsigned char** data, int* size, int* endOfFile) +{ + return stream->read(data, size, endOfFile); +} diff --git a/modules/highgui/src/cap_ffmpeg_impl_v2.hpp b/modules/highgui/src/cap_ffmpeg_impl_v2.hpp index 829820838..555ff8f0b 100755 --- a/modules/highgui/src/cap_ffmpeg_impl_v2.hpp +++ b/modules/highgui/src/cap_ffmpeg_impl_v2.hpp @@ -1907,3 +1907,232 @@ void write_OutputMediaStream_FFMPEG(struct OutputMediaStream_FFMPEG* stream, uns { stream->write(data, size, keyFrame); } + +/* + * For CUDA decoder + */ + +enum +{ + VideoCodec_MPEG1 = 0, + VideoCodec_MPEG2, + VideoCodec_MPEG4, + VideoCodec_VC1, + VideoCodec_H264, + VideoCodec_JPEG, + VideoCodec_H264_SVC, + VideoCodec_H264_MVC, + + // Uncompressed YUV + VideoCodec_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V')), // Y,U,V (4:2:0) + VideoCodec_YV12 = (('Y'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,V,U (4:2:0) + VideoCodec_NV12 = (('N'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,UV (4:2:0) + VideoCodec_YUYV = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V')), // YUYV/YUY2 (4:2:2) + VideoCodec_UYVY = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y')), // UYVY (4:2:2) +}; + +enum +{ + VideoChromaFormat_Monochrome = 0, + VideoChromaFormat_YUV420, + VideoChromaFormat_YUV422, + VideoChromaFormat_YUV444, +}; + +struct InputMediaStream_FFMPEG +{ +public: + bool open(const char* fileName, int* codec, int* chroma_format, int* width, int* height); + void close(); + + bool read(unsigned char** data, int* size, int* endOfFile); + +private: + InputMediaStream_FFMPEG(const InputMediaStream_FFMPEG&); + InputMediaStream_FFMPEG& operator =(const InputMediaStream_FFMPEG&); + + AVFormatContext* ctx_; + int video_stream_id_; + AVPacket pkt_; +}; + +bool InputMediaStream_FFMPEG::open(const char* fileName, int* codec, int* chroma_format, int* width, int* height) +{ + int err; + + ctx_ = 0; + video_stream_id_ = -1; + memset(&pkt_, 0, sizeof(AVPacket)); + + #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 13, 0) + avformat_network_init(); + #endif + + // register all codecs, demux and protocols + av_register_all(); + + av_log_set_level(AV_LOG_ERROR); + + #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 6, 0) + err = avformat_open_input(&ctx_, fileName, 0, 0); + #else + err = av_open_input_file(&ctx_, fileName, 0, 0, 0); + #endif + if (err < 0) + return false; + + #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 3, 0) + err = avformat_find_stream_info(ctx_, 0); + #else + err = av_find_stream_info(ctx_); + #endif + if (err < 0) + return false; + + for (unsigned int i = 0; i < ctx_->nb_streams; ++i) + { + #if LIBAVFORMAT_BUILD > 4628 + AVCodecContext *enc = ctx_->streams[i]->codec; + #else + AVCodecContext *enc = &ctx_->streams[i]->codec; + #endif + + if (enc->codec_type == AVMEDIA_TYPE_VIDEO) + { + video_stream_id_ = static_cast(i); + + switch (enc->codec_id) + { + case CODEC_ID_MPEG1VIDEO: + *codec = ::VideoCodec_MPEG1; + break; + + case CODEC_ID_MPEG2VIDEO: + *codec = ::VideoCodec_MPEG2; + break; + + case CODEC_ID_MPEG4: + *codec = ::VideoCodec_MPEG4; + break; + + case CODEC_ID_VC1: + *codec = ::VideoCodec_VC1; + break; + + case CODEC_ID_H264: + *codec = ::VideoCodec_H264; + break; + + default: + return false; + }; + + switch (enc->pix_fmt) + { + case PIX_FMT_YUV420P: + *chroma_format = ::VideoChromaFormat_YUV420; + break; + + case PIX_FMT_YUV422P: + *chroma_format = ::VideoChromaFormat_YUV422; + break; + + case PIX_FMT_YUV444P: + *chroma_format = ::VideoChromaFormat_YUV444; + break; + + default: + return false; + } + + *width = enc->coded_width; + *height = enc->coded_height; + + break; + } + } + + if (video_stream_id_ < 0) + return false; + + av_init_packet(&pkt_); + + return true; +} + +void InputMediaStream_FFMPEG::close() +{ + if (ctx_) + { + #if LIBAVFORMAT_BUILD >= CALC_FFMPEG_VERSION(53, 24, 2) + avformat_close_input(&ctx_); + #else + av_close_input_file(ctx_); + #endif + } + + // free last packet if exist + if (pkt_.data) + av_free_packet(&pkt_); +} + +bool InputMediaStream_FFMPEG::read(unsigned char** data, int* size, int* endOfFile) +{ + // free last packet if exist + if (pkt_.data) + av_free_packet(&pkt_); + + // get the next frame + for (;;) + { + int ret = av_read_frame(ctx_, &pkt_); + + if (ret == AVERROR(EAGAIN)) + continue; + + if (ret < 0) + { + if (ret == AVERROR_EOF) + *endOfFile = true; + return false; + } + + if (pkt_.stream_index != video_stream_id_) + { + av_free_packet(&pkt_); + continue; + } + + break; + } + + *data = pkt_.data; + *size = pkt_.size; + *endOfFile = false; + + return true; +} + +InputMediaStream_FFMPEG* create_InputMediaStream_FFMPEG(const char* fileName, int* codec, int* chroma_format, int* width, int* height) +{ + InputMediaStream_FFMPEG* stream = (InputMediaStream_FFMPEG*) malloc(sizeof(InputMediaStream_FFMPEG)); + + if (stream && stream->open(fileName, codec, chroma_format, width, height)) + return stream; + + stream->close(); + free(stream); + + return 0; +} + +void release_InputMediaStream_FFMPEG(InputMediaStream_FFMPEG* stream) +{ + stream->close(); + free(stream); +} + +int read_InputMediaStream_FFMPEG(InputMediaStream_FFMPEG* stream, unsigned char** data, int* size, int* endOfFile) +{ + return stream->read(data, size, endOfFile); +} diff --git a/samples/gpu/brox_optical_flow.cpp b/samples/gpu/brox_optical_flow.cpp index b6104117b..824ab272c 100644 --- a/samples/gpu/brox_optical_flow.cpp +++ b/samples/gpu/brox_optical_flow.cpp @@ -215,7 +215,7 @@ int main(int argc, const char* argv[]) switch (key) { case 27: - break; + return 0; case 'A': if (currentFrame > 0) @@ -243,8 +243,6 @@ int main(int argc, const char* argv[]) cerr << "Unknow error" << endl; return -1; } - - return 0; } template inline T clamp (T x, T a, T b) diff --git a/samples/gpu/video_reader.cpp b/samples/gpu/video_reader.cpp new file mode 100644 index 000000000..3e86b4d85 --- /dev/null +++ b/samples/gpu/video_reader.cpp @@ -0,0 +1,71 @@ +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +int main(int argc, const char* argv[]) +{ + if (argc != 2) + return -1; + + const std::string fname(argv[1]); + + cv::namedWindow("CPU", cv::WINDOW_NORMAL); + cv::namedWindow("GPU", cv::WINDOW_OPENGL); + cv::gpu::setGlDevice(); + + cv::Mat frame; + cv::VideoCapture reader(fname); + + cv::gpu::GpuMat d_frame; + cv::gpu::VideoReader_GPU d_reader(fname); + d_reader.dumpFormat(std::cout); + + cv::TickMeter tm; + std::vector cpu_times; + std::vector gpu_times; + + for (;;) + { + tm.reset(); tm.start(); + if (!reader.read(frame)) + break; + tm.stop(); + cpu_times.push_back(tm.getTimeMilli()); + + tm.reset(); tm.start(); + if (!d_reader.read(d_frame)) + break; + tm.stop(); + gpu_times.push_back(tm.getTimeMilli()); + + cv::imshow("CPU", frame); + cv::imshow("GPU", frame); + + if (cv::waitKey(3) > 0) + break; + } + + if (!cpu_times.empty() && !gpu_times.empty()) + { + std::cout << std::endl << "Results:" << std::endl; + + std::sort(cpu_times.begin(), cpu_times.end()); + std::sort(gpu_times.begin(), gpu_times.end()); + + double cpu_avg = std::accumulate(cpu_times.begin(), cpu_times.end(), 0.0) / cpu_times.size(); + double gpu_avg = std::accumulate(gpu_times.begin(), gpu_times.end(), 0.0) / gpu_times.size(); + + std::cout << "CPU : Avg : " << cpu_avg << " ms FPS : " << 1000.0 / cpu_avg << std::endl; + std::cout << "GPU : Avg : " << gpu_avg << " ms FPS : " << 1000.0 / gpu_avg << std::endl; + } + + return 0; +}