generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

renamed gpucodec -> cudacodec

Browse Source
This commit is contained in:
Vladislav Vinogradov
2013-07-23 14:29:21 +04:00
parent cfe4a71dc6
commit dbeb3e2968
39 changed files with 33 additions and 33 deletions
Download Patch File Download Diff File Expand all files Collapse all files

188
modules/cudacodec/src/cuda/nv12_to_rgb.cu Normal file
View File

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

170
modules/cudacodec/src/cuda/rgb_to_yv12.cu Normal file
View File

@@ -0,0 +1,170 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "opencv2/core/cuda/common.hpp"
#include "opencv2/core/cuda/vec_traits.hpp"
namespace cv { namespace cuda { namespace device
{
__device__ __forceinline__ void rgb_to_y(const uchar b, const uchar g, const uchar r, uchar& y)
{
y = static_cast<uchar>(((int)(30 * r) + (int)(59 * g) + (int)(11 * b)) / 100);
}
__device__ __forceinline__ void rgb_to_yuv(const uchar b, const uchar g, const uchar r, uchar& y, uchar& u, uchar& v)
{
rgb_to_y(b, g, r, y);
u = static_cast<uchar>(((int)(-17 * r) - (int)(33 * g) + (int)(50 * b) + 12800) / 100);
v = static_cast<uchar>(((int)(50 * r) - (int)(42 * g) - (int)(8 * b) + 12800) / 100);
}
__global__ void Gray_to_YV12(const PtrStepSzb src, PtrStepb dst)
{
const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
if (x + 1 >= src.cols || y + 1 >= src.rows)
return;
// get pointers to the data
const size_t planeSize = src.rows * dst.step;
PtrStepb y_plane(dst.data, dst.step);
PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
uchar pix;
uchar y_val, u_val, v_val;
pix = src(y, x);
rgb_to_y(pix, pix, pix, y_val);
y_plane(y, x) = y_val;
pix = src(y, x + 1);
rgb_to_y(pix, pix, pix, y_val);
y_plane(y, x + 1) = y_val;
pix = src(y + 1, x);
rgb_to_y(pix, pix, pix, y_val);
y_plane(y + 1, x) = y_val;
pix = src(y + 1, x + 1);
rgb_to_yuv(pix, pix, pix, y_val, u_val, v_val);
y_plane(y + 1, x + 1) = y_val;
u_plane(y / 2, x / 2) = u_val;
v_plane(y / 2, x / 2) = v_val;
}
template <typename T>
__global__ void RGB_to_YV12(const PtrStepSz<T> src, PtrStepb dst)
{
const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
if (x + 1 >= src.cols || y + 1 >= src.rows)
return;
// get pointers to the data
const size_t planeSize = src.rows * dst.step;
PtrStepb y_plane(dst.data, dst.step);
PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
T pix;
uchar y_val, u_val, v_val;
pix = src(y, x);
rgb_to_y(pix.z, pix.y, pix.x, y_val);
y_plane(y, x) = y_val;
pix = src(y, x + 1);
rgb_to_y(pix.z, pix.y, pix.x, y_val);
y_plane(y, x + 1) = y_val;
pix = src(y + 1, x);
rgb_to_y(pix.z, pix.y, pix.x, y_val);
y_plane(y + 1, x) = y_val;
pix = src(y + 1, x + 1);
rgb_to_yuv(pix.z, pix.y, pix.x, y_val, u_val, v_val);
y_plane(y + 1, x + 1) = y_val;
u_plane(y / 2, x / 2) = u_val;
v_plane(y / 2, x / 2) = v_val;
}
void Gray_to_YV12_caller(const PtrStepSzb src, PtrStepb dst, cudaStream_t stream)
{
dim3 block(32, 8);
dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
Gray_to_YV12<<<grid, block, 0, stream>>>(src, dst);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
template <int cn>
void RGB_to_YV12_caller(const PtrStepSzb src, PtrStepb dst, cudaStream_t stream)
{
typedef typename TypeVec<uchar, cn>::vec_type src_t;
dim3 block(32, 8);
dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
RGB_to_YV12<<<grid, block, 0, stream>>>(static_cast< PtrStepSz<src_t> >(src), dst);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
void RGB_to_YV12(const PtrStepSzb src, int cn, PtrStepSzb dst, cudaStream_t stream)
{
typedef void (*func_t)(const PtrStepSzb src, PtrStepb dst, cudaStream_t stream);
static const func_t funcs[] =
{
0, Gray_to_YV12_caller, 0, RGB_to_YV12_caller<3>, RGB_to_YV12_caller<4>
};
funcs[cn](src, dst, stream);
}
}}}

114
modules/cudacodec/src/cuvid_video_source.cpp Normal file
View File

@@ -0,0 +1,114 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_NVCUVID
using namespace cv;
using namespace cv::cudacodec;
using namespace cv::cudacodec::detail;
cv::cudacodec::detail::CuvidVideoSource::CuvidVideoSource(const String& fname)
{
CUVIDSOURCEPARAMS params;
std::memset(&params, 0, sizeof(CUVIDSOURCEPARAMS));
// Fill parameter struct
params.pUserData = this; // will be passed to data handlers
params.pfnVideoDataHandler = HandleVideoData; // our local video-handler callback
params.pfnAudioDataHandler = 0;
// now create the actual source
CUresult cuRes = cuvidCreateVideoSource(&videoSource_, fname.c_str(), &params);
if (cuRes == CUDA_ERROR_INVALID_SOURCE)
throw std::runtime_error("");
cuSafeCall( cuRes );
CUVIDEOFORMAT vidfmt;
cuSafeCall( cuvidGetSourceVideoFormat(videoSource_, &vidfmt, 0) );
format_.codec = static_cast<Codec>(vidfmt.codec);
format_.chromaFormat = static_cast<ChromaFormat>(vidfmt.chroma_format);
format_.width = vidfmt.coded_width;
format_.height = vidfmt.coded_height;
}
cv::cudacodec::detail::CuvidVideoSource::~CuvidVideoSource()
{
cuvidDestroyVideoSource(videoSource_);
}
FormatInfo cv::cudacodec::detail::CuvidVideoSource::format() const
{
return format_;
}
void cv::cudacodec::detail::CuvidVideoSource::start()
{
cuSafeCall( cuvidSetVideoSourceState(videoSource_, cudaVideoState_Started) );
}
void cv::cudacodec::detail::CuvidVideoSource::stop()
{
cuSafeCall( cuvidSetVideoSourceState(videoSource_, cudaVideoState_Stopped) );
}
bool cv::cudacodec::detail::CuvidVideoSource::isStarted() const
{
return (cuvidGetVideoSourceState(videoSource_) == cudaVideoState_Started);
}
bool cv::cudacodec::detail::CuvidVideoSource::hasError() const
{
return (cuvidGetVideoSourceState(videoSource_) == cudaVideoState_Error);
}
int CUDAAPI cv::cudacodec::detail::CuvidVideoSource::HandleVideoData(void* userData, CUVIDSOURCEDATAPACKET* packet)
{
CuvidVideoSource* thiz = static_cast<CuvidVideoSource*>(userData);
return thiz->parseVideoData(packet->payload, packet->payload_size, (packet->flags & CUVID_PKT_ENDOFSTREAM) != 0);
}
#endif // HAVE_NVCUVID

87
modules/cudacodec/src/cuvid_video_source.hpp Normal file
View File

@@ -0,0 +1,87 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, 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 __CUVID_VIDEO_SOURCE_HPP__
#define __CUVID_VIDEO_SOURCE_HPP__
#include <nvcuvid.h>
#include "opencv2/core/private.cuda.hpp"
#include "opencv2/cudacodec.hpp"
#include "video_source.hpp"
namespace cv { namespace cudacodec { namespace detail
{
class CuvidVideoSource : public VideoSource
{
public:
explicit CuvidVideoSource(const String& fname);
~CuvidVideoSource();
FormatInfo format() const;
void start();
void stop();
bool isStarted() const;
bool hasError() const;
private:
// Callback for handling packages of demuxed video data.
//
// Parameters:
// pUserData - Pointer to user data. We must pass a pointer to a
// VideoSourceData struct here, that contains a valid CUvideoparser
// and FrameQueue.
// pPacket - video-source data packet.
//
// NOTE: called from a different thread that doesn't not have a cuda context
//
static int CUDAAPI HandleVideoData(void* pUserData, CUVIDSOURCEDATAPACKET* pPacket);
CUvideosource videoSource_;
FormatInfo format_;
};
}}}
#endif // __CUVID_VIDEO_SOURCE_HPP__

139
modules/cudacodec/src/ffmpeg_video_source.cpp Normal file
View File

@@ -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.
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_NVCUVID
using namespace cv;
using namespace cv::cudacodec;
using namespace cv::cudacodec::detail;
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"
CVAUX_STR(CV_VERSION_EPOCH) CVAUX_STR(CV_VERSION_MAJOR) CVAUX_STR(CV_VERSION_MINOR)
#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::cudacodec::detail::FFmpegVideoSource::FFmpegVideoSource(const 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(Error::StsUnsupportedFormat, "Unsupported video source");
format_.codec = static_cast<Codec>(codec);
format_.chromaFormat = static_cast<ChromaFormat>(chroma_format);
format_.width = width;
format_.height = height;
}
cv::cudacodec::detail::FFmpegVideoSource::~FFmpegVideoSource()
{
if (stream_)
release_InputMediaStream_FFMPEG_p(stream_);
}
FormatInfo cv::cudacodec::detail::FFmpegVideoSource::format() const
{
return format_;
}
bool cv::cudacodec::detail::FFmpegVideoSource::getNextPacket(unsigned char** data, int* size, bool* bEndOfFile)
{
int endOfFile;
int res = read_InputMediaStream_FFMPEG_p(stream_, data, size, &endOfFile);
*bEndOfFile = (endOfFile != 0);
return res != 0;
}
#endif // HAVE_CUDA

71
modules/cudacodec/src/ffmpeg_video_source.hpp Normal file
View File

@@ -0,0 +1,71 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, 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_HPP__
#define __FFMPEG_VIDEO_SOURCE_HPP__
#include "opencv2/cudacodec.hpp"
struct InputMediaStream_FFMPEG;
namespace cv { namespace cudacodec { namespace detail {
class FFmpegVideoSource : public RawVideoSource
{
public:
FFmpegVideoSource(const String& fname);
~FFmpegVideoSource();
bool getNextPacket(unsigned char** data, int* size, bool* endOfFile);
FormatInfo format() const;
private:
FormatInfo format_;
InputMediaStream_FFMPEG* stream_;
};
}}}
#endif // __FFMPEG_VIDEO_SOURCE_HPP__

118
modules/cudacodec/src/frame_queue.cpp Normal file
View File

@@ -0,0 +1,118 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_NVCUVID
cv::cudacodec::detail::FrameQueue::FrameQueue() :
endOfDecode_(0),
framesInQueue_(0),
readPosition_(0)
{
std::memset(displayQueue_, 0, sizeof(displayQueue_));
std::memset((void*) isFrameInUse_, 0, sizeof(isFrameInUse_));
}
bool cv::cudacodec::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::cudacodec::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;
{
AutoLock autoLock(mtx_);
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::cudacodec::detail::FrameQueue::dequeue(CUVIDPARSERDISPINFO& displayInfo)
{
AutoLock autoLock(mtx_);
if (framesInQueue_ > 0)
{
int entry = readPosition_;
displayInfo = displayQueue_[entry];
readPosition_ = (entry + 1) % MaximumSize;
framesInQueue_--;
return true;
}
return false;
}
#endif // HAVE_NVCUVID

98
modules/cudacodec/src/frame_queue.hpp Normal file
View File

@@ -0,0 +1,98 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, 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_HPP__
#define __FRAME_QUEUE_HPP__
#include "opencv2/core/utility.hpp"
#include "opencv2/core/private.cuda.hpp"
#include <nvcuvid.h>
namespace cv { namespace cudacodec { 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:
bool isInUse(int pictureIndex) const { return isFrameInUse_[pictureIndex] != 0; }
Mutex mtx_;
volatile int isFrameInUse_[MaximumSize];
volatile int endOfDecode_;
int framesInQueue_;
int readPosition_;
CUVIDPARSERDISPINFO displayQueue_[MaximumSize];
};
}}}
#endif // __FRAME_QUEUE_HPP__

43
modules/cudacodec/src/precomp.cpp Normal file
View File

@@ -0,0 +1,43 @@
/*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"

81
modules/cudacodec/src/precomp.hpp Normal file
View File

@@ -0,0 +1,81 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#ifndef __OPENCV_PRECOMP_H__
#define __OPENCV_PRECOMP_H__
#include <cstdlib>
#include <cstring>
#include <deque>
#include <utility>
#include <stdexcept>
#include <iostream>
#include "opencv2/cudacodec.hpp"
#include "opencv2/core/private.cuda.hpp"
#ifdef HAVE_NVCUVID
#include <nvcuvid.h>
#ifdef WIN32
#define NOMINMAX
#include <windows.h>
#include <NVEncoderAPI.h>
#else
#include <pthread.h>
#include <unistd.h>
#endif
#include "thread.hpp"
#include "video_source.hpp"
#include "ffmpeg_video_source.hpp"
#include "cuvid_video_source.hpp"
#include "frame_queue.hpp"
#include "video_decoder.hpp"
#include "video_parser.hpp"
#include "../src/cap_ffmpeg_api.hpp"
#endif
#endif /* __OPENCV_PRECOMP_H__ */

175
modules/cudacodec/src/thread.cpp Normal file
View File

@@ -0,0 +1,175 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_NVCUVID
using namespace cv::cudacodec::detail;
#ifdef WIN32
namespace
{
struct UserData
{
Thread::Func func;
void* param;
};
DWORD WINAPI WinThreadFunction(LPVOID lpParam)
{
UserData* userData = static_cast<UserData*>(lpParam);
userData->func(userData->param);
return 0;
}
}
class cv::cudacodec::detail::Thread::Impl
{
public:
Impl(Thread::Func func, void* userData)
{
userData_.func = func;
userData_.param = userData;
thread_ = CreateThread(
NULL, // default security attributes
0, // use default stack size
WinThreadFunction, // thread function name
&userData_, // argument to thread function
0, // use default creation flags
&threadId_); // returns the thread identifier
}
~Impl()
{
CloseHandle(thread_);
}
void wait()
{
WaitForSingleObject(thread_, INFINITE);
}
private:
UserData userData_;
HANDLE thread_;
DWORD threadId_;
};
#else
namespace
{
struct UserData
{
Thread::Func func;
void* param;
};
void* PThreadFunction(void* lpParam)
{
UserData* userData = static_cast<UserData*>(lpParam);
userData->func(userData->param);
return 0;
}
}
class cv::cudacodec::detail::Thread::Impl
{
public:
Impl(Thread::Func func, void* userData)
{
userData_.func = func;
userData_.param = userData;
pthread_create(&thread_, NULL, PThreadFunction, &userData_);
}
~Impl()
{
pthread_detach(thread_);
}
void wait()
{
pthread_join(thread_, NULL);
}
private:
pthread_t thread_;
UserData userData_;
};
#endif
cv::cudacodec::detail::Thread::Thread(Func func, void* userData) :
impl_(new Impl(func, userData))
{
}
void cv::cudacodec::detail::Thread::wait()
{
impl_->wait();
}
void cv::cudacodec::detail::Thread::sleep(int ms)
{
#ifdef WIN32
::Sleep(ms);
#else
::usleep(ms * 1000);
#endif
}
template <> void cv::Ptr<cv::cudacodec::detail::Thread::Impl>::delete_obj()
{
if (obj) delete obj;
}
#endif // HAVE_NVCUVID

74
modules/cudacodec/src/thread.hpp Normal file
View File

@@ -0,0 +1,74 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, 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_HPP__
#define __THREAD_WRAPPERS_HPP__
#include "opencv2/core.hpp"
namespace cv { namespace cudacodec { namespace detail {
class Thread
{
public:
typedef void (*Func)(void* userData);
explicit Thread(Func func, void* userData = 0);
void wait();
static void sleep(int ms);
class Impl;
private:
cv::Ptr<Impl> impl_;
};
}}}
namespace cv {
template <> void Ptr<cv::cudacodec::detail::Thread::Impl>::delete_obj();
}
#endif // __THREAD_WRAPPERS_HPP__

116
modules/cudacodec/src/video_decoder.cpp Normal file
View File

@@ -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.
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_NVCUVID
void cv::cudacodec::detail::VideoDecoder::create(const FormatInfo& videoFormat)
{
release();
cudaVideoCodec _codec = static_cast<cudaVideoCodec>(videoFormat.codec);
cudaVideoChromaFormat _chromaFormat = static_cast<cudaVideoChromaFormat>(videoFormat.chromaFormat);
cudaVideoCreateFlags videoCreateFlags = (_codec == cudaVideoCodec_JPEG || _codec == cudaVideoCodec_MPEG2) ?
cudaVideoCreate_PreferCUDA :
cudaVideoCreate_PreferCUVID;
// Validate video format. These are the currently supported formats via NVCUVID
CV_Assert(cudaVideoCodec_MPEG1 == _codec ||
cudaVideoCodec_MPEG2 == _codec ||
cudaVideoCodec_MPEG4 == _codec ||
cudaVideoCodec_VC1 == _codec ||
cudaVideoCodec_H264 == _codec ||
cudaVideoCodec_JPEG == _codec ||
cudaVideoCodec_YUV420== _codec ||
cudaVideoCodec_YV12 == _codec ||
cudaVideoCodec_NV12 == _codec ||
cudaVideoCodec_YUYV == _codec ||
cudaVideoCodec_UYVY == _codec );
CV_Assert(cudaVideoChromaFormat_Monochrome == _chromaFormat ||
cudaVideoChromaFormat_420 == _chromaFormat ||
cudaVideoChromaFormat_422 == _chromaFormat ||
cudaVideoChromaFormat_444 == _chromaFormat);
// Fill the decoder-create-info struct from the given video-format struct.
std::memset(&createInfo_, 0, sizeof(CUVIDDECODECREATEINFO));
// Create video decoder
createInfo_.CodecType = _codec;
createInfo_.ulWidth = videoFormat.width;
createInfo_.ulHeight = videoFormat.height;
createInfo_.ulNumDecodeSurfaces = FrameQueue::MaximumSize;
// Limit decode memory to 24MB (16M pixels at 4:2:0 = 24M bytes)
while (createInfo_.ulNumDecodeSurfaces * videoFormat.width * videoFormat.height > 16 * 1024 * 1024)
createInfo_.ulNumDecodeSurfaces--;
createInfo_.ChromaFormat = _chromaFormat;
createInfo_.OutputFormat = cudaVideoSurfaceFormat_NV12;
createInfo_.DeinterlaceMode = cudaVideoDeinterlaceMode_Adaptive;
// No scaling
static const int MAX_FRAME_COUNT = 2;
createInfo_.ulTargetWidth = createInfo_.ulWidth;
createInfo_.ulTargetHeight = createInfo_.ulHeight;
createInfo_.ulNumOutputSurfaces = MAX_FRAME_COUNT; // We won't simultaneously map more than 8 surfaces
createInfo_.ulCreationFlags = videoCreateFlags;
createInfo_.vidLock = lock_;
// create the decoder
cuSafeCall( cuvidCreateDecoder(&decoder_, &createInfo_) );
}
void cv::cudacodec::detail::VideoDecoder::release()
{
if (decoder_)
{
cuvidDestroyDecoder(decoder_);
decoder_ = 0;
}
}
#endif // HAVE_NVCUVID

112
modules/cudacodec/src/video_decoder.hpp Normal file
View File

@@ -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.
// Copyright (C) 2013, OpenCV Foundation, 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_HPP__
#define __VIDEO_DECODER_HPP__
#include <nvcuvid.h>
#include "opencv2/core/private.cuda.hpp"
#include "opencv2/cudacodec.hpp"
namespace cv { namespace cudacodec { namespace detail
{
class VideoDecoder
{
public:
VideoDecoder(const FormatInfo& videoFormat, CUvideoctxlock lock) : lock_(lock), decoder_(0)
{
create(videoFormat);
}
~VideoDecoder()
{
release();
}
void create(const 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;
}
cuda::GpuMat mapFrame(int picIdx, CUVIDPROCPARAMS& videoProcParams)
{
CUdeviceptr ptr;
unsigned int pitch;
cuSafeCall( cuvidMapVideoFrame(decoder_, picIdx, &ptr, &pitch, &videoProcParams) );
return cuda::GpuMat(targetHeight() * 3 / 2, targetWidth(), CV_8UC1, (void*) ptr, pitch);
}
void unmapFrame(cuda::GpuMat& frame)
{
cuSafeCall( cuvidUnmapVideoFrame(decoder_, (CUdeviceptr) frame.data) );
frame.release();
}
private:
CUvideoctxlock lock_;
CUVIDDECODECREATEINFO createInfo_;
CUvideodecoder decoder_;
};
}}}
#endif // __VIDEO_DECODER_HPP__

162
modules/cudacodec/src/video_parser.cpp Normal file
View File

@@ -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.
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_NVCUVID
cv::cudacodec::detail::VideoParser::VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue) :
videoDecoder_(videoDecoder), frameQueue_(frameQueue), unparsedPackets_(0), hasError_(false)
{
CUVIDPARSERPARAMS params;
std::memset(&params, 0, sizeof(CUVIDPARSERPARAMS));
params.CodecType = videoDecoder->codec();
params.ulMaxNumDecodeSurfaces = videoDecoder->maxDecodeSurfaces();
params.ulMaxDisplayDelay = 1; // this flag is needed so the parser will push frames out to the decoder as quickly as it can
params.pUserData = this;
params.pfnSequenceCallback = HandleVideoSequence; // Called before decoding frames and/or whenever there is a format change
params.pfnDecodePicture = HandlePictureDecode; // Called when a picture is ready to be decoded (decode order)
params.pfnDisplayPicture = HandlePictureDisplay; // Called whenever a picture is ready to be displayed (display order)
cuSafeCall( cuvidCreateVideoParser(&parser_, &params) );
}
bool cv::cudacodec::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 = static_cast<unsigned long>(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::cudacodec::detail::VideoParser::HandleVideoSequence(void* userData, CUVIDEOFORMAT* format)
{
VideoParser* thiz = static_cast<VideoParser*>(userData);
thiz->unparsedPackets_ = 0;
if (format->codec != thiz->videoDecoder_->codec() ||
format->coded_width != thiz->videoDecoder_->frameWidth() ||
format->coded_height != thiz->videoDecoder_->frameHeight() ||
format->chroma_format != thiz->videoDecoder_->chromaFormat())
{
FormatInfo newFormat;
newFormat.codec = static_cast<Codec>(format->codec);
newFormat.chromaFormat = static_cast<ChromaFormat>(format->chroma_format);
newFormat.width = format->coded_width;
newFormat.height = format->coded_height;
try
{
thiz->videoDecoder_->create(newFormat);
}
catch (const cv::Exception&)
{
thiz->hasError_ = true;
return false;
}
}
return true;
}
int CUDAAPI cv::cudacodec::detail::VideoParser::HandlePictureDecode(void* userData, CUVIDPICPARAMS* picParams)
{
VideoParser* thiz = static_cast<VideoParser*>(userData);
thiz->unparsedPackets_ = 0;
bool isFrameAvailable = thiz->frameQueue_->waitUntilFrameAvailable(picParams->CurrPicIdx);
if (!isFrameAvailable)
return false;
if (!thiz->videoDecoder_->decodePicture(picParams))
{
thiz->hasError_ = true;
return false;
}
return true;
}
int CUDAAPI cv::cudacodec::detail::VideoParser::HandlePictureDisplay(void* userData, CUVIDPARSERDISPINFO* picParams)
{
VideoParser* thiz = static_cast<VideoParser*>(userData);
thiz->unparsedPackets_ = 0;
thiz->frameQueue_->enqueue(picParams);
return true;
}
#endif // HAVE_NVCUVID

95
modules/cudacodec/src/video_parser.hpp Normal file
View File

@@ -0,0 +1,95 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, 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_HPP__
#define __VIDEO_PARSER_HPP__
#include <nvcuvid.h>
#include "opencv2/core/private.cuda.hpp"
#include "opencv2/cudacodec.hpp"
#include "frame_queue.hpp"
#include "video_decoder.hpp"
namespace cv { namespace cudacodec { 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:
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 // __VIDEO_PARSER_HPP__

236
modules/cudacodec/src/video_reader.cpp Normal file
View File

@@ -0,0 +1,236 @@
/*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"
using namespace cv;
using namespace cv::cuda;
using namespace cv::cudacodec;
#ifndef HAVE_NVCUVID
Ptr<VideoReader> cv::cudacodec::createVideoReader(const String&) { throw_no_cuda(); return Ptr<VideoReader>(); }
Ptr<VideoReader> cv::cudacodec::createVideoReader(const Ptr<RawVideoSource>&) { throw_no_cuda(); return Ptr<VideoReader>(); }
#else // HAVE_NVCUVID
namespace cv { namespace cuda { namespace device
{
void NV12_to_RGB(const PtrStepb decodedFrame, PtrStepSz<uint> interopFrame, cudaStream_t stream = 0);
}}}
namespace
{
class VideoReaderImpl : public VideoReader
{
public:
explicit VideoReaderImpl(const Ptr<detail::VideoSource>& source);
~VideoReaderImpl();
bool nextFrame(OutputArray frame);
FormatInfo format() const;
private:
Ptr<detail::VideoSource> videoSource_;
Ptr<detail::FrameQueue> frameQueue_;
Ptr<detail::VideoDecoder> videoDecoder_;
Ptr<detail::VideoParser> videoParser_;
CUvideoctxlock lock_;
std::deque< std::pair<CUVIDPARSERDISPINFO, CUVIDPROCPARAMS> > frames_;
};
FormatInfo VideoReaderImpl::format() const
{
return videoSource_->format();
}
VideoReaderImpl::VideoReaderImpl(const Ptr<detail::VideoSource>& source) :
videoSource_(source),
lock_(0)
{
// init context
GpuMat temp(1, 1, CV_8UC1);
temp.release();
CUcontext ctx;
cuSafeCall( cuCtxGetCurrent(&ctx) );
cuSafeCall( cuvidCtxLockCreate(&lock_, ctx) );
frameQueue_ = new detail::FrameQueue;
videoDecoder_ = new detail::VideoDecoder(videoSource_->format(), lock_);
videoParser_ = new detail::VideoParser(videoDecoder_, frameQueue_);
videoSource_->setVideoParser(videoParser_);
videoSource_->start();
}
VideoReaderImpl::~VideoReaderImpl()
{
frameQueue_->endDecode();
videoSource_->stop();
}
class VideoCtxAutoLock
{
public:
VideoCtxAutoLock(CUvideoctxlock lock) : m_lock(lock) { cuSafeCall( cuvidCtxLock(m_lock, 0) ); }
~VideoCtxAutoLock() { cuvidCtxUnlock(m_lock, 0); }
private:
CUvideoctxlock m_lock;
};
void cudaPostProcessFrame(const GpuMat& decodedFrame, OutputArray _outFrame, int width, int height)
{
using namespace cv::cuda::device;
// Final Stage: NV12toARGB color space conversion
_outFrame.create(height, width, CV_8UC4);
GpuMat outFrame = _outFrame.getGpuMat();
NV12_to_RGB(decodedFrame, outFrame);
}
bool VideoReaderImpl::nextFrame(OutputArray frame)
{
if (videoSource_->hasError() || videoParser_->hasError())
CV_Error(Error::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(Error::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<CUVIDPARSERDISPINFO, CUVIDPROCPARAMS> frameInfo = frames_.front();
frames_.pop_front();
{
VideoCtxAutoLock autoLock(lock_);
// map decoded video frame to CUDA surface
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;
}
}
Ptr<VideoReader> cv::cudacodec::createVideoReader(const String& filename)
{
CV_Assert( !filename.empty() );
Ptr<detail::VideoSource> videoSource;
try
{
videoSource = new detail::CuvidVideoSource(filename);
}
catch (...)
{
Ptr<RawVideoSource> source(new detail::FFmpegVideoSource(filename));
videoSource = new detail::RawVideoSourceWrapper(source);
}
return new VideoReaderImpl(videoSource);
}
Ptr<VideoReader> cv::cudacodec::createVideoReader(const Ptr<RawVideoSource>& source)
{
Ptr<detail::VideoSource> videoSource(new detail::RawVideoSourceWrapper(source));
return new VideoReaderImpl(videoSource);
}
#endif // HAVE_NVCUVID

121
modules/cudacodec/src/video_source.cpp Normal file
View File

@@ -0,0 +1,121 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_NVCUVID
using namespace cv;
using namespace cv::cudacodec;
using namespace cv::cudacodec::detail;
bool cv::cudacodec::detail::VideoSource::parseVideoData(const unsigned char* data, size_t size, bool endOfStream)
{
return videoParser_->parseVideoData(data, size, endOfStream);
}
cv::cudacodec::detail::RawVideoSourceWrapper::RawVideoSourceWrapper(const Ptr<RawVideoSource>& source) :
source_(source)
{
CV_Assert( !source_.empty() );
}
cv::cudacodec::FormatInfo cv::cudacodec::detail::RawVideoSourceWrapper::format() const
{
return source_->format();
}
void cv::cudacodec::detail::RawVideoSourceWrapper::start()
{
stop_ = false;
hasError_ = false;
thread_ = new Thread(readLoop, this);
}
void cv::cudacodec::detail::RawVideoSourceWrapper::stop()
{
stop_ = true;
thread_->wait();
thread_.release();
}
bool cv::cudacodec::detail::RawVideoSourceWrapper::isStarted() const
{
return !stop_;
}
bool cv::cudacodec::detail::RawVideoSourceWrapper::hasError() const
{
return hasError_;
}
void cv::cudacodec::detail::RawVideoSourceWrapper::readLoop(void* userData)
{
RawVideoSourceWrapper* thiz = static_cast<RawVideoSourceWrapper*>(userData);
for (;;)
{
unsigned char* data;
int size;
bool endOfFile;
if (!thiz->source_->getNextPacket(&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_NVCUVID

99
modules/cudacodec/src/video_source.hpp Normal file
View File

@@ -0,0 +1,99 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, 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 __CUDACODEC_VIDEO_SOURCE_H__
#define __CUDACODEC_VIDEO_SOURCE_H__
#include "opencv2/core/private.cuda.hpp"
#include "opencv2/cudacodec.hpp"
#include "thread.hpp"
namespace cv { namespace cudacodec { namespace detail
{
class VideoParser;
class VideoSource
{
public:
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 setVideoParser(detail::VideoParser* videoParser) { videoParser_ = videoParser; }
protected:
bool parseVideoData(const uchar* data, size_t size, bool endOfStream = false);
private:
detail::VideoParser* videoParser_;
};
class RawVideoSourceWrapper : public VideoSource
{
public:
RawVideoSourceWrapper(const Ptr<RawVideoSource>& source);
FormatInfo format() const;
void start();
void stop();
bool isStarted() const;
bool hasError() const;
private:
Ptr<RawVideoSource> source_;
Ptr<Thread> thread_;
volatile bool stop_;
volatile bool hasError_;
static void readLoop(void* userData);
};
}}}
#endif // __CUDACODEC_VIDEO_SOURCE_H__

919
modules/cudacodec/src/video_writer.cpp Normal file
View File

@@ -0,0 +1,919 @@
/*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.
// Copyright (C) 2013, OpenCV Foundation, 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"
using namespace cv;
using namespace cv::cuda;
using namespace cv::cudacodec;
#if !defined(HAVE_NVCUVID) || !defined(WIN32)
cv::cudacodec::EncoderParams::EncoderParams() { throw_no_cuda(); }
cv::cudacodec::EncoderParams::EncoderParams(const String&) { throw_no_cuda(); }
void cv::cudacodec::EncoderParams::load(const String&) { throw_no_cuda(); }
void cv::cudacodec::EncoderParams::save(const String&) const { throw_no_cuda(); }
Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const String&, Size, double, SurfaceFormat) { throw_no_cuda(); return Ptr<VideoWriter>(); }
Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const String&, Size, double, const EncoderParams&, SurfaceFormat) { throw_no_cuda(); return Ptr<VideoWriter>(); }
Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>&, Size, double, SurfaceFormat) { throw_no_cuda(); return Ptr<VideoWriter>(); }
Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>&, Size, double, const EncoderParams&, SurfaceFormat) { throw_no_cuda(); return Ptr<VideoWriter>(); }
#else // !defined HAVE_CUDA || !defined WIN32
namespace cv { namespace cuda { namespace device
{
void RGB_to_YV12(const PtrStepSzb src, int cn, PtrStepSzb dst, cudaStream_t stream = 0);
}}}
///////////////////////////////////////////////////////////////////////////
// VideoWriterImpl
namespace
{
class NVEncoderWrapper
{
public:
NVEncoderWrapper() : encoder_(0)
{
int err;
err = NVGetHWEncodeCaps();
if (err)
CV_Error(Error::GpuNotSupported, "No CUDA capability present");
// Create the Encoder API Interface
err = NVCreateEncoder(&encoder_);
CV_Assert( err == 0 );
}
~NVEncoderWrapper()
{
if (encoder_)
NVDestroyEncoder(encoder_);
}
operator NVEncoder() const
{
return encoder_;
}
private:
NVEncoder encoder_;
};
enum CodecType
{
MPEG1, // not supported yet
MPEG2, // not supported yet
MPEG4, // not supported yet
H264
};
class VideoWriterImpl : public VideoWriter
{
public:
VideoWriterImpl(const Ptr<EncoderCallBack>& callback, Size frameSize, double fps, SurfaceFormat format, CodecType codec = H264);
VideoWriterImpl(const Ptr<EncoderCallBack>& callback, Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format, CodecType codec = H264);
void write(InputArray frame, bool lastFrame = false);
EncoderParams getEncoderParams() const;
private:
void initEncoder(double fps);
void setEncodeParams(const EncoderParams& params);
void initGpuMemory();
void initCallBacks();
void createHWEncoder();
Ptr<EncoderCallBack> callback_;
Size frameSize_;
CodecType codec_;
SurfaceFormat inputFormat_;
NVVE_SurfaceFormat surfaceFormat_;
NVEncoderWrapper encoder_;
GpuMat videoFrame_;
CUvideoctxlock cuCtxLock_;
// CallBacks
static unsigned char* NVENCAPI HandleAcquireBitStream(int* pBufferSize, void* pUserdata);
static void NVENCAPI HandleReleaseBitStream(int nBytesInBuffer, unsigned char* cb, void* pUserdata);
static void NVENCAPI HandleOnBeginFrame(const NVVE_BeginFrameInfo* pbfi, void* pUserdata);
static void NVENCAPI HandleOnEndFrame(const NVVE_EndFrameInfo* pefi, void* pUserdata);
};
VideoWriterImpl::VideoWriterImpl(const Ptr<EncoderCallBack>& callback, Size frameSize, double fps, SurfaceFormat format, CodecType codec) :
callback_(callback),
frameSize_(frameSize),
codec_(codec),
inputFormat_(format),
cuCtxLock_(0)
{
surfaceFormat_ = (inputFormat_ == SF_BGR ? YV12 : static_cast<NVVE_SurfaceFormat>(inputFormat_));
initEncoder(fps);
initGpuMemory();
initCallBacks();
createHWEncoder();
}
VideoWriterImpl::VideoWriterImpl(const Ptr<EncoderCallBack>& callback, Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format, CodecType codec) :
callback_(callback),
frameSize_(frameSize),
codec_(codec),
inputFormat_(format),
cuCtxLock_(0)
{
surfaceFormat_ = (inputFormat_ == SF_BGR ? YV12 : static_cast<NVVE_SurfaceFormat>(inputFormat_));
initEncoder(fps);
setEncodeParams(params);
initGpuMemory();
initCallBacks();
createHWEncoder();
}
void VideoWriterImpl::initEncoder(double fps)
{
int err;
// Set codec
static const unsigned long codecs_id[] =
{
NV_CODEC_TYPE_MPEG1, NV_CODEC_TYPE_MPEG2, NV_CODEC_TYPE_MPEG4, NV_CODEC_TYPE_H264, NV_CODEC_TYPE_VC1
};
err = NVSetCodec(encoder_, codecs_id[codec_]);
if (err)
CV_Error(Error::StsNotImplemented, "Codec format is not supported");
// Set default params
err = NVSetDefaultParam(encoder_);
CV_Assert( err == 0 );
// Set some common params
int inputSize[] = { frameSize_.width, frameSize_.height };
err = NVSetParamValue(encoder_, NVVE_IN_SIZE, &inputSize);
CV_Assert( err == 0 );
err = NVSetParamValue(encoder_, NVVE_OUT_SIZE, &inputSize);
CV_Assert( err == 0 );
int aspectRatio[] = { frameSize_.width, frameSize_.height, ASPECT_RATIO_DAR };
err = NVSetParamValue(encoder_, NVVE_ASPECT_RATIO, &aspectRatio);
CV_Assert( err == 0 );
// FPS
int frame_rate = static_cast<int>(fps + 0.5);
int frame_rate_base = 1;
while (fabs(static_cast<double>(frame_rate) / frame_rate_base) - fps > 0.001)
{
frame_rate_base *= 10;
frame_rate = static_cast<int>(fps*frame_rate_base + 0.5);
}
int FrameRate[] = { frame_rate, frame_rate_base };
err = NVSetParamValue(encoder_, NVVE_FRAME_RATE, &FrameRate);
CV_Assert( err == 0 );
// Select device for encoding
int gpuID = getDevice();
err = NVSetParamValue(encoder_, NVVE_FORCE_GPU_SELECTION, &gpuID);
CV_Assert( err == 0 );
}
void VideoWriterImpl::setEncodeParams(const EncoderParams& params)
{
int err;
int P_Interval = params.P_Interval;
err = NVSetParamValue(encoder_, NVVE_P_INTERVAL, &P_Interval);
CV_Assert( err == 0 );
int IDR_Period = params.IDR_Period;
err = NVSetParamValue(encoder_, NVVE_IDR_PERIOD, &IDR_Period);
CV_Assert( err == 0 );
int DynamicGOP = params.DynamicGOP;
err = NVSetParamValue(encoder_, NVVE_DYNAMIC_GOP, &DynamicGOP);
CV_Assert( err == 0 );
NVVE_RateCtrlType RCType = static_cast<NVVE_RateCtrlType>(params.RCType);
err = NVSetParamValue(encoder_, NVVE_RC_TYPE, &RCType);
CV_Assert( err == 0 );
int AvgBitrate = params.AvgBitrate;
err = NVSetParamValue(encoder_, NVVE_AVG_BITRATE, &AvgBitrate);
CV_Assert( err == 0 );
int PeakBitrate = params.PeakBitrate;
err = NVSetParamValue(encoder_, NVVE_PEAK_BITRATE, &PeakBitrate);
CV_Assert( err == 0 );
int QP_Level_Intra = params.QP_Level_Intra;
err = NVSetParamValue(encoder_, NVVE_QP_LEVEL_INTRA, &QP_Level_Intra);
CV_Assert( err == 0 );
int QP_Level_InterP = params.QP_Level_InterP;
err = NVSetParamValue(encoder_, NVVE_QP_LEVEL_INTER_P, &QP_Level_InterP);
CV_Assert( err == 0 );
int QP_Level_InterB = params.QP_Level_InterB;
err = NVSetParamValue(encoder_, NVVE_QP_LEVEL_INTER_B, &QP_Level_InterB);
CV_Assert( err == 0 );
int DeblockMode = params.DeblockMode;
err = NVSetParamValue(encoder_, NVVE_DEBLOCK_MODE, &DeblockMode);
CV_Assert( err == 0 );
int ProfileLevel = params.ProfileLevel;
err = NVSetParamValue(encoder_, NVVE_PROFILE_LEVEL, &ProfileLevel);
CV_Assert( err == 0 );
int ForceIntra = params.ForceIntra;
err = NVSetParamValue(encoder_, NVVE_FORCE_INTRA, &ForceIntra);
CV_Assert( err == 0 );
int ForceIDR = params.ForceIDR;
err = NVSetParamValue(encoder_, NVVE_FORCE_IDR, &ForceIDR);
CV_Assert( err == 0 );
int ClearStat = params.ClearStat;
err = NVSetParamValue(encoder_, NVVE_CLEAR_STAT, &ClearStat);
CV_Assert( err == 0 );
NVVE_DI_MODE DIMode = static_cast<NVVE_DI_MODE>(params.DIMode);
err = NVSetParamValue(encoder_, NVVE_SET_DEINTERLACE, &DIMode);
CV_Assert( err == 0 );
if (params.Presets != -1)
{
NVVE_PRESETS_TARGET Presets = static_cast<NVVE_PRESETS_TARGET>(params.Presets);
err = NVSetParamValue(encoder_, NVVE_PRESETS, &Presets);
CV_Assert( err == 0 );
}
int DisableCabac = params.DisableCabac;
err = NVSetParamValue(encoder_, NVVE_DISABLE_CABAC, &DisableCabac);
CV_Assert( err == 0 );
int NaluFramingType = params.NaluFramingType;
err = NVSetParamValue(encoder_, NVVE_CONFIGURE_NALU_FRAMING_TYPE, &NaluFramingType);
CV_Assert( err == 0 );
int DisableSPSPPS = params.DisableSPSPPS;
err = NVSetParamValue(encoder_, NVVE_DISABLE_SPS_PPS, &DisableSPSPPS);
CV_Assert( err == 0 );
}
EncoderParams VideoWriterImpl::getEncoderParams() const
{
int err;
EncoderParams params;
int P_Interval;
err = NVGetParamValue(encoder_, NVVE_P_INTERVAL, &P_Interval);
CV_Assert( err == 0 );
params.P_Interval = P_Interval;
int IDR_Period;
err = NVGetParamValue(encoder_, NVVE_IDR_PERIOD, &IDR_Period);
CV_Assert( err == 0 );
params.IDR_Period = IDR_Period;
int DynamicGOP;
err = NVGetParamValue(encoder_, NVVE_DYNAMIC_GOP, &DynamicGOP);
CV_Assert( err == 0 );
params.DynamicGOP = DynamicGOP;
NVVE_RateCtrlType RCType;
err = NVGetParamValue(encoder_, NVVE_RC_TYPE, &RCType);
CV_Assert( err == 0 );
params.RCType = RCType;
int AvgBitrate;
err = NVGetParamValue(encoder_, NVVE_AVG_BITRATE, &AvgBitrate);
CV_Assert( err == 0 );
params.AvgBitrate = AvgBitrate;
int PeakBitrate;
err = NVGetParamValue(encoder_, NVVE_PEAK_BITRATE, &PeakBitrate);
CV_Assert( err == 0 );
params.PeakBitrate = PeakBitrate;
int QP_Level_Intra;
err = NVGetParamValue(encoder_, NVVE_QP_LEVEL_INTRA, &QP_Level_Intra);
CV_Assert( err == 0 );
params.QP_Level_Intra = QP_Level_Intra;
int QP_Level_InterP;
err = NVGetParamValue(encoder_, NVVE_QP_LEVEL_INTER_P, &QP_Level_InterP);
CV_Assert( err == 0 );
params.QP_Level_InterP = QP_Level_InterP;
int QP_Level_InterB;
err = NVGetParamValue(encoder_, NVVE_QP_LEVEL_INTER_B, &QP_Level_InterB);
CV_Assert( err == 0 );
params.QP_Level_InterB = QP_Level_InterB;
int DeblockMode;
err = NVGetParamValue(encoder_, NVVE_DEBLOCK_MODE, &DeblockMode);
CV_Assert( err == 0 );
params.DeblockMode = DeblockMode;
int ProfileLevel;
err = NVGetParamValue(encoder_, NVVE_PROFILE_LEVEL, &ProfileLevel);
CV_Assert( err == 0 );
params.ProfileLevel = ProfileLevel;
int ForceIntra;
err = NVGetParamValue(encoder_, NVVE_FORCE_INTRA, &ForceIntra);
CV_Assert( err == 0 );
params.ForceIntra = ForceIntra;
int ForceIDR;
err = NVGetParamValue(encoder_, NVVE_FORCE_IDR, &ForceIDR);
CV_Assert( err == 0 );
params.ForceIDR = ForceIDR;
int ClearStat;
err = NVGetParamValue(encoder_, NVVE_CLEAR_STAT, &ClearStat);
CV_Assert( err == 0 );
params.ClearStat = ClearStat;
NVVE_DI_MODE DIMode;
err = NVGetParamValue(encoder_, NVVE_SET_DEINTERLACE, &DIMode);
CV_Assert( err == 0 );
params.DIMode = DIMode;
params.Presets = -1;
int DisableCabac;
err = NVGetParamValue(encoder_, NVVE_DISABLE_CABAC, &DisableCabac);
CV_Assert( err == 0 );
params.DisableCabac = DisableCabac;
int NaluFramingType;
err = NVGetParamValue(encoder_, NVVE_CONFIGURE_NALU_FRAMING_TYPE, &NaluFramingType);
CV_Assert( err == 0 );
params.NaluFramingType = NaluFramingType;
int DisableSPSPPS;
err = NVGetParamValue(encoder_, NVVE_DISABLE_SPS_PPS, &DisableSPSPPS);
CV_Assert( err == 0 );
params.DisableSPSPPS = DisableSPSPPS;
return params;
}
void VideoWriterImpl::initGpuMemory()
{
int err;
// initialize context
GpuMat temp(1, 1, CV_8U);
temp.release();
static const int bpp[] =
{
16, // UYVY, 4:2:2
16, // YUY2, 4:2:2
12, // YV12, 4:2:0
12, // NV12, 4:2:0
12, // IYUV, 4:2:0
};
CUcontext cuContext;
cuSafeCall( cuCtxGetCurrent(&cuContext) );
// Allocate the CUDA memory Pitched Surface
if (surfaceFormat_ == UYVY || surfaceFormat_ == YUY2)
videoFrame_.create(frameSize_.height, (frameSize_.width * bpp[surfaceFormat_]) / 8, CV_8UC1);
else
videoFrame_.create((frameSize_.height * bpp[surfaceFormat_]) / 8, frameSize_.width, CV_8UC1);
// Create the Video Context Lock (used for synchronization)
cuSafeCall( cuvidCtxLockCreate(&cuCtxLock_, cuContext) );
// If we are using GPU Device Memory with NVCUVENC, it is necessary to create a
// CUDA Context with a Context Lock cuvidCtxLock. The Context Lock needs to be passed to NVCUVENC
int iUseDeviceMem = 1;
err = NVSetParamValue(encoder_, NVVE_DEVICE_MEMORY_INPUT, &iUseDeviceMem);
CV_Assert( err == 0 );
err = NVSetParamValue(encoder_, NVVE_DEVICE_CTX_LOCK, &cuCtxLock_);
CV_Assert( err == 0 );
}
void VideoWriterImpl::initCallBacks()
{
NVVE_CallbackParams cb;
memset(&cb, 0, sizeof(NVVE_CallbackParams));
cb.pfnacquirebitstream = HandleAcquireBitStream;
cb.pfnonbeginframe = HandleOnBeginFrame;
cb.pfnonendframe = HandleOnEndFrame;
cb.pfnreleasebitstream = HandleReleaseBitStream;
NVRegisterCB(encoder_, cb, this);
}
void VideoWriterImpl::createHWEncoder()
{
int err;
// Create the NVIDIA HW resources for Encoding on NVIDIA hardware
err = NVCreateHWEncoder(encoder_);
CV_Assert( err == 0 );
}
// UYVY/YUY2 are both 4:2:2 formats (16bpc)
// Luma, U, V are interleaved, chroma is subsampled (w/2,h)
void copyUYVYorYUY2Frame(Size frameSize, const GpuMat& src, GpuMat& dst)
{
// Source is YUVY/YUY2 4:2:2, the YUV data in a packed and interleaved
// YUV Copy setup
CUDA_MEMCPY2D stCopyYUV422;
memset(&stCopyYUV422, 0, sizeof(CUDA_MEMCPY2D));
stCopyYUV422.srcXInBytes = 0;
stCopyYUV422.srcY = 0;
stCopyYUV422.srcMemoryType = CU_MEMORYTYPE_DEVICE;
stCopyYUV422.srcHost = 0;
stCopyYUV422.srcDevice = (CUdeviceptr) src.data;
stCopyYUV422.srcArray = 0;
stCopyYUV422.srcPitch = src.step;
stCopyYUV422.dstXInBytes = 0;
stCopyYUV422.dstY = 0;
stCopyYUV422.dstMemoryType = CU_MEMORYTYPE_DEVICE;
stCopyYUV422.dstHost = 0;
stCopyYUV422.dstDevice = (CUdeviceptr) dst.data;
stCopyYUV422.dstArray = 0;
stCopyYUV422.dstPitch = dst.step;
stCopyYUV422.WidthInBytes = frameSize.width * 2;
stCopyYUV422.Height = frameSize.height;
// DMA Luma/Chroma
cuSafeCall( cuMemcpy2D(&stCopyYUV422) );
}
// YV12/IYUV are both 4:2:0 planar formats (12bpc)
// Luma, U, V chroma planar (12bpc), chroma is subsampled (w/2,h/2)
void copyYV12orIYUVFrame(Size frameSize, const GpuMat& src, GpuMat& dst)
{
// Source is YV12/IYUV, this native format is converted to NV12 format by the video encoder
// (1) luma copy setup
CUDA_MEMCPY2D stCopyLuma;
memset(&stCopyLuma, 0, sizeof(CUDA_MEMCPY2D));
stCopyLuma.srcXInBytes = 0;
stCopyLuma.srcY = 0;
stCopyLuma.srcMemoryType = CU_MEMORYTYPE_DEVICE;
stCopyLuma.srcHost = 0;
stCopyLuma.srcDevice = (CUdeviceptr) src.data;
stCopyLuma.srcArray = 0;
stCopyLuma.srcPitch = src.step;
stCopyLuma.dstXInBytes = 0;
stCopyLuma.dstY = 0;
stCopyLuma.dstMemoryType = CU_MEMORYTYPE_DEVICE;
stCopyLuma.dstHost = 0;
stCopyLuma.dstDevice = (CUdeviceptr) dst.data;
stCopyLuma.dstArray = 0;
stCopyLuma.dstPitch = dst.step;
stCopyLuma.WidthInBytes = frameSize.width;
stCopyLuma.Height = frameSize.height;
// (2) chroma copy setup, U/V can be done together
CUDA_MEMCPY2D stCopyChroma;
memset(&stCopyChroma, 0, sizeof(CUDA_MEMCPY2D));
stCopyChroma.srcXInBytes = 0;
stCopyChroma.srcY = frameSize.height << 1; // U/V chroma offset
stCopyChroma.srcMemoryType = CU_MEMORYTYPE_DEVICE;
stCopyChroma.srcHost = 0;
stCopyChroma.srcDevice = (CUdeviceptr) src.data;
stCopyChroma.srcArray = 0;
stCopyChroma.srcPitch = src.step >> 1; // chroma is subsampled by 2 (but it has U/V are next to each other)
stCopyChroma.dstXInBytes = 0;
stCopyChroma.dstY = frameSize.height << 1; // chroma offset (srcY*srcPitch now points to the chroma planes)
stCopyChroma.dstMemoryType = CU_MEMORYTYPE_DEVICE;
stCopyChroma.dstHost = 0;
stCopyChroma.dstDevice = (CUdeviceptr) dst.data;
stCopyChroma.dstArray = 0;
stCopyChroma.dstPitch = dst.step >> 1;
stCopyChroma.WidthInBytes = frameSize.width >> 1;
stCopyChroma.Height = frameSize.height; // U/V are sent together
// DMA Luma
cuSafeCall( cuMemcpy2D(&stCopyLuma) );
// DMA Chroma channels (UV side by side)
cuSafeCall( cuMemcpy2D(&stCopyChroma) );
}
// NV12 is 4:2:0 format (12bpc)
// Luma followed by U/V chroma interleaved (12bpc), chroma is subsampled (w/2,h/2)
void copyNV12Frame(Size frameSize, const GpuMat& src, GpuMat& dst)
{
// Source is NV12 in pitch linear memory
// Because we are assume input is NV12 (if we take input in the native format), the encoder handles NV12 as a native format in pitch linear memory
// Luma/Chroma can be done in a single transfer
CUDA_MEMCPY2D stCopyNV12;
memset(&stCopyNV12, 0, sizeof(CUDA_MEMCPY2D));
stCopyNV12.srcXInBytes = 0;
stCopyNV12.srcY = 0;
stCopyNV12.srcMemoryType = CU_MEMORYTYPE_DEVICE;
stCopyNV12.srcHost = 0;
stCopyNV12.srcDevice = (CUdeviceptr) src.data;
stCopyNV12.srcArray = 0;
stCopyNV12.srcPitch = src.step;
stCopyNV12.dstXInBytes = 0;
stCopyNV12.dstY = 0;
stCopyNV12.dstMemoryType = CU_MEMORYTYPE_DEVICE;
stCopyNV12.dstHost = 0;
stCopyNV12.dstDevice = (CUdeviceptr) dst.data;
stCopyNV12.dstArray = 0;
stCopyNV12.dstPitch = dst.step;
stCopyNV12.WidthInBytes = frameSize.width;
stCopyNV12.Height = (frameSize.height * 3) >> 1;
// DMA Luma/Chroma
cuSafeCall( cuMemcpy2D(&stCopyNV12) );
}
void VideoWriterImpl::write(InputArray _frame, bool lastFrame)
{
GpuMat frame = _frame.getGpuMat();
if (inputFormat_ == SF_BGR)
{
CV_Assert( frame.size() == frameSize_ );
CV_Assert( frame.type() == CV_8UC1 || frame.type() == CV_8UC3 || frame.type() == CV_8UC4 );
}
else
{
CV_Assert( frame.size() == videoFrame_.size() );
CV_Assert( frame.type() == videoFrame_.type() );
}
NVVE_EncodeFrameParams efparams;
efparams.Width = frameSize_.width;
efparams.Height = frameSize_.height;
efparams.Pitch = static_cast<int>(videoFrame_.step);
efparams.SurfFmt = surfaceFormat_;
efparams.PictureStruc = FRAME_PICTURE;
efparams.topfieldfirst = 0;
efparams.repeatFirstField = 0;
efparams.progressiveFrame = (surfaceFormat_ == NV12) ? 1 : 0;
efparams.bLast = lastFrame;
efparams.picBuf = 0; // Must be set to NULL in order to support device memory input
// Don't forget we need to lock/unlock between memcopies
cuSafeCall( cuvidCtxLock(cuCtxLock_, 0) );
if (inputFormat_ == SF_BGR)
{
device::RGB_to_YV12(frame, frame.channels(), videoFrame_);
}
else
{
switch (surfaceFormat_)
{
case UYVY: // UYVY (4:2:2)
case YUY2: // YUY2 (4:2:2)
copyUYVYorYUY2Frame(frameSize_, frame, videoFrame_);
break;
case YV12: // YV12 (4:2:0), Y V U
case IYUV: // IYUV (4:2:0), Y U V
copyYV12orIYUVFrame(frameSize_, frame, videoFrame_);
break;
case NV12: // NV12 (4:2:0)
copyNV12Frame(frameSize_, frame, videoFrame_);
break;
}
}
cuSafeCall( cuvidCtxUnlock(cuCtxLock_, 0) );
int err = NVEncodeFrame(encoder_, &efparams, 0, videoFrame_.data);
CV_Assert( err == 0 );
}
unsigned char* NVENCAPI VideoWriterImpl::HandleAcquireBitStream(int* pBufferSize, void* pUserdata)
{
VideoWriterImpl* thiz = static_cast<VideoWriterImpl*>(pUserdata);
return thiz->callback_->acquireBitStream(pBufferSize);
}
void NVENCAPI VideoWriterImpl::HandleReleaseBitStream(int nBytesInBuffer, unsigned char* cb, void* pUserdata)
{
VideoWriterImpl* thiz = static_cast<VideoWriterImpl*>(pUserdata);
thiz->callback_->releaseBitStream(cb, nBytesInBuffer);
}
void NVENCAPI VideoWriterImpl::HandleOnBeginFrame(const NVVE_BeginFrameInfo* pbfi, void* pUserdata)
{
VideoWriterImpl* thiz = static_cast<VideoWriterImpl*>(pUserdata);
thiz->callback_->onBeginFrame(pbfi->nFrameNumber, static_cast<EncoderCallBack::PicType>(pbfi->nPicType));
}
void NVENCAPI VideoWriterImpl::HandleOnEndFrame(const NVVE_EndFrameInfo* pefi, void* pUserdata)
{
VideoWriterImpl* thiz = static_cast<VideoWriterImpl*>(pUserdata);
thiz->callback_->onEndFrame(pefi->nFrameNumber, static_cast<EncoderCallBack::PicType>(pefi->nPicType));
}
///////////////////////////////////////////////////////////////////////////
// FFMPEG
class EncoderCallBackFFMPEG : public EncoderCallBack
{
public:
EncoderCallBackFFMPEG(const String& fileName, Size frameSize, double fps);
~EncoderCallBackFFMPEG();
unsigned char* acquireBitStream(int* bufferSize);
void releaseBitStream(unsigned char* data, int size);
void onBeginFrame(int frameNumber, PicType picType);
void onEndFrame(int frameNumber, PicType picType);
private:
static bool init_MediaStream_FFMPEG();
struct OutputMediaStream_FFMPEG* stream_;
std::vector<uchar> buf_;
bool isKeyFrame_;
static Create_OutputMediaStream_FFMPEG_Plugin create_OutputMediaStream_FFMPEG_p;
static Release_OutputMediaStream_FFMPEG_Plugin release_OutputMediaStream_FFMPEG_p;
static Write_OutputMediaStream_FFMPEG_Plugin write_OutputMediaStream_FFMPEG_p;
};
Create_OutputMediaStream_FFMPEG_Plugin EncoderCallBackFFMPEG::create_OutputMediaStream_FFMPEG_p = 0;
Release_OutputMediaStream_FFMPEG_Plugin EncoderCallBackFFMPEG::release_OutputMediaStream_FFMPEG_p = 0;
Write_OutputMediaStream_FFMPEG_Plugin EncoderCallBackFFMPEG::write_OutputMediaStream_FFMPEG_p = 0;
bool EncoderCallBackFFMPEG::init_MediaStream_FFMPEG()
{
static bool initialized = false;
if (!initialized)
{
#if defined(WIN32) || defined(_WIN32)
const char* module_name = "opencv_ffmpeg"
CVAUX_STR(CV_VERSION_EPOCH) CVAUX_STR(CV_VERSION_MAJOR) CVAUX_STR(CV_VERSION_MINOR)
#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
}
return initialized;
}
EncoderCallBackFFMPEG::EncoderCallBackFFMPEG(const String& fileName, Size frameSize, double fps) :
stream_(0), isKeyFrame_(false)
{
int buf_size = std::max(frameSize.area() * 4, 1024 * 1024);
buf_.resize(buf_size);
CV_Assert( init_MediaStream_FFMPEG() );
stream_ = create_OutputMediaStream_FFMPEG_p(fileName.c_str(), frameSize.width, frameSize.height, fps);
CV_Assert( stream_ != 0 );
}
EncoderCallBackFFMPEG::~EncoderCallBackFFMPEG()
{
release_OutputMediaStream_FFMPEG_p(stream_);
}
unsigned char* EncoderCallBackFFMPEG::acquireBitStream(int* bufferSize)
{
*bufferSize = static_cast<int>(buf_.size());
return &buf_[0];
}
void EncoderCallBackFFMPEG::releaseBitStream(unsigned char* data, int size)
{
write_OutputMediaStream_FFMPEG_p(stream_, data, size, isKeyFrame_);
}
void EncoderCallBackFFMPEG::onBeginFrame(int frameNumber, PicType picType)
{
(void) frameNumber;
isKeyFrame_ = (picType == IFRAME);
}
void EncoderCallBackFFMPEG::onEndFrame(int frameNumber, PicType picType)
{
(void) frameNumber;
(void) picType;
}
}
///////////////////////////////////////////////////////////////////////////
// EncoderParams
cv::cudacodec::EncoderParams::EncoderParams()
{
P_Interval = 3;
IDR_Period = 15;
DynamicGOP = 0;
RCType = 1;
AvgBitrate = 4000000;
PeakBitrate = 10000000;
QP_Level_Intra = 25;
QP_Level_InterP = 28;
QP_Level_InterB = 31;
DeblockMode = 1;
ProfileLevel = 65357;
ForceIntra = 0;
ForceIDR = 0;
ClearStat = 0;
DIMode = 1;
Presets = 2;
DisableCabac = 0;
NaluFramingType = 0;
DisableSPSPPS = 0;
}
cv::cudacodec::EncoderParams::EncoderParams(const String& configFile)
{
load(configFile);
}
void cv::cudacodec::EncoderParams::load(const String& configFile)
{
FileStorage fs(configFile, FileStorage::READ);
CV_Assert( fs.isOpened() );
read(fs["P_Interval" ], P_Interval, 3);
read(fs["IDR_Period" ], IDR_Period, 15);
read(fs["DynamicGOP" ], DynamicGOP, 0);
read(fs["RCType" ], RCType, 1);
read(fs["AvgBitrate" ], AvgBitrate, 4000000);
read(fs["PeakBitrate" ], PeakBitrate, 10000000);
read(fs["QP_Level_Intra" ], QP_Level_Intra, 25);
read(fs["QP_Level_InterP"], QP_Level_InterP, 28);
read(fs["QP_Level_InterB"], QP_Level_InterB, 31);
read(fs["DeblockMode" ], DeblockMode, 1);
read(fs["ProfileLevel" ], ProfileLevel, 65357);
read(fs["ForceIntra" ], ForceIntra, 0);
read(fs["ForceIDR" ], ForceIDR, 0);
read(fs["ClearStat" ], ClearStat, 0);
read(fs["DIMode" ], DIMode, 1);
read(fs["Presets" ], Presets, 2);
read(fs["DisableCabac" ], DisableCabac, 0);
read(fs["NaluFramingType"], NaluFramingType, 0);
read(fs["DisableSPSPPS" ], DisableSPSPPS, 0);
}
void cv::cudacodec::EncoderParams::save(const String& configFile) const
{
FileStorage fs(configFile, FileStorage::WRITE);
CV_Assert( fs.isOpened() );
write(fs, "P_Interval" , P_Interval);
write(fs, "IDR_Period" , IDR_Period);
write(fs, "DynamicGOP" , DynamicGOP);
write(fs, "RCType" , RCType);
write(fs, "AvgBitrate" , AvgBitrate);
write(fs, "PeakBitrate" , PeakBitrate);
write(fs, "QP_Level_Intra" , QP_Level_Intra);
write(fs, "QP_Level_InterP", QP_Level_InterP);
write(fs, "QP_Level_InterB", QP_Level_InterB);
write(fs, "DeblockMode" , DeblockMode);
write(fs, "ProfileLevel" , ProfileLevel);
write(fs, "ForceIntra" , ForceIntra);
write(fs, "ForceIDR" , ForceIDR);
write(fs, "ClearStat" , ClearStat);
write(fs, "DIMode" , DIMode);
write(fs, "Presets" , Presets);
write(fs, "DisableCabac" , DisableCabac);
write(fs, "NaluFramingType", NaluFramingType);
write(fs, "DisableSPSPPS" , DisableSPSPPS);
}
///////////////////////////////////////////////////////////////////////////
// createVideoWriter
Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const String& fileName, Size frameSize, double fps, SurfaceFormat format)
{
Ptr<EncoderCallBack> encoderCallback(new EncoderCallBackFFMPEG(fileName, frameSize, fps));
return createVideoWriter(encoderCallback, frameSize, fps, format);
}
Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const String& fileName, Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format)
{
Ptr<EncoderCallBack> encoderCallback(new EncoderCallBackFFMPEG(fileName, frameSize, fps));
return createVideoWriter(encoderCallback, frameSize, fps, params, format);
}
Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>& encoderCallback, Size frameSize, double fps, SurfaceFormat format)
{
return new VideoWriterImpl(encoderCallback, frameSize, fps, format);
}
Ptr<VideoWriter> cv::cudacodec::createVideoWriter(const Ptr<EncoderCallBack>& encoderCallback, Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format)
{
return new VideoWriterImpl(encoderCallback, frameSize, fps, params, format);
}
#endif // !defined HAVE_CUDA || !defined WIN32