Skip to content

O
opencv
Commit 1b00a3ed authored by Vladislav Vinogradov's avatar Vladislav Vinogradov
Browse files
Options

gpucodec module for video decoding/encoding

parent 02131ffb
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 2114 additions and 1286 deletions
cmake/OpenCVDetectCUDA.cmake
View file @ 1b00a3ed
......@@ -28,6 +28,9 @@ if(CUDA_FOUND)
if(WITH_NVCUVID)
find_cuda_helper_libs(nvcuvid)
if(WIN32)
find_cuda_helper_libs(nvcuvenc)
endif()
set(HAVE_NVCUVID 1)
endif()
......
modules/gpu/CMakeLists.txt
View file @ 1b00a3ed
......@@ -39,19 +39,6 @@ if(HAVE_CUDA)
ocv_cuda_compile(cuda_objs ${lib_cuda} ${ncv_cuda})
set(cuda_link_libs ${CUDA_LIBRARIES} ${CUDA_npp_LIBRARY})
if(WITH_NVCUVID)
set(cuda_link_libs ${cuda_link_libs} ${CUDA_nvcuvid_LIBRARY})
endif()
if(WIN32)
find_cuda_helper_libs(nvcuvenc)
set(cuda_link_libs ${cuda_link_libs} ${CUDA_nvcuvenc_LIBRARY})
endif()
if(WITH_FFMPEG)
set(cuda_link_libs ${cuda_link_libs} ${HIGHGUI_LIBRARIES})
endif()
else()
set(lib_cuda "")
set(cuda_objs "")
......
modules/gpu/doc/video.rst
View file @ 1b00a3ed
This diff is collapsed.
modules/gpu/include/opencv2/gpu.hpp
View file @ 1b00a3ed
......@@ -2156,211 +2156,6 @@ private:
GpuMat buf_;
};
////////////////////////////////// Video Encoding //////////////////////////////////
// Works only under Windows
// Supports olny H264 video codec and AVI files
class CV_EXPORTS VideoWriter_GPU
{
public:
struct EncoderParams;
// Callbacks for video encoder, use it if you want to work with raw video stream
class EncoderCallBack;
enum SurfaceFormat
{
SF_UYVY = 0,
SF_YUY2,
SF_YV12,
SF_NV12,
SF_IYUV,
SF_BGR,
SF_GRAY = SF_BGR
};
VideoWriter_GPU();
VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
~VideoWriter_GPU();
// all methods throws cv::Exception if error occurs
void open(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
void open(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
void open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
void open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
bool isOpened() const;
void close();
void write(const cv::gpu::GpuMat& image, bool lastFrame = false);
struct CV_EXPORTS EncoderParams
{
int P_Interval; // NVVE_P_INTERVAL,
int IDR_Period; // NVVE_IDR_PERIOD,
int DynamicGOP; // NVVE_DYNAMIC_GOP,
int RCType; // NVVE_RC_TYPE,
int AvgBitrate; // NVVE_AVG_BITRATE,
int PeakBitrate; // NVVE_PEAK_BITRATE,
int QP_Level_Intra; // NVVE_QP_LEVEL_INTRA,
int QP_Level_InterP; // NVVE_QP_LEVEL_INTER_P,
int QP_Level_InterB; // NVVE_QP_LEVEL_INTER_B,
int DeblockMode; // NVVE_DEBLOCK_MODE,
int ProfileLevel; // NVVE_PROFILE_LEVEL,
int ForceIntra; // NVVE_FORCE_INTRA,
int ForceIDR; // NVVE_FORCE_IDR,
int ClearStat; // NVVE_CLEAR_STAT,
int DIMode; // NVVE_SET_DEINTERLACE,
int Presets; // NVVE_PRESETS,
int DisableCabac; // NVVE_DISABLE_CABAC,
int NaluFramingType; // NVVE_CONFIGURE_NALU_FRAMING_TYPE
int DisableSPSPPS; // NVVE_DISABLE_SPS_PPS
EncoderParams();
explicit EncoderParams(const String& configFile);
void load(const String& configFile);
void save(const String& configFile) const;
};
EncoderParams getParams() const;
class CV_EXPORTS EncoderCallBack
{
public:
enum PicType
{
IFRAME = 1,
PFRAME = 2,
BFRAME = 3
};
virtual ~EncoderCallBack() {}
// callback function to signal the start of bitstream that is to be encoded
// must return pointer to buffer
virtual uchar* acquireBitStream(int* bufferSize) = 0;
// callback function to signal that the encoded bitstream is ready to be written to file
virtual void releaseBitStream(unsigned char* data, int size) = 0;
// callback function to signal that the encoding operation on the frame has started
virtual void onBeginFrame(int frameNumber, PicType picType) = 0;
// callback function signals that the encoding operation on the frame has finished
virtual void onEndFrame(int frameNumber, PicType picType) = 0;
};
private:
VideoWriter_GPU(const VideoWriter_GPU&);
VideoWriter_GPU& operator=(const VideoWriter_GPU&);
class Impl;
std::auto_ptr<Impl> impl_;
};
////////////////////////////////// Video Decoding //////////////////////////////////////////
namespace detail
{
class FrameQueue;
class VideoParser;
}
class CV_EXPORTS VideoReader_GPU
{
public:
enum Codec
{
MPEG1 = 0,
MPEG2,
MPEG4,
VC1,
H264,
JPEG,
H264_SVC,
H264_MVC,
Uncompressed_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V')), // Y,U,V (4:2:0)
Uncompressed_YV12 = (('Y'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,V,U (4:2:0)
Uncompressed_NV12 = (('N'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,UV (4:2:0)
Uncompressed_YUYV = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V')), // YUYV/YUY2 (4:2:2)
Uncompressed_UYVY = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y')), // UYVY (4:2:2)
};
enum ChromaFormat
{
Monochrome=0,
YUV420,
YUV422,
YUV444,
};
struct FormatInfo
{
Codec codec;
ChromaFormat chromaFormat;
int width;
int height;
};
class VideoSource;
VideoReader_GPU();
explicit VideoReader_GPU(const String& filename);
explicit VideoReader_GPU(const cv::Ptr<VideoSource>& source);
~VideoReader_GPU();
void open(const String& filename);
void open(const cv::Ptr<VideoSource>& source);
bool isOpened() const;
void close();
bool read(GpuMat& image);
FormatInfo format() const;
void dumpFormat(std::ostream& st);
class CV_EXPORTS VideoSource
{
public:
VideoSource() : frameQueue_(0), videoParser_(0) {}
virtual ~VideoSource() {}
virtual FormatInfo format() const = 0;
virtual void start() = 0;
virtual void stop() = 0;
virtual bool isStarted() const = 0;
virtual bool hasError() const = 0;
void setFrameQueue(detail::FrameQueue* frameQueue) { frameQueue_ = frameQueue; }
void setVideoParser(detail::VideoParser* videoParser) { videoParser_ = videoParser; }
protected:
bool parseVideoData(const uchar* data, size_t size, bool endOfStream = false);
private:
VideoSource(const VideoSource&);
VideoSource& operator =(const VideoSource&);
detail::FrameQueue* frameQueue_;
detail::VideoParser* videoParser_;
};
private:
VideoReader_GPU(const VideoReader_GPU&);
VideoReader_GPU& operator =(const VideoReader_GPU&);
class Impl;
std::auto_ptr<Impl> impl_;
};
//! removes points (CV_32FC2, single row matrix) with zero mask value
CV_EXPORTS void compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask);
......
modules/gpu/perf/perf_video.cpp
View file @ 1b00a3ed
......@@ -1005,103 +1005,3 @@ PERF_TEST_P(Video_Cn_MaxFeatures, Video_GMG,
}
#endif
//////////////////////////////////////////////////////
// VideoReader
#if defined(HAVE_NVCUVID) && BUILD_WITH_VIDEO_INPUT_SUPPORT
PERF_TEST_P(Video, DISABLED_Video_VideoReader, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
{
declare.time(20);
const string inputFile = perf::TestBase::getDataPath(GetParam());
if (PERF_RUN_GPU())
{
cv::gpu::VideoReader_GPU d_reader(inputFile);
ASSERT_TRUE( d_reader.isOpened() );
cv::gpu::GpuMat frame;
TEST_CYCLE_N(10) d_reader.read(frame);
GPU_SANITY_CHECK(frame);
}
else
{
cv::VideoCapture reader(inputFile);
ASSERT_TRUE( reader.isOpened() );
cv::Mat frame;
TEST_CYCLE_N(10) reader >> frame;
CPU_SANITY_CHECK(frame);
}
}
#endif
//////////////////////////////////////////////////////
// VideoWriter
#if defined(HAVE_NVCUVID) && defined(WIN32)
PERF_TEST_P(Video, DISABLED_Video_VideoWriter, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
{
declare.time(30);
const string inputFile = perf::TestBase::getDataPath(GetParam());
const string outputFile = cv::tempfile(".avi");
const double FPS = 25.0;
cv::VideoCapture reader(inputFile);
ASSERT_TRUE( reader.isOpened() );
cv::Mat frame;
if (PERF_RUN_GPU())
{
cv::gpu::VideoWriter_GPU d_writer;
cv::gpu::GpuMat d_frame;
for (int i = 0; i < 10; ++i)
{
reader >> frame;
ASSERT_FALSE(frame.empty());
d_frame.upload(frame);
if (!d_writer.isOpened())
d_writer.open(outputFile, frame.size(), FPS);
startTimer(); next();
d_writer.write(d_frame);
stopTimer();
}
}
else
{
cv::VideoWriter writer;
for (int i = 0; i < 10; ++i)
{
reader >> frame;
ASSERT_FALSE(frame.empty());
if (!writer.isOpened())
writer.open(outputFile, CV_FOURCC('X', 'V', 'I', 'D'), FPS, frame.size());
startTimer(); next();
writer.write(frame);
stopTimer();
}
}
SANITY_CHECK(frame);
}
#endif
modules/gpu/src/cuda/NV12ToARGB.cu deleted 100644 → 0
View file @ 02131ffb
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
/*
* NV12ToARGB color space conversion CUDA kernel
*
* This sample uses CUDA to perform a simple NV12 (YUV 4:2:0 planar)
* source and converts to output in ARGB format
*/
#if !defined CUDA_DISABLER
#include "opencv2/core/cuda/common.hpp"
namespace cv { namespace gpu { namespace cudev {
namespace video_decoding
{
__constant__ uint constAlpha = ((uint)0xff << 24);
__constant__ float constHueColorSpaceMat[9];
void loadHueCSC(float hueCSC[9])
{
cudaSafeCall( cudaMemcpyToSymbol(constHueColorSpaceMat, hueCSC, 9 * sizeof(float)) );
}
__device__ void YUV2RGB(const uint* yuvi, float* red, float* green, float* blue)
{
float luma, chromaCb, chromaCr;
// Prepare for hue adjustment
luma = (float)yuvi[0];
chromaCb = (float)((int)yuvi[1] - 512.0f);
chromaCr = (float)((int)yuvi[2] - 512.0f);
// Convert YUV To RGB with hue adjustment
*red = (luma * constHueColorSpaceMat[0]) +
(chromaCb * constHueColorSpaceMat[1]) +
(chromaCr * constHueColorSpaceMat[2]);
*green = (luma * constHueColorSpaceMat[3]) +
(chromaCb * constHueColorSpaceMat[4]) +
(chromaCr * constHueColorSpaceMat[5]);
*blue = (luma * constHueColorSpaceMat[6]) +
(chromaCb * constHueColorSpaceMat[7]) +
(chromaCr * constHueColorSpaceMat[8]);
}
__device__ uint RGBAPACK_10bit(float red, float green, float blue, uint alpha)
{
uint ARGBpixel = 0;
// Clamp final 10 bit results
red = ::fmin(::fmax(red, 0.0f), 1023.f);
green = ::fmin(::fmax(green, 0.0f), 1023.f);
blue = ::fmin(::fmax(blue, 0.0f), 1023.f);
// Convert to 8 bit unsigned integers per color component
ARGBpixel = (((uint)blue >> 2) |
(((uint)green >> 2) << 8) |
(((uint)red >> 2) << 16) |
(uint)alpha);
return ARGBpixel;
}
// CUDA kernel for outputing the final ARGB output from NV12
#define COLOR_COMPONENT_BIT_SIZE 10
#define COLOR_COMPONENT_MASK 0x3FF
__global__ void NV12ToARGB(uchar* srcImage, size_t nSourcePitch,
uint* dstImage, size_t nDestPitch,
uint width, uint height)
{
// Pad borders with duplicate pixels, and we multiply by 2 because we process 2 pixels per thread
const int x = blockIdx.x * (blockDim.x << 1) + (threadIdx.x << 1);
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= width || y >= height)
return;
// Read 2 Luma components at a time, so we don't waste processing since CbCr are decimated this way.
// if we move to texture we could read 4 luminance values
uint yuv101010Pel[2];
yuv101010Pel[0] = (srcImage[y * nSourcePitch + x ]) << 2;
yuv101010Pel[1] = (srcImage[y * nSourcePitch + x + 1]) << 2;
const size_t chromaOffset = nSourcePitch * height;
const int y_chroma = y >> 1;
if (y & 1) // odd scanline ?
{
uint chromaCb = srcImage[chromaOffset + y_chroma * nSourcePitch + x ];
uint chromaCr = srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1];
if (y_chroma < ((height >> 1) - 1)) // interpolate chroma vertically
{
chromaCb = (chromaCb + srcImage[chromaOffset + (y_chroma + 1) * nSourcePitch + x ] + 1) >> 1;
chromaCr = (chromaCr + srcImage[chromaOffset + (y_chroma + 1) * nSourcePitch + x + 1] + 1) >> 1;
}
yuv101010Pel[0] |= (chromaCb << ( COLOR_COMPONENT_BIT_SIZE + 2));
yuv101010Pel[0] |= (chromaCr << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
yuv101010Pel[1] |= (chromaCb << ( COLOR_COMPONENT_BIT_SIZE + 2));
yuv101010Pel[1] |= (chromaCr << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
}
else
{
yuv101010Pel[0] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x ] << ( COLOR_COMPONENT_BIT_SIZE + 2));
yuv101010Pel[0] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1] << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
yuv101010Pel[1] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x ] << ( COLOR_COMPONENT_BIT_SIZE + 2));
yuv101010Pel[1] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1] << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
}
// this steps performs the color conversion
uint yuvi[6];
float red[2], green[2], blue[2];
yuvi[0] = (yuv101010Pel[0] & COLOR_COMPONENT_MASK );
yuvi[1] = ((yuv101010Pel[0] >> COLOR_COMPONENT_BIT_SIZE) & COLOR_COMPONENT_MASK);
yuvi[2] = ((yuv101010Pel[0] >> (COLOR_COMPONENT_BIT_SIZE << 1)) & COLOR_COMPONENT_MASK);
yuvi[3] = (yuv101010Pel[1] & COLOR_COMPONENT_MASK );
yuvi[4] = ((yuv101010Pel[1] >> COLOR_COMPONENT_BIT_SIZE) & COLOR_COMPONENT_MASK);
yuvi[5] = ((yuv101010Pel[1] >> (COLOR_COMPONENT_BIT_SIZE << 1)) & COLOR_COMPONENT_MASK);
// YUV to RGB Transformation conversion
YUV2RGB(&yuvi[0], &red[0], &green[0], &blue[0]);
YUV2RGB(&yuvi[3], &red[1], &green[1], &blue[1]);
// Clamp the results to RGBA
const size_t dstImagePitch = nDestPitch >> 2;
dstImage[y * dstImagePitch + x ] = RGBAPACK_10bit(red[0], green[0], blue[0], constAlpha);
dstImage[y * dstImagePitch + x + 1 ] = RGBAPACK_10bit(red[1], green[1], blue[1], constAlpha);
}
void NV12ToARGB_gpu(const PtrStepb decodedFrame, PtrStepSz<uint> interopFrame, cudaStream_t stream)
{
dim3 block(32, 8);
dim3 grid(divUp(interopFrame.cols, 2 * block.x), divUp(interopFrame.rows, block.y));
NV12ToARGB<<<grid, block, 0, stream>>>(decodedFrame.data, decodedFrame.step, interopFrame.data, interopFrame.step,
interopFrame.cols, interopFrame.rows);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
}
}}}
#endif /* CUDA_DISABLER */
modules/gpu/src/precomp.hpp
View file @ 1b00a3ed
......@@ -80,20 +80,6 @@
#include <cublas.h>
#endif
#ifdef HAVE_NVCUVID
#include <nvcuvid.h>
#ifdef WIN32
#include <windows.h>
#undef small
#undef min
#undef max
#undef abs
#include <NVEncoderAPI.h>
#endif
#endif
#include "internal_shared.hpp"
#include "opencv2/core/stream_accessor.hpp"
......
modules/gpu/src/thread_wrappers.cpp deleted 100644 → 0
View file @ 02131ffb
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "thread_wrappers.h"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#ifdef WIN32
#define NOMINMAX
#include <windows.h>
#else
#include <pthread.h>
#include <unistd.h>
#endif
#ifdef WIN32
class cv::gpu::detail::CriticalSection::Impl
{
public:
Impl()
{
InitializeCriticalSection(&criticalSection_);
}
~Impl()
{
DeleteCriticalSection(&criticalSection_);
}
void enter()
{
EnterCriticalSection(&criticalSection_);
}
void leave()
{
LeaveCriticalSection(&criticalSection_);
}
private:
CRITICAL_SECTION criticalSection_;
};
#else
class cv::gpu::detail::CriticalSection::Impl
{
public:
Impl()
{
pthread_mutexattr_t mutex_attribute;
pthread_mutexattr_init(&mutex_attribute);
pthread_mutexattr_settype(&mutex_attribute, PTHREAD_MUTEX_RECURSIVE);
pthread_mutex_init(&mutex_, 0);
pthread_mutexattr_destroy(&mutex_attribute);
}
~Impl()
{
pthread_mutex_destroy(&mutex_);
}
void enter()
{
pthread_mutex_lock(&mutex_);
}
void leave()
{
pthread_mutex_unlock(&mutex_);
}
private:
pthread_mutex_t mutex_;
};
#endif
cv::gpu::detail::CriticalSection::CriticalSection() :
impl_(new Impl)
{
}
cv::gpu::detail::CriticalSection::~CriticalSection()
{
}
void cv::gpu::detail::CriticalSection::enter()
{
impl_->enter();
}
void cv::gpu::detail::CriticalSection::leave()
{
impl_->leave();
}
#ifdef WIN32
namespace
{
struct UserData
{
void (*func)(void* userData);
void* param;
};
DWORD WINAPI WinThreadFunction(LPVOID lpParam)
{
UserData* userData = static_cast<UserData*>(lpParam);
userData->func(userData->param);
return 0;
}
}
class cv::gpu::detail::Thread::Impl
{
public:
Impl(void (*func)(void* userData), void* userData)
{
userData_.func = func;
userData_.param = userData;
thread_ = CreateThread(
NULL, // default security attributes
0, // use default stack size
WinThreadFunction, // thread function name
&userData_, // argument to thread function
0, // use default creation flags
&threadId_); // returns the thread identifier
}
~Impl()
{
CloseHandle(thread_);
}
void wait()
{
WaitForSingleObject(thread_, INFINITE);
}
private:
UserData userData_;
HANDLE thread_;
DWORD threadId_;
};
#else
namespace
{
struct UserData
{
void (*func)(void* userData);
void* param;
};
void* PThreadFunction(void* lpParam)
{
UserData* userData = static_cast<UserData*>(lpParam);
userData->func(userData->param);
return 0;
}
}
class cv::gpu::detail::Thread::Impl
{
public:
Impl(void (*func)(void* userData), void* userData)
{
userData_.func = func;
userData_.param = userData;
pthread_create(&thread_, NULL, PThreadFunction, &userData_);
}
~Impl()
{
pthread_detach(thread_);
}
void wait()
{
pthread_join(thread_, NULL);
}
private:
pthread_t thread_;
UserData userData_;
};
#endif
cv::gpu::detail::Thread::Thread(void (*func)(void* userData), void* userData) :
impl_(new Impl(func, userData))
{
}
cv::gpu::detail::Thread::~Thread()
{
}
void cv::gpu::detail::Thread::wait()
{
impl_->wait();
}
void cv::gpu::detail::Thread::sleep(int ms)
{
#ifdef WIN32
::Sleep(ms);
#else
::usleep(ms * 1000);
#endif
}
#endif // HAVE_CUDA
modules/gpucodec/CMakeLists.txt 0 → 100644
View file @ 1b00a3ed
if(ANDROID OR IOS)
ocv_module_disable(gpucodec)
endif()
set(the_description "GPU-accelerated Video Encoding/Decoding")
ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations)
ocv_add_module(gpucodec opencv_highgui)
ocv_module_include_directories()
ocv_glob_module_sources()
set(extra_libs ${HIGHGUI_LIBRARIES})
if(HAVE_NVCUVID)
list(APPEND extra_libs ${CUDA_nvcuvid_LIBRARY})
if(WIN32)
list(APPEND extra_libs ${CUDA_nvcuvenc_LIBRARY})
endif()
endif()
ocv_create_module(${extra_libs})
ocv_add_precompiled_headers(${the_module})
ocv_add_accuracy_tests()
ocv_add_perf_tests()
modules/gpucodec/doc/gpucodec.rst 0 → 100644
View file @ 1b00a3ed
*************************************************
gpucodec. GPU-accelerated Video Encoding/Decoding
*************************************************
.. toctree::
:maxdepth: 1
videodec
videoenc
modules/gpucodec/doc/videodec.rst 0 → 100644
View file @ 1b00a3ed
Video Decoding
==============
.. highlight:: cpp
gpu::VideoReader_GPU
--------------------
Video reader class.
.. ocv:class:: gpu::VideoReader_GPU
gpu::VideoReader_GPU::Codec
---------------------------
Video codecs supported by :ocv:class:`gpu::VideoReader_GPU` .
.. ocv:enum:: gpu::VideoReader_GPU::Codec
.. ocv:emember:: MPEG1 = 0
.. ocv:emember:: MPEG2
.. ocv:emember:: MPEG4
.. ocv:emember:: VC1
.. ocv:emember:: H264
.. ocv:emember:: JPEG
.. ocv:emember:: H264_SVC
.. ocv:emember:: H264_MVC
.. ocv:emember:: Uncompressed_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V'))
Y,U,V (4:2:0)
.. ocv:emember:: Uncompressed_YV12 = (('Y'<<24)|('V'<<16)|('1'<<8)|('2'))
Y,V,U (4:2:0)
.. ocv:emember:: Uncompressed_NV12 = (('N'<<24)|('V'<<16)|('1'<<8)|('2'))
Y,UV (4:2:0)
.. ocv:emember:: Uncompressed_YUYV = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V'))
YUYV/YUY2 (4:2:2)
.. ocv:emember:: Uncompressed_UYVY = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y'))
UYVY (4:2:2)
gpu::VideoReader_GPU::ChromaFormat
----------------------------------
Chroma formats supported by :ocv:class:`gpu::VideoReader_GPU` .
.. ocv:enum:: gpu::VideoReader_GPU::ChromaFormat
.. ocv:emember:: Monochrome = 0
.. ocv:emember:: YUV420
.. ocv:emember:: YUV422
.. ocv:emember:: YUV444
gpu::VideoReader_GPU::FormatInfo
--------------------------------
.. ocv:struct:: gpu::VideoReader_GPU::FormatInfo
Struct providing information about video file format. ::
struct FormatInfo
{
Codec codec;
ChromaFormat chromaFormat;
int width;
int height;
};
gpu::VideoReader_GPU::VideoReader_GPU
-------------------------------------
Constructors.
.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU()
.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU(const String& filename)
.. ocv:function:: gpu::VideoReader_GPU::VideoReader_GPU(const cv::Ptr<VideoSource>& source)
:param filename: Name of the input video file.
:param source: Video file parser implemented by user.
The constructors initialize video reader. FFMPEG is used to read videos. User can implement own demultiplexing with :ocv:class:`gpu::VideoReader_GPU::VideoSource` .
gpu::VideoReader_GPU::open
--------------------------
Initializes or reinitializes video reader.
.. ocv:function:: void gpu::VideoReader_GPU::open(const String& filename)
.. ocv:function:: void gpu::VideoReader_GPU::open(const cv::Ptr<VideoSource>& source)
The method opens video reader. Parameters are the same as in the constructor :ocv:func:`gpu::VideoReader_GPU::VideoReader_GPU` . The method throws :ocv:class:`Exception` if error occurs.
gpu::VideoReader_GPU::isOpened
------------------------------
Returns true if video reader has been successfully initialized.
.. ocv:function:: bool gpu::VideoReader_GPU::isOpened() const
gpu::VideoReader_GPU::close
---------------------------
Releases the video reader.
.. ocv:function:: void gpu::VideoReader_GPU::close()
gpu::VideoReader_GPU::read
--------------------------
Grabs, decodes and returns the next video frame.
.. ocv:function:: bool gpu::VideoReader_GPU::read(GpuMat& image)
If no frames has been grabbed (there are no more frames in video file), the methods return ``false`` . The method throws :ocv:class:`Exception` if error occurs.
gpu::VideoReader_GPU::format
----------------------------
Returns information about video file format.
.. ocv:function:: FormatInfo gpu::VideoReader_GPU::format() const
The method throws :ocv:class:`Exception` if video reader wasn't initialized.
gpu::VideoReader_GPU::dumpFormat
--------------------------------
Dump information about video file format to specified stream.
.. ocv:function:: void gpu::VideoReader_GPU::dumpFormat(std::ostream& st)
:param st: Output stream.
The method throws :ocv:class:`Exception` if video reader wasn't initialized.
gpu::VideoReader_GPU::VideoSource
-----------------------------------
.. ocv:class:: gpu::VideoReader_GPU::VideoSource
Interface for video demultiplexing. ::
class VideoSource
{
public:
VideoSource();
virtual ~VideoSource() {}
virtual FormatInfo format() const = 0;
virtual void start() = 0;
virtual void stop() = 0;
virtual bool isStarted() const = 0;
virtual bool hasError() const = 0;
protected:
bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream = false);
};
User can implement own demultiplexing by implementing this interface.
gpu::VideoReader_GPU::VideoSource::format
-----------------------------------------
Returns information about video file format.
.. ocv:function:: virtual FormatInfo gpu::VideoReader_GPU::VideoSource::format() const = 0
gpu::VideoReader_GPU::VideoSource::start
----------------------------------------
Starts processing.
.. ocv:function:: virtual void gpu::VideoReader_GPU::VideoSource::start() = 0
Implementation must create own thread with video processing and call periodic :ocv:func:`gpu::VideoReader_GPU::VideoSource::parseVideoData` .
gpu::VideoReader_GPU::VideoSource::stop
---------------------------------------
Stops processing.
.. ocv:function:: virtual void gpu::VideoReader_GPU::VideoSource::stop() = 0
gpu::VideoReader_GPU::VideoSource::isStarted
--------------------------------------------
Returns ``true`` if processing was successfully started.
.. ocv:function:: virtual bool gpu::VideoReader_GPU::VideoSource::isStarted() const = 0
gpu::VideoReader_GPU::VideoSource::hasError
-------------------------------------------
Returns ``true`` if error occured during processing.
.. ocv:function:: virtual bool gpu::VideoReader_GPU::VideoSource::hasError() const = 0
gpu::VideoReader_GPU::VideoSource::parseVideoData
-------------------------------------------------
Parse next video frame. Implementation must call this method after new frame was grabbed.
.. ocv:function:: bool gpu::VideoReader_GPU::VideoSource::parseVideoData(const uchar* data, size_t size, bool endOfStream = false)
:param data: Pointer to frame data. Can be ``NULL`` if ``endOfStream`` if ``true`` .
:param size: Size in bytes of current frame.
:param endOfStream: Indicates that it is end of stream.
modules/gpucodec/doc/videoenc.rst 0 → 100644
View file @ 1b00a3ed
Video Encoding
==============
.. highlight:: cpp
gpu::VideoWriter_GPU
---------------------
Video writer class.
.. ocv:class:: gpu::VideoWriter_GPU
The class uses H264 video codec.
.. note:: Currently only Windows platform is supported.
gpu::VideoWriter_GPU::VideoWriter_GPU
-------------------------------------
Constructors.
.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU()
.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
.. ocv:function:: gpu::VideoWriter_GPU::VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
:param fileName: Name of the output video file. Only AVI file format is supported.
:param frameSize: Size of the input video frames.
:param fps: Framerate of the created video stream.
:param params: Encoder parameters. See :ocv:struct:`gpu::VideoWriter_GPU::EncoderParams` .
:param format: Surface format of input frames ( ``SF_UYVY`` , ``SF_YUY2`` , ``SF_YV12`` , ``SF_NV12`` , ``SF_IYUV`` , ``SF_BGR`` or ``SF_GRAY``). BGR or gray frames will be converted to YV12 format before encoding, frames with other formats will be used as is.
:param encoderCallback: Callbacks for video encoder. See :ocv:class:`gpu::VideoWriter_GPU::EncoderCallBack` . Use it if you want to work with raw video stream.
The constructors initialize video writer. FFMPEG is used to write videos. User can implement own multiplexing with :ocv:class:`gpu::VideoWriter_GPU::EncoderCallBack` .
gpu::VideoWriter_GPU::open
--------------------------
Initializes or reinitializes video writer.
.. ocv:function:: void gpu::VideoWriter_GPU::open(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
.. ocv:function:: void gpu::VideoWriter_GPU::open(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
.. ocv:function:: void gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR)
.. ocv:function:: void gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR)
The method opens video writer. Parameters are the same as in the constructor :ocv:func:`gpu::VideoWriter_GPU::VideoWriter_GPU` . The method throws :ocv:class:`Exception` if error occurs.
gpu::VideoWriter_GPU::isOpened
------------------------------
Returns true if video writer has been successfully initialized.
.. ocv:function:: bool gpu::VideoWriter_GPU::isOpened() const
gpu::VideoWriter_GPU::close
---------------------------
Releases the video writer.
.. ocv:function:: void gpu::VideoWriter_GPU::close()
gpu::VideoWriter_GPU::write
---------------------------
Writes the next video frame.
.. ocv:function:: void gpu::VideoWriter_GPU::write(const cv::gpu::GpuMat& image, bool lastFrame = false)
:param image: The written frame.
:param lastFrame: Indicates that it is end of stream. The parameter can be ignored.
The method write the specified image to video file. The image must have the same size and the same surface format as has been specified when opening the video writer.
gpu::VideoWriter_GPU::EncoderParams
-----------------------------------
.. ocv:struct:: gpu::VideoWriter_GPU::EncoderParams
Different parameters for CUDA video encoder. ::
struct EncoderParams
{
int P_Interval; // NVVE_P_INTERVAL,
int IDR_Period; // NVVE_IDR_PERIOD,
int DynamicGOP; // NVVE_DYNAMIC_GOP,
int RCType; // NVVE_RC_TYPE,
int AvgBitrate; // NVVE_AVG_BITRATE,
int PeakBitrate; // NVVE_PEAK_BITRATE,
int QP_Level_Intra; // NVVE_QP_LEVEL_INTRA,
int QP_Level_InterP; // NVVE_QP_LEVEL_INTER_P,
int QP_Level_InterB; // NVVE_QP_LEVEL_INTER_B,
int DeblockMode; // NVVE_DEBLOCK_MODE,
int ProfileLevel; // NVVE_PROFILE_LEVEL,
int ForceIntra; // NVVE_FORCE_INTRA,
int ForceIDR; // NVVE_FORCE_IDR,
int ClearStat; // NVVE_CLEAR_STAT,
int DIMode; // NVVE_SET_DEINTERLACE,
int Presets; // NVVE_PRESETS,
int DisableCabac; // NVVE_DISABLE_CABAC,
int NaluFramingType; // NVVE_CONFIGURE_NALU_FRAMING_TYPE
int DisableSPSPPS; // NVVE_DISABLE_SPS_PPS
EncoderParams();
explicit EncoderParams(const String& configFile);
void load(const String& configFile);
void save(const String& configFile) const;
};
gpu::VideoWriter_GPU::EncoderParams::EncoderParams
--------------------------------------------------
Constructors.
.. ocv:function:: gpu::VideoWriter_GPU::EncoderParams::EncoderParams()
.. ocv:function:: gpu::VideoWriter_GPU::EncoderParams::EncoderParams(const String& configFile)
:param configFile: Config file name.
Creates default parameters or reads parameters from config file.
gpu::VideoWriter_GPU::EncoderParams::load
-----------------------------------------
Reads parameters from config file.
.. ocv:function:: void gpu::VideoWriter_GPU::EncoderParams::load(const String& configFile)
:param configFile: Config file name.
gpu::VideoWriter_GPU::EncoderParams::save
-----------------------------------------
Saves parameters to config file.
.. ocv:function:: void gpu::VideoWriter_GPU::EncoderParams::save(const String& configFile) const
:param configFile: Config file name.
gpu::VideoWriter_GPU::EncoderCallBack
-------------------------------------
.. ocv:class:: gpu::VideoWriter_GPU::EncoderCallBack
Callbacks for CUDA video encoder. ::
class EncoderCallBack
{
public:
enum PicType
{
IFRAME = 1,
PFRAME = 2,
BFRAME = 3
};
virtual ~EncoderCallBack() {}
virtual unsigned char* acquireBitStream(int* bufferSize) = 0;
virtual void releaseBitStream(unsigned char* data, int size) = 0;
virtual void onBeginFrame(int frameNumber, PicType picType) = 0;
virtual void onEndFrame(int frameNumber, PicType picType) = 0;
};
gpu::VideoWriter_GPU::EncoderCallBack::acquireBitStream
-------------------------------------------------------
Callback function to signal the start of bitstream that is to be encoded.
.. ocv:function:: virtual uchar* gpu::VideoWriter_GPU::EncoderCallBack::acquireBitStream(int* bufferSize) = 0
Callback must allocate buffer for CUDA encoder and return pointer to it and it's size.
gpu::VideoWriter_GPU::EncoderCallBack::releaseBitStream
-------------------------------------------------------
Callback function to signal that the encoded bitstream is ready to be written to file.
.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::releaseBitStream(unsigned char* data, int size) = 0
gpu::VideoWriter_GPU::EncoderCallBack::onBeginFrame
---------------------------------------------------
Callback function to signal that the encoding operation on the frame has started.
.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::onBeginFrame(int frameNumber, PicType picType) = 0
:param picType: Specify frame type (I-Frame, P-Frame or B-Frame).
gpu::VideoWriter_GPU::EncoderCallBack::onEndFrame
-------------------------------------------------
Callback function signals that the encoding operation on the frame has finished.
.. ocv:function:: virtual void gpu::VideoWriter_GPU::EncoderCallBack::onEndFrame(int frameNumber, PicType picType) = 0
:param picType: Specify frame type (I-Frame, P-Frame or B-Frame).
modules/gpucodec/include/opencv2/gpucodec.hpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#ifndef __OPENCV_GPUCODEC_HPP__
#define __OPENCV_GPUCODEC_HPP__
#ifndef __cplusplus
# error gpucodec.hpp header must be compiled as C++
#endif
#include <iosfwd>
#include "opencv2/core/gpumat.hpp"
namespace cv { namespace gpu {
////////////////////////////////// Video Encoding //////////////////////////////////
// Works only under Windows
// Supports olny H264 video codec and AVI files
class CV_EXPORTS VideoWriter_GPU
{
public:
struct EncoderParams;
// Callbacks for video encoder, use it if you want to work with raw video stream
class EncoderCallBack;
enum SurfaceFormat
{
SF_UYVY = 0,
SF_YUY2,
SF_YV12,
SF_NV12,
SF_IYUV,
SF_BGR,
SF_GRAY = SF_BGR
};
VideoWriter_GPU();
VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
VideoWriter_GPU(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
VideoWriter_GPU(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
~VideoWriter_GPU();
// all methods throws cv::Exception if error occurs
void open(const String& fileName, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
void open(const String& fileName, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
void open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format = SF_BGR);
void open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format = SF_BGR);
bool isOpened() const;
void close();
void write(const cv::gpu::GpuMat& image, bool lastFrame = false);
struct CV_EXPORTS EncoderParams
{
int P_Interval; // NVVE_P_INTERVAL,
int IDR_Period; // NVVE_IDR_PERIOD,
int DynamicGOP; // NVVE_DYNAMIC_GOP,
int RCType; // NVVE_RC_TYPE,
int AvgBitrate; // NVVE_AVG_BITRATE,
int PeakBitrate; // NVVE_PEAK_BITRATE,
int QP_Level_Intra; // NVVE_QP_LEVEL_INTRA,
int QP_Level_InterP; // NVVE_QP_LEVEL_INTER_P,
int QP_Level_InterB; // NVVE_QP_LEVEL_INTER_B,
int DeblockMode; // NVVE_DEBLOCK_MODE,
int ProfileLevel; // NVVE_PROFILE_LEVEL,
int ForceIntra; // NVVE_FORCE_INTRA,
int ForceIDR; // NVVE_FORCE_IDR,
int ClearStat; // NVVE_CLEAR_STAT,
int DIMode; // NVVE_SET_DEINTERLACE,
int Presets; // NVVE_PRESETS,
int DisableCabac; // NVVE_DISABLE_CABAC,
int NaluFramingType; // NVVE_CONFIGURE_NALU_FRAMING_TYPE
int DisableSPSPPS; // NVVE_DISABLE_SPS_PPS
EncoderParams();
explicit EncoderParams(const String& configFile);
void load(const String& configFile);
void save(const String& configFile) const;
};
EncoderParams getParams() const;
class CV_EXPORTS EncoderCallBack
{
public:
enum PicType
{
IFRAME = 1,
PFRAME = 2,
BFRAME = 3
};
virtual ~EncoderCallBack() {}
// callback function to signal the start of bitstream that is to be encoded
// must return pointer to buffer
virtual uchar* acquireBitStream(int* bufferSize) = 0;
// callback function to signal that the encoded bitstream is ready to be written to file
virtual void releaseBitStream(unsigned char* data, int size) = 0;
// callback function to signal that the encoding operation on the frame has started
virtual void onBeginFrame(int frameNumber, PicType picType) = 0;
// callback function signals that the encoding operation on the frame has finished
virtual void onEndFrame(int frameNumber, PicType picType) = 0;
};
class Impl;
private:
cv::Ptr<Impl> impl_;
};
////////////////////////////////// Video Decoding //////////////////////////////////////////
namespace detail
{
class FrameQueue;
class VideoParser;
}
class CV_EXPORTS VideoReader_GPU
{
public:
enum Codec
{
MPEG1 = 0,
MPEG2,
MPEG4,
VC1,
H264,
JPEG,
H264_SVC,
H264_MVC,
Uncompressed_YUV420 = (('I'<<24)|('Y'<<16)|('U'<<8)|('V')), // Y,U,V (4:2:0)
Uncompressed_YV12 = (('Y'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,V,U (4:2:0)
Uncompressed_NV12 = (('N'<<24)|('V'<<16)|('1'<<8)|('2')), // Y,UV (4:2:0)
Uncompressed_YUYV = (('Y'<<24)|('U'<<16)|('Y'<<8)|('V')), // YUYV/YUY2 (4:2:2)
Uncompressed_UYVY = (('U'<<24)|('Y'<<16)|('V'<<8)|('Y')), // UYVY (4:2:2)
};
enum ChromaFormat
{
Monochrome=0,
YUV420,
YUV422,
YUV444,
};
struct FormatInfo
{
Codec codec;
ChromaFormat chromaFormat;
int width;
int height;
};
class VideoSource;
VideoReader_GPU();
explicit VideoReader_GPU(const String& filename);
explicit VideoReader_GPU(const cv::Ptr<VideoSource>& source);
~VideoReader_GPU();
void open(const String& filename);
void open(const cv::Ptr<VideoSource>& source);
bool isOpened() const;
void close();
bool read(GpuMat& image);
FormatInfo format() const;
void dumpFormat(std::ostream& st);
class CV_EXPORTS VideoSource
{
public:
VideoSource() : frameQueue_(0), videoParser_(0) {}
virtual ~VideoSource() {}
virtual FormatInfo format() const = 0;
virtual void start() = 0;
virtual void stop() = 0;
virtual bool isStarted() const = 0;
virtual bool hasError() const = 0;
void setFrameQueue(detail::FrameQueue* frameQueue) { frameQueue_ = frameQueue; }
void setVideoParser(detail::VideoParser* videoParser) { videoParser_ = videoParser; }
protected:
bool parseVideoData(const uchar* data, size_t size, bool endOfStream = false);
private:
VideoSource(const VideoSource&);
VideoSource& operator =(const VideoSource&);
detail::FrameQueue* frameQueue_;
detail::VideoParser* videoParser_;
};
class Impl;
private:
cv::Ptr<Impl> impl_;
};
}} // namespace cv { namespace gpu {
namespace cv {
template <> CV_EXPORTS void Ptr<cv::gpu::VideoWriter_GPU::Impl>::delete_obj();
template <> CV_EXPORTS void Ptr<cv::gpu::VideoReader_GPU::Impl>::delete_obj();
}
#endif /* __OPENCV_GPUCODEC_HPP__ */
modules/gpucodec/perf/perf_main.cpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "perf_precomp.hpp"
using namespace perf;
CV_PERF_TEST_MAIN(gpucodec, printCudaInfo())
modules/gpucodec/perf/perf_precomp.cpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "perf_precomp.hpp"
modules/gpucodec/perf/perf_precomp.hpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#ifdef __GNUC__
# pragma GCC diagnostic ignored "-Wmissing-declarations"
# if defined __clang__ || defined __APPLE__
# pragma GCC diagnostic ignored "-Wmissing-prototypes"
# pragma GCC diagnostic ignored "-Wextra"
# endif
#endif
#ifndef __OPENCV_PERF_PRECOMP_HPP__
#define __OPENCV_PERF_PRECOMP_HPP__
#include "opencv2/ts.hpp"
#include "opencv2/ts/gpu_perf.hpp"
#include "opencv2/gpucodec.hpp"
#include "opencv2/highgui.hpp"
#ifdef GTEST_CREATE_SHARED_LIBRARY
#error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined
#endif
#endif
modules/gpu/src/video_decoder.h → modules/gpucodec/perf/perf_video.cpp
View file @ 1b00a3ed
......@@ -40,77 +40,123 @@
//
//M*/
#ifndef __VIDEO_DECODER_H__
#define __VIDEO_DECODER_H__
#include "perf_precomp.hpp"
#include "opencv2/highgui/highgui_c.h"
#include "precomp.hpp"
using namespace std;
using namespace testing;
using namespace perf;
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#if defined(HAVE_XINE) || \
defined(HAVE_GSTREAMER) || \
defined(HAVE_QUICKTIME) || \
defined(HAVE_AVFOUNDATION) || \
defined(HAVE_FFMPEG) || \
defined(WIN32) /* assume that we have ffmpeg */
namespace cv { namespace gpu
# define BUILD_WITH_VIDEO_INPUT_SUPPORT 1
#else
# define BUILD_WITH_VIDEO_INPUT_SUPPORT 0
#endif
DEF_PARAM_TEST_1(FileName, string);
//////////////////////////////////////////////////////
// VideoReader
#if defined(HAVE_NVCUVID) && BUILD_WITH_VIDEO_INPUT_SUPPORT
PERF_TEST_P(FileName, VideoReader, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
{
namespace detail
declare.time(20);
const string inputFile = perf::TestBase::getDataPath(GetParam());
if (PERF_RUN_GPU())
{
class VideoDecoder
{
public:
VideoDecoder(const VideoReader_GPU::FormatInfo& videoFormat, CUvideoctxlock lock) : lock_(lock), decoder_(0)
{
create(videoFormat);
}
cv::gpu::VideoReader_GPU d_reader(inputFile);
ASSERT_TRUE( d_reader.isOpened() );
cv::gpu::GpuMat frame;
TEST_CYCLE_N(10) d_reader.read(frame);
GPU_SANITY_CHECK(frame);
}
else
{
cv::VideoCapture reader(inputFile);
ASSERT_TRUE( reader.isOpened() );
cv::Mat frame;
~VideoDecoder()
{
release();
}
TEST_CYCLE_N(10) reader >> frame;
void create(const VideoReader_GPU::FormatInfo& videoFormat);
void release();
CPU_SANITY_CHECK(frame);
}
}
#endif
// Get the code-type currently used.
cudaVideoCodec codec() const { return createInfo_.CodecType; }
unsigned long maxDecodeSurfaces() const { return createInfo_.ulNumDecodeSurfaces; }
//////////////////////////////////////////////////////
// VideoWriter
unsigned long frameWidth() const { return createInfo_.ulWidth; }
unsigned long frameHeight() const { return createInfo_.ulHeight; }
#if defined(HAVE_NVCUVID) && defined(WIN32)
unsigned long targetWidth() const { return createInfo_.ulTargetWidth; }
unsigned long targetHeight() const { return createInfo_.ulTargetHeight; }
PERF_TEST_P(FileName, VideoWriter, Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"))
{
declare.time(30);
cudaVideoChromaFormat chromaFormat() const { return createInfo_.ChromaFormat; }
const string inputFile = perf::TestBase::getDataPath(GetParam());
const string outputFile = cv::tempfile(".avi");
bool decodePicture(CUVIDPICPARAMS* picParams)
{
return cuvidDecodePicture(decoder_, picParams) == CUDA_SUCCESS;
}
const double FPS = 25.0;
cv::gpu::GpuMat mapFrame(int picIdx, CUVIDPROCPARAMS& videoProcParams)
{
CUdeviceptr ptr;
unsigned int pitch;
cv::VideoCapture reader(inputFile);
ASSERT_TRUE( reader.isOpened() );
cuSafeCall( cuvidMapVideoFrame(decoder_, picIdx, &ptr, &pitch, &videoProcParams) );
cv::Mat frame;
return GpuMat(targetHeight() * 3 / 2, targetWidth(), CV_8UC1, (void*) ptr, pitch);
}
if (PERF_RUN_GPU())
{
cv::gpu::VideoWriter_GPU d_writer;
cv::gpu::GpuMat d_frame;
for (int i = 0; i < 10; ++i)
{
reader >> frame;
ASSERT_FALSE(frame.empty());
void unmapFrame(cv::gpu::GpuMat& frame)
{
cuSafeCall( cuvidUnmapVideoFrame(decoder_, (CUdeviceptr) frame.data) );
frame.release();
}
d_frame.upload(frame);
if (!d_writer.isOpened())
d_writer.open(outputFile, frame.size(), FPS);
startTimer(); next();
d_writer.write(d_frame);
stopTimer();
}
}
else
{
cv::VideoWriter writer;
for (int i = 0; i < 10; ++i)
{
reader >> frame;
ASSERT_FALSE(frame.empty());
private:
VideoDecoder(const VideoDecoder&);
VideoDecoder& operator =(const VideoDecoder&);
if (!writer.isOpened())
writer.open(outputFile, CV_FOURCC('X', 'V', 'I', 'D'), FPS, frame.size());
CUvideoctxlock lock_;
CUVIDDECODECREATEINFO createInfo_;
CUvideodecoder decoder_;
};
startTimer(); next();
writer.write(frame);
stopTimer();
}
}
}}
#endif // HAVE_CUDA
SANITY_CHECK(frame);
}
#endif // __VIDEO_DECODER_H__
#endif
modules/gpucodec/src/cuda/nv12_to_rgb.cu 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
/*
* NV12ToARGB color space conversion CUDA kernel
*
* This sample uses CUDA to perform a simple NV12 (YUV 4:2:0 planar)
* source and converts to output in ARGB format
*/
#include "opencv2/core/cuda/common.hpp"
namespace cv { namespace gpu { namespace cudev
{
__constant__ float constHueColorSpaceMat[9];
void loadHueCSC(float hueCSC[9])
{
cudaSafeCall( cudaMemcpyToSymbol(constHueColorSpaceMat, hueCSC, 9 * sizeof(float)) );
}
__device__ void YUV2RGB(const uint* yuvi, float* red, float* green, float* blue)
{
float luma, chromaCb, chromaCr;
// Prepare for hue adjustment
luma = (float)yuvi[0];
chromaCb = (float)((int)yuvi[1] - 512.0f);
chromaCr = (float)((int)yuvi[2] - 512.0f);
// Convert YUV To RGB with hue adjustment
*red = (luma * constHueColorSpaceMat[0]) +
(chromaCb * constHueColorSpaceMat[1]) +
(chromaCr * constHueColorSpaceMat[2]);
*green = (luma * constHueColorSpaceMat[3]) +
(chromaCb * constHueColorSpaceMat[4]) +
(chromaCr * constHueColorSpaceMat[5]);
*blue = (luma * constHueColorSpaceMat[6]) +
(chromaCb * constHueColorSpaceMat[7]) +
(chromaCr * constHueColorSpaceMat[8]);
}
__device__ uint RGBA_pack_10bit(float red, float green, float blue, uint alpha)
{
uint ARGBpixel = 0;
// Clamp final 10 bit results
red = ::fmin(::fmax(red, 0.0f), 1023.f);
green = ::fmin(::fmax(green, 0.0f), 1023.f);
blue = ::fmin(::fmax(blue, 0.0f), 1023.f);
// Convert to 8 bit unsigned integers per color component
ARGBpixel = (((uint)blue >> 2) |
(((uint)green >> 2) << 8) |
(((uint)red >> 2) << 16) |
(uint)alpha);
return ARGBpixel;
}
// CUDA kernel for outputing the final ARGB output from NV12
#define COLOR_COMPONENT_BIT_SIZE 10
#define COLOR_COMPONENT_MASK 0x3FF
__global__ void NV12_to_RGB(uchar* srcImage, size_t nSourcePitch,
uint* dstImage, size_t nDestPitch,
uint width, uint height)
{
// Pad borders with duplicate pixels, and we multiply by 2 because we process 2 pixels per thread
const int x = blockIdx.x * (blockDim.x << 1) + (threadIdx.x << 1);
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= width || y >= height)
return;
// Read 2 Luma components at a time, so we don't waste processing since CbCr are decimated this way.
// if we move to texture we could read 4 luminance values
uint yuv101010Pel[2];
yuv101010Pel[0] = (srcImage[y * nSourcePitch + x ]) << 2;
yuv101010Pel[1] = (srcImage[y * nSourcePitch + x + 1]) << 2;
const size_t chromaOffset = nSourcePitch * height;
const int y_chroma = y >> 1;
if (y & 1) // odd scanline ?
{
uint chromaCb = srcImage[chromaOffset + y_chroma * nSourcePitch + x ];
uint chromaCr = srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1];
if (y_chroma < ((height >> 1) - 1)) // interpolate chroma vertically
{
chromaCb = (chromaCb + srcImage[chromaOffset + (y_chroma + 1) * nSourcePitch + x ] + 1) >> 1;
chromaCr = (chromaCr + srcImage[chromaOffset + (y_chroma + 1) * nSourcePitch + x + 1] + 1) >> 1;
}
yuv101010Pel[0] |= (chromaCb << ( COLOR_COMPONENT_BIT_SIZE + 2));
yuv101010Pel[0] |= (chromaCr << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
yuv101010Pel[1] |= (chromaCb << ( COLOR_COMPONENT_BIT_SIZE + 2));
yuv101010Pel[1] |= (chromaCr << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
}
else
{
yuv101010Pel[0] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x ] << ( COLOR_COMPONENT_BIT_SIZE + 2));
yuv101010Pel[0] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1] << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
yuv101010Pel[1] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x ] << ( COLOR_COMPONENT_BIT_SIZE + 2));
yuv101010Pel[1] |= ((uint)srcImage[chromaOffset + y_chroma * nSourcePitch + x + 1] << ((COLOR_COMPONENT_BIT_SIZE << 1) + 2));
}
// this steps performs the color conversion
uint yuvi[6];
float red[2], green[2], blue[2];
yuvi[0] = (yuv101010Pel[0] & COLOR_COMPONENT_MASK );
yuvi[1] = ((yuv101010Pel[0] >> COLOR_COMPONENT_BIT_SIZE) & COLOR_COMPONENT_MASK);
yuvi[2] = ((yuv101010Pel[0] >> (COLOR_COMPONENT_BIT_SIZE << 1)) & COLOR_COMPONENT_MASK);
yuvi[3] = (yuv101010Pel[1] & COLOR_COMPONENT_MASK );
yuvi[4] = ((yuv101010Pel[1] >> COLOR_COMPONENT_BIT_SIZE) & COLOR_COMPONENT_MASK);
yuvi[5] = ((yuv101010Pel[1] >> (COLOR_COMPONENT_BIT_SIZE << 1)) & COLOR_COMPONENT_MASK);
// YUV to RGB Transformation conversion
YUV2RGB(&yuvi[0], &red[0], &green[0], &blue[0]);
YUV2RGB(&yuvi[3], &red[1], &green[1], &blue[1]);
// Clamp the results to RGBA
const size_t dstImagePitch = nDestPitch >> 2;
dstImage[y * dstImagePitch + x ] = RGBA_pack_10bit(red[0], green[0], blue[0], ((uint)0xff << 24));
dstImage[y * dstImagePitch + x + 1 ] = RGBA_pack_10bit(red[1], green[1], blue[1], ((uint)0xff << 24));
}
void NV12_to_RGB(const PtrStepb decodedFrame, PtrStepSz<uint> interopFrame, cudaStream_t stream)
{
dim3 block(32, 8);
dim3 grid(divUp(interopFrame.cols, 2 * block.x), divUp(interopFrame.rows, block.y));
NV12_to_RGB<<<grid, block, 0, stream>>>(decodedFrame.data, decodedFrame.step, interopFrame.data, interopFrame.step,
interopFrame.cols, interopFrame.rows);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
}}}
modules/gpu/src/cuda/rgb_to_yv12.cu → modules/gpucodec/src/cuda/rgb_to_yv12.cu
View file @ 1b00a3ed
......@@ -40,136 +40,131 @@
//
//M*/
#if !defined CUDA_DISABLER
#include "opencv2/core/cuda/common.hpp"
#include "opencv2/core/cuda/vec_traits.hpp"
namespace cv { namespace gpu { namespace cudev
{
namespace video_encoding
__device__ __forceinline__ void rgb_to_y(const uchar b, const uchar g, const uchar r, uchar& y)
{
__device__ __forceinline__ void rgbtoy(const uchar b, const uchar g, const uchar r, uchar& y)
{
y = static_cast<uchar>(((int)(30 * r) + (int)(59 * g) + (int)(11 * b)) / 100);
}
__device__ __forceinline__ void rgbtoyuv(const uchar b, const uchar g, const uchar r, uchar& y, uchar& u, uchar& v)
{
rgbtoy(b, g, r, y);
u = static_cast<uchar>(((int)(-17 * r) - (int)(33 * g) + (int)(50 * b) + 12800) / 100);
v = static_cast<uchar>(((int)(50 * r) - (int)(42 * g) - (int)(8 * b) + 12800) / 100);
}
__global__ void Gray_to_YV12(const PtrStepSzb src, PtrStepb dst)
{
const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
if (x + 1 >= src.cols || y + 1 >= src.rows)
return;
// get pointers to the data
const size_t planeSize = src.rows * dst.step;
PtrStepb y_plane(dst.data, dst.step);
PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
uchar pix;
uchar y_val, u_val, v_val;
pix = src(y, x);
rgbtoy(pix, pix, pix, y_val);
y_plane(y, x) = y_val;
pix = src(y, x + 1);
rgbtoy(pix, pix, pix, y_val);
y_plane(y, x + 1) = y_val;
pix = src(y + 1, x);
rgbtoy(pix, pix, pix, y_val);
y_plane(y + 1, x) = y_val;
pix = src(y + 1, x + 1);
rgbtoyuv(pix, pix, pix, y_val, u_val, v_val);
y_plane(y + 1, x + 1) = y_val;
u_plane(y / 2, x / 2) = u_val;
v_plane(y / 2, x / 2) = v_val;
}
template <typename T>
__global__ void BGR_to_YV12(const PtrStepSz<T> src, PtrStepb dst)
{
const int x = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
const int y = (blockIdx.y * blockDim.y + threadIdx.y) * 2;
if (x + 1 >= src.cols || y + 1 >= src.rows)
return;
// get pointers to the data
const size_t planeSize = src.rows * dst.step;
PtrStepb y_plane(dst.data, dst.step);
PtrStepb u_plane(y_plane.data + planeSize, dst.step / 2);
PtrStepb v_plane(u_plane.data + (planeSize / 4), dst.step / 2);
T pix;
uchar y_val, u_val, v_val;
pix = src(y, x);
rgbtoy(pix.z, pix.y, pix.x, y_val);
y_plane(y, x) = y_val;
y = static_cast<uchar>(((int)(30 * r) + (int)(59 * g) + (int)(11 * b)) / 100);
}
pix = src(y, x + 1);
rgbtoy(pix.z, pix.y, pix.x, y_val);
y_plane(y, x + 1) = y_val;
__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);
}
pix = src(y + 1, x);
rgbtoy(pix.z, pix.y, pix.x, y_val);
y_plane(y + 1, x) = y_val;
__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;
}
pix = src(y + 1, x + 1);
rgbtoyuv(pix.z, pix.y, pix.x, y_val, u_val, v_val);
y_plane(y + 1, x + 1) = y_val;
u_plane(y / 2, x / 2) = u_val;
v_plane(y / 2, x / 2) = v_val;
}
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)
{
dim3 block(32, 8);
dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
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>>>(src, dst);
cudaSafeCall( cudaGetLastError() );
Gray_to_YV12<<<grid, block, 0, stream>>>(src, dst);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
template <int cn>
void BGR_to_YV12_caller(const PtrStepSzb src, PtrStepb dst)
{
typedef typename TypeVec<uchar, cn>::vec_type src_t;
}
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));
dim3 block(32, 8);
dim3 grid(divUp(src.cols, block.x * 2), divUp(src.rows, block.y * 2));
BGR_to_YV12<<<grid, block>>>(static_cast< PtrStepSz<src_t> >(src), dst);
cudaSafeCall( cudaGetLastError() );
RGB_to_YV12<<<grid, block, 0, stream>>>(static_cast< PtrStepSz<src_t> >(src), dst);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
}
void YV12_gpu(const PtrStepSzb src, int cn, PtrStepSzb dst)
{
typedef void (*func_t)(const PtrStepSzb src, PtrStepb dst);
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, BGR_to_YV12_caller<3>, BGR_to_YV12_caller<4>
};
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);
}
funcs[cn](src, dst, stream);
}
}}}
#endif /* CUDA_DISABLER */
modules/gpu/src/cuvid_video_source.cpp → modules/gpucodec/src/cuvid_video_source.cpp
View file @ 1b00a3ed
......@@ -40,9 +40,9 @@
//
//M*/
#include "cuvid_video_source.h"
#include "precomp.hpp"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#ifdef HAVE_NVCUVID
cv::gpu::detail::CuvidVideoSource::CuvidVideoSource(const String& fname)
{
......@@ -69,6 +69,11 @@ cv::gpu::detail::CuvidVideoSource::CuvidVideoSource(const String& fname)
format_.height = vidfmt.coded_height;
}
cv::gpu::detail::CuvidVideoSource::~CuvidVideoSource()
{
cuvidDestroyVideoSource(videoSource_);
}
cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::detail::CuvidVideoSource::format() const
{
return format_;
......@@ -101,4 +106,4 @@ int CUDAAPI cv::gpu::detail::CuvidVideoSource::HandleVideoData(void* userData, C
return thiz->parseVideoData(packet->payload, packet->payload_size, (packet->flags & CUVID_PKT_ENDOFSTREAM) != 0);
}
#endif // defined(HAVE_CUDA) && !defined(__APPLE__)
#endif // HAVE_NVCUVID
modules/gpu/src/cuvid_video_source.h → modules/gpucodec/src/cuvid_video_source.h
View file @ 1b00a3ed
......@@ -43,48 +43,44 @@
#ifndef __CUVUD_VIDEO_SOURCE_H__
#define __CUVUD_VIDEO_SOURCE_H__
#include "precomp.hpp"
#include "opencv2/core/gpu_private.hpp"
#include "opencv2/gpucodec.hpp"
#include "thread.h"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#include <nvcuvid.h>
namespace cv { namespace gpu
namespace cv { namespace gpu { namespace detail
{
namespace detail
{
class CuvidVideoSource : public VideoReader_GPU::VideoSource
{
public:
explicit CuvidVideoSource(const String& fname);
~CuvidVideoSource() { cuvidDestroyVideoSource(videoSource_); }
VideoReader_GPU::FormatInfo format() const;
void start();
void stop();
bool isStarted() const;
bool hasError() const;
class CuvidVideoSource : public VideoReader_GPU::VideoSource
{
public:
explicit CuvidVideoSource(const String& fname);
~CuvidVideoSource();
private:
CuvidVideoSource(const CuvidVideoSource&);
CuvidVideoSource& operator =(const CuvidVideoSource&);
VideoReader_GPU::FormatInfo format() const;
void start();
void stop();
bool isStarted() const;
bool hasError() const;
// Callback for handling packages of demuxed video data.
//
// Parameters:
// pUserData - Pointer to user data. We must pass a pointer to a
// VideoSourceData struct here, that contains a valid CUvideoparser
// and FrameQueue.
// pPacket - video-source data packet.
//
// NOTE: called from a different thread that doesn't not have a cuda context
//
static int CUDAAPI HandleVideoData(void* pUserData, CUVIDSOURCEDATAPACKET* pPacket);
private:
// Callback for handling packages of demuxed video data.
//
// Parameters:
// pUserData - Pointer to user data. We must pass a pointer to a
// VideoSourceData struct here, that contains a valid CUvideoparser
// and FrameQueue.
// pPacket - video-source data packet.
//
// NOTE: called from a different thread that doesn't not have a cuda context
//
static int CUDAAPI HandleVideoData(void* pUserData, CUVIDSOURCEDATAPACKET* pPacket);
CUvideosource videoSource_;
VideoReader_GPU::FormatInfo format_;
};
}
}}
CUvideosource videoSource_;
VideoReader_GPU::FormatInfo format_;
};
#endif // defined(HAVE_CUDA) && !defined(__APPLE__)
}}}
#endif // __CUVUD_VIDEO_SOURCE_H__
modules/gpu/src/ffmpeg_video_source.cpp → modules/gpucodec/src/ffmpeg_video_source.cpp
View file @ 1b00a3ed
......@@ -40,14 +40,12 @@
//
//M*/
#include "ffmpeg_video_source.h"
#include "precomp.hpp"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#ifdef HAVE_NVCUVID
#if defined(HAVE_FFMPEG) && defined(BUILD_SHARED_LIBS)
#if defined(HAVE_FFMPEG) && defined(BUILD_SHARED_LIBS) && !defined(WIN32)
#include "../src/cap_ffmpeg_impl.hpp"
#else
#include "../src/cap_ffmpeg_api.hpp"
#endif
namespace
......@@ -116,11 +114,6 @@ cv::gpu::detail::FFmpegVideoSource::FFmpegVideoSource(const String& fname) :
format_.height = height;
}
cv::gpu::detail::FFmpegVideoSource::~FFmpegVideoSource()
{
release_InputMediaStream_FFMPEG_p(stream_);
}
cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::detail::FFmpegVideoSource::format() const
{
return format_;
......@@ -130,14 +123,14 @@ void cv::gpu::detail::FFmpegVideoSource::start()
{
stop_ = false;
hasError_ = false;
thread_.reset(new Thread(readLoop, this));
thread_ = new Thread(readLoop, this);
}
void cv::gpu::detail::FFmpegVideoSource::stop()
{
stop_ = true;
thread_->wait();
thread_.reset();
thread_.release();
}
bool cv::gpu::detail::FFmpegVideoSource::isStarted() const
......@@ -179,4 +172,9 @@ void cv::gpu::detail::FFmpegVideoSource::readLoop(void* userData)
thiz->parseVideoData(0, 0, true);
}
template <> void cv::Ptr<InputMediaStream_FFMPEG>::delete_obj()
{
if (obj) release_InputMediaStream_FFMPEG_p(obj);
}
#endif // HAVE_CUDA
modules/gpu/src/ffmpeg_video_source.h → modules/gpucodec/src/ffmpeg_video_source.h
View file @ 1b00a3ed
......@@ -43,46 +43,40 @@
#ifndef __FFMPEG_VIDEO_SOURCE_H__
#define __FFMPEG_VIDEO_SOURCE_H__
#include "precomp.hpp"
#include "thread_wrappers.h"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#include "opencv2/gpucodec.hpp"
#include "thread.h"
struct InputMediaStream_FFMPEG;
namespace cv { namespace gpu
namespace cv { namespace gpu { namespace detail {
class FFmpegVideoSource : public VideoReader_GPU::VideoSource
{
namespace detail
{
class FFmpegVideoSource : public VideoReader_GPU::VideoSource
{
public:
FFmpegVideoSource(const String& fname);
~FFmpegVideoSource();
public:
FFmpegVideoSource(const String& fname);
VideoReader_GPU::FormatInfo format() const;
void start();
void stop();
bool isStarted() const;
bool hasError() const;
VideoReader_GPU::FormatInfo format() const;
void start();
void stop();
bool isStarted() const;
bool hasError() const;
private:
FFmpegVideoSource(const FFmpegVideoSource&);
FFmpegVideoSource& operator =(const FFmpegVideoSource&);
private:
VideoReader_GPU::FormatInfo format_;
VideoReader_GPU::FormatInfo format_;
cv::Ptr<InputMediaStream_FFMPEG> stream_;
InputMediaStream_FFMPEG* stream_;
cv::Ptr<Thread> thread_;
volatile bool stop_;
volatile bool hasError_;
std::auto_ptr<Thread> thread_;
volatile bool stop_;
volatile bool hasError_;
static void readLoop(void* userData);
};
static void readLoop(void* userData);
};
}
}}
}}}
#endif // HAVE_CUDA
namespace cv {
template <> void Ptr<InputMediaStream_FFMPEG>::delete_obj();
}
#endif // __CUVUD_VIDEO_SOURCE_H__
#endif // __FFMPEG_VIDEO_SOURCE_H__
modules/gpu/src/frame_queue.cpp → modules/gpucodec/src/frame_queue.cpp
View file @ 1b00a3ed
......@@ -40,9 +40,9 @@
//
//M*/
#include "frame_queue.h"
#include "precomp.hpp"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#ifdef HAVE_NVCUVID
cv::gpu::detail::FrameQueue::FrameQueue() :
endOfDecode_(0),
......@@ -79,7 +79,7 @@ void cv::gpu::detail::FrameQueue::enqueue(const CUVIDPARSERDISPINFO* picParams)
bool isFramePlaced = false;
{
CriticalSection::AutoLock autoLock(criticalSection_);
AutoLock autoLock(mtx_);
if (framesInQueue_ < MaximumSize)
{
......@@ -100,7 +100,7 @@ void cv::gpu::detail::FrameQueue::enqueue(const CUVIDPARSERDISPINFO* picParams)
bool cv::gpu::detail::FrameQueue::dequeue(CUVIDPARSERDISPINFO& displayInfo)
{
CriticalSection::AutoLock autoLock(criticalSection_);
AutoLock autoLock(mtx_);
if (framesInQueue_ > 0)
{
......@@ -114,4 +114,4 @@ bool cv::gpu::detail::FrameQueue::dequeue(CUVIDPARSERDISPINFO& displayInfo)
return false;
}
#endif // HAVE_CUDA
#endif // HAVE_NVCUVID
modules/gpu/src/frame_queue.h → modules/gpucodec/src/frame_queue.h
View file @ 1b00a3ed
......@@ -43,61 +43,55 @@
#ifndef __FRAME_QUEUE_H__
#define __FRAME_QUEUE_H__
#include "precomp.hpp"
#include "thread_wrappers.h"
#include "opencv2/core/utility.hpp"
#include "opencv2/core/gpu_private.hpp"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#include <nvcuvid.h>
namespace cv { namespace gpu
namespace cv { namespace gpu { namespace detail
{
namespace detail
{
class FrameQueue
{
public:
static const int MaximumSize = 20; // MAX_FRM_CNT;
FrameQueue();
class FrameQueue
{
public:
static const int MaximumSize = 20; // MAX_FRM_CNT;
void endDecode() { endOfDecode_ = true; }
bool isEndOfDecode() const { return endOfDecode_ != 0;}
FrameQueue();
// Spins until frame becomes available or decoding gets canceled.
// If the requested frame is available the method returns true.
// If decoding was interupted before the requested frame becomes
// available, the method returns false.
bool waitUntilFrameAvailable(int pictureIndex);
void endDecode() { endOfDecode_ = true; }
bool isEndOfDecode() const { return endOfDecode_ != 0;}
void enqueue(const CUVIDPARSERDISPINFO* picParams);
// Spins until frame becomes available or decoding gets canceled.
// If the requested frame is available the method returns true.
// If decoding was interupted before the requested frame becomes
// available, the method returns false.
bool waitUntilFrameAvailable(int pictureIndex);
// Deque the next frame.
// Parameters:
// displayInfo - New frame info gets placed into this object.
// Returns:
// true, if a new frame was returned,
// false, if the queue was empty and no new frame could be returned.
bool dequeue(CUVIDPARSERDISPINFO& displayInfo);
void enqueue(const CUVIDPARSERDISPINFO* picParams);
void releaseFrame(const CUVIDPARSERDISPINFO& picParams) { isFrameInUse_[picParams.picture_index] = false; }
// Deque the next frame.
// Parameters:
// displayInfo - New frame info gets placed into this object.
// Returns:
// true, if a new frame was returned,
// false, if the queue was empty and no new frame could be returned.
bool dequeue(CUVIDPARSERDISPINFO& displayInfo);
private:
FrameQueue(const FrameQueue&);
FrameQueue& operator =(const FrameQueue&);
void releaseFrame(const CUVIDPARSERDISPINFO& picParams) { isFrameInUse_[picParams.picture_index] = false; }
bool isInUse(int pictureIndex) const { return isFrameInUse_[pictureIndex] != 0; }
private:
bool isInUse(int pictureIndex) const { return isFrameInUse_[pictureIndex] != 0; }
CriticalSection criticalSection_;
Mutex mtx_;
volatile int isFrameInUse_[MaximumSize];
volatile int endOfDecode_;
volatile int isFrameInUse_[MaximumSize];
volatile int endOfDecode_;
int framesInQueue_;
int readPosition_;
CUVIDPARSERDISPINFO displayQueue_[MaximumSize];
};
}
}}
int framesInQueue_;
int readPosition_;
CUVIDPARSERDISPINFO displayQueue_[MaximumSize];
};
#endif // HAVE_CUDA
}}}
#endif // __FRAME_QUEUE_H__
modules/gpucodec/src/precomp.cpp 0 → 100644
View file @ 1b00a3ed
/*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"
modules/gpucodec/src/precomp.hpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#ifndef __OPENCV_PRECOMP_H__
#define __OPENCV_PRECOMP_H__
#include <cstdlib>
#include <cstring>
#include <deque>
#include <utility>
#include <stdexcept>
#include <iostream>
#include "opencv2/gpucodec.hpp"
#include "opencv2/core/gpu_private.hpp"
#ifdef HAVE_NVCUVID
#include <nvcuvid.h>
#ifdef WIN32
#define NOMINMAX
#include <windows.h>
#include <NVEncoderAPI.h>
#else
#include <pthread.h>
#include <unistd.h>
#endif
#include "thread.h"
#include "ffmpeg_video_source.h"
#include "cuvid_video_source.h"
#include "frame_queue.h"
#include "video_decoder.h"
#include "video_parser.h"
#include "../src/cap_ffmpeg_api.hpp"
#endif
#endif /* __OPENCV_PRECOMP_H__ */
modules/gpucodec/src/thread.cpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#ifdef HAVE_NVCUVID
using namespace cv::gpu::detail;
#ifdef WIN32
namespace
{
struct UserData
{
Thread::Func func;
void* param;
};
DWORD WINAPI WinThreadFunction(LPVOID lpParam)
{
UserData* userData = static_cast<UserData*>(lpParam);
userData->func(userData->param);
return 0;
}
}
class cv::gpu::detail::Thread::Impl
{
public:
Impl(Thread::Func func, void* userData)
{
userData_.func = func;
userData_.param = userData;
thread_ = CreateThread(
NULL, // default security attributes
0, // use default stack size
WinThreadFunction, // thread function name
&userData_, // argument to thread function
0, // use default creation flags
&threadId_); // returns the thread identifier
}
~Impl()
{
CloseHandle(thread_);
}
void wait()
{
WaitForSingleObject(thread_, INFINITE);
}
private:
UserData userData_;
HANDLE thread_;
DWORD threadId_;
};
#else
namespace
{
struct UserData
{
Thread::Func func;
void* param;
};
void* PThreadFunction(void* lpParam)
{
UserData* userData = static_cast<UserData*>(lpParam);
userData->func(userData->param);
return 0;
}
}
class cv::gpu::detail::Thread::Impl
{
public:
Impl(Thread::Func func, void* userData)
{
userData_.func = func;
userData_.param = userData;
pthread_create(&thread_, NULL, PThreadFunction, &userData_);
}
~Impl()
{
pthread_detach(thread_);
}
void wait()
{
pthread_join(thread_, NULL);
}
private:
pthread_t thread_;
UserData userData_;
};
#endif
cv::gpu::detail::Thread::Thread(Func func, void* userData) :
impl_(new Impl(func, userData))
{
}
void cv::gpu::detail::Thread::wait()
{
impl_->wait();
}
void cv::gpu::detail::Thread::sleep(int ms)
{
#ifdef WIN32
::Sleep(ms);
#else
::usleep(ms * 1000);
#endif
}
template <> void cv::Ptr<cv::gpu::detail::Thread::Impl>::delete_obj()
{
if (obj) delete obj;
}
#endif // HAVE_NVCUVID
modules/gpu/src/thread_wrappers.h → modules/gpucodec/src/thread.h
View file @ 1b00a3ed
......@@ -43,70 +43,31 @@
#ifndef __THREAD_WRAPPERS_H__
#define __THREAD_WRAPPERS_H__
#include "precomp.hpp"
#include "opencv2/core.hpp"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
namespace cv { namespace gpu { namespace detail {
namespace cv { namespace gpu
class Thread
{
namespace detail
{
class CriticalSection
{
public:
CriticalSection();
~CriticalSection();
public:
typedef void (*Func)(void* userData);
void enter();
void leave();
explicit Thread(Func func, void* userData = 0);
class AutoLock
{
public:
explicit AutoLock(CriticalSection& criticalSection) :
criticalSection_(criticalSection)
{
criticalSection_.enter();
}
void wait();
~AutoLock()
{
criticalSection_.leave();
}
static void sleep(int ms);
private:
CriticalSection& criticalSection_;
};
class Impl;
private:
CriticalSection(const CriticalSection&);
CriticalSection& operator=(const CriticalSection&);
private:
cv::Ptr<Impl> impl_;
};
class Impl;
std::auto_ptr<Impl> impl_;
};
}}}
class Thread
{
public:
explicit Thread(void (*func)(void* userData), void* userData = 0);
~Thread();
void wait();
static void sleep(int ms);
private:
Thread(const Thread&);
Thread& operator=(const Thread&);
class Impl;
std::auto_ptr<Impl> impl_;
};
}
}}
#endif // HAVE_CUDA
namespace cv {
template <> void Ptr<cv::gpu::detail::Thread::Impl>::delete_obj();
}
#endif // __THREAD_WRAPPERS_H__
modules/gpu/src/video_decoder.cpp → modules/gpucodec/src/video_decoder.cpp
View file @ 1b00a3ed
......@@ -40,10 +40,9 @@
//
//M*/
#include "video_decoder.h"
#include "frame_queue.h"
#include "precomp.hpp"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#ifdef HAVE_NVCUVID
void cv::gpu::detail::VideoDecoder::create(const VideoReader_GPU::FormatInfo& videoFormat)
{
......@@ -113,4 +112,4 @@ void cv::gpu::detail::VideoDecoder::release()
}
}
#endif // HAVE_CUDA
#endif // HAVE_NVCUVID
modules/gpucodec/src/video_decoder.h 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#ifndef __VIDEO_DECODER_H__
#define __VIDEO_DECODER_H__
#include "opencv2/core/gpu_private.hpp"
#include "opencv2/gpucodec.hpp"
#include <nvcuvid.h>
namespace cv { namespace gpu { namespace detail
{
class VideoDecoder
{
public:
VideoDecoder(const VideoReader_GPU::FormatInfo& videoFormat, CUvideoctxlock lock) : lock_(lock), decoder_(0)
{
create(videoFormat);
}
~VideoDecoder()
{
release();
}
void create(const VideoReader_GPU::FormatInfo& videoFormat);
void release();
// Get the code-type currently used.
cudaVideoCodec codec() const { return createInfo_.CodecType; }
unsigned long maxDecodeSurfaces() const { return createInfo_.ulNumDecodeSurfaces; }
unsigned long frameWidth() const { return createInfo_.ulWidth; }
unsigned long frameHeight() const { return createInfo_.ulHeight; }
unsigned long targetWidth() const { return createInfo_.ulTargetWidth; }
unsigned long targetHeight() const { return createInfo_.ulTargetHeight; }
cudaVideoChromaFormat chromaFormat() const { return createInfo_.ChromaFormat; }
bool decodePicture(CUVIDPICPARAMS* picParams)
{
return cuvidDecodePicture(decoder_, picParams) == CUDA_SUCCESS;
}
cv::gpu::GpuMat mapFrame(int picIdx, CUVIDPROCPARAMS& videoProcParams)
{
CUdeviceptr ptr;
unsigned int pitch;
cuSafeCall( cuvidMapVideoFrame(decoder_, picIdx, &ptr, &pitch, &videoProcParams) );
return GpuMat(targetHeight() * 3 / 2, targetWidth(), CV_8UC1, (void*) ptr, pitch);
}
void unmapFrame(cv::gpu::GpuMat& frame)
{
cuSafeCall( cuvidUnmapVideoFrame(decoder_, (CUdeviceptr) frame.data) );
frame.release();
}
private:
CUvideoctxlock lock_;
CUVIDDECODECREATEINFO createInfo_;
CUvideodecoder decoder_;
};
}}}
#endif // __VIDEO_DECODER_H__
modules/gpu/src/video_parser.cpp → modules/gpucodec/src/video_parser.cpp
View file @ 1b00a3ed
......@@ -40,9 +40,9 @@
//
//M*/
#include "video_parser.h"
#include "precomp.hpp"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#ifdef HAVE_NVCUVID
cv::gpu::detail::VideoParser::VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue) :
videoDecoder_(videoDecoder), frameQueue_(frameQueue), unparsedPackets_(0), hasError_(false)
......@@ -158,4 +158,4 @@ int CUDAAPI cv::gpu::detail::VideoParser::HandlePictureDisplay(void* userData, C
return true;
}
#endif // HAVE_CUDA
#endif // HAVE_NVCUVID
modules/gpu/src/video_parser.h → modules/gpucodec/src/video_parser.h
View file @ 1b00a3ed
......@@ -43,58 +43,52 @@
#ifndef __VIDEO_PARSER_H__
#define __VIDEO_PARSER_H__
#include "precomp.hpp"
#include "opencv2/core/gpu_private.hpp"
#include "opencv2/gpucodec.hpp"
#include "frame_queue.h"
#include "video_decoder.h"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#include <nvcuvid.h>
namespace cv { namespace gpu
namespace cv { namespace gpu { namespace detail
{
namespace detail
{
class VideoParser
{
public:
VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue);
~VideoParser()
{
cuvidDestroyVideoParser(parser_);
}
class VideoParser
{
public:
VideoParser(VideoDecoder* videoDecoder, FrameQueue* frameQueue);
bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream);
~VideoParser()
{
cuvidDestroyVideoParser(parser_);
}
bool hasError() const { return hasError_; }
bool parseVideoData(const unsigned char* data, size_t size, bool endOfStream);
private:
VideoParser(const VideoParser&);
VideoParser& operator =(const VideoParser&);
bool hasError() const { return hasError_; }
VideoDecoder* videoDecoder_;
FrameQueue* frameQueue_;
CUvideoparser parser_;
int unparsedPackets_;
volatile bool hasError_;
private:
VideoDecoder* videoDecoder_;
FrameQueue* frameQueue_;
CUvideoparser parser_;
int unparsedPackets_;
volatile bool hasError_;
// Called when the decoder encounters a video format change (or initial sequence header)
// This particular implementation of the callback returns 0 in case the video format changes
// to something different than the original format. Returning 0 causes a stop of the app.
static int CUDAAPI HandleVideoSequence(void* pUserData, CUVIDEOFORMAT* pFormat);
// Called when the decoder encounters a video format change (or initial sequence header)
// This particular implementation of the callback returns 0 in case the video format changes
// to something different than the original format. Returning 0 causes a stop of the app.
static int CUDAAPI HandleVideoSequence(void* pUserData, CUVIDEOFORMAT* pFormat);
// Called by the video parser to decode a single picture
// Since the parser will deliver data as fast as it can, we need to make sure that the picture
// index we're attempting to use for decode is no longer used for display
static int CUDAAPI HandlePictureDecode(void* pUserData, CUVIDPICPARAMS* pPicParams);
// Called by the video parser to decode a single picture
// Since the parser will deliver data as fast as it can, we need to make sure that the picture
// index we're attempting to use for decode is no longer used for display
static int CUDAAPI HandlePictureDecode(void* pUserData, CUVIDPICPARAMS* pPicParams);
// Called by the video parser to display a video frame (in the case of field pictures, there may be
// 2 decode calls per 1 display call, since two fields make up one frame)
static int CUDAAPI HandlePictureDisplay(void* pUserData, CUVIDPARSERDISPINFO* pPicParams);
};
}
}}
// Called by the video parser to display a video frame (in the case of field pictures, there may be
// 2 decode calls per 1 display call, since two fields make up one frame)
static int CUDAAPI HandlePictureDisplay(void* pUserData, CUVIDPARSERDISPINFO* pPicParams);
};
#endif // HAVE_CUDA
}}}
#endif // __VIDEO_PARSER_H__
modules/gpu/src/video_reader.cpp → modules/gpucodec/src/video_reader.cpp
View file @ 1b00a3ed
......@@ -42,7 +42,7 @@
#include "precomp.hpp"
#if !defined(HAVE_CUDA) || defined(CUDA_DISABLER) || !defined(HAVE_NVCUVID)
#ifndef HAVE_NVCUVID
class cv::gpu::VideoReader_GPU::Impl
{
......@@ -61,14 +61,7 @@ cv::gpu::VideoReader_GPU::FormatInfo cv::gpu::VideoReader_GPU::format() const {
bool cv::gpu::VideoReader_GPU::VideoSource::parseVideoData(const unsigned char*, size_t, bool) { throw_no_cuda(); return false; }
void cv::gpu::VideoReader_GPU::dumpFormat(std::ostream&) { throw_no_cuda(); }
#else // HAVE_CUDA
#include "frame_queue.h"
#include "video_decoder.h"
#include "video_parser.h"
#include "cuvid_video_source.h"
#include "ffmpeg_video_source.h"
#else // HAVE_NVCUVID
class cv::gpu::VideoReader_GPU::Impl
{
......@@ -81,14 +74,11 @@ public:
cv::gpu::VideoReader_GPU::FormatInfo format() const { return videoSource_->format(); }
private:
Impl(const Impl&);
Impl& operator =(const Impl&);
cv::Ptr<cv::gpu::VideoReader_GPU::VideoSource> videoSource_;
std::auto_ptr<cv::gpu::detail::FrameQueue> frameQueue_;
std::auto_ptr<cv::gpu::detail::VideoDecoder> videoDecoder_;
std::auto_ptr<cv::gpu::detail::VideoParser> videoParser_;
cv::Ptr<cv::gpu::detail::FrameQueue> frameQueue_;
cv::Ptr<cv::gpu::detail::VideoDecoder> videoDecoder_;
cv::Ptr<cv::gpu::detail::VideoParser> videoParser_;
CUvideoctxlock lock_;
......@@ -110,12 +100,12 @@ cv::gpu::VideoReader_GPU::Impl::Impl(const cv::Ptr<VideoSource>& source) :
cuSafeCall( cuCtxGetCurrent(&ctx) );
cuSafeCall( cuvidCtxLockCreate(&lock_, ctx) );
frameQueue_.reset(new detail::FrameQueue);
videoDecoder_.reset(new detail::VideoDecoder(videoSource_->format(), lock_));
videoParser_.reset(new detail::VideoParser(videoDecoder_.get(), frameQueue_.get()));
frameQueue_ = new detail::FrameQueue;
videoDecoder_ = new detail::VideoDecoder(videoSource_->format(), lock_);
videoParser_ = new detail::VideoParser(videoDecoder_, frameQueue_);
videoSource_->setFrameQueue(frameQueue_.get());
videoSource_->setVideoParser(videoParser_.get());
videoSource_->setFrameQueue(frameQueue_);
videoSource_->setVideoParser(videoParser_);
videoSource_->start();
}
......@@ -126,12 +116,10 @@ cv::gpu::VideoReader_GPU::Impl::~Impl()
videoSource_->stop();
}
namespace cv { namespace gpu { namespace cudev {
namespace video_decoding
{
void loadHueCSC(float hueCSC[9]);
void NV12ToARGB_gpu(const PtrStepb decodedFrame, PtrStepSz<unsigned int> interopFrame, cudaStream_t stream = 0);
}
namespace cv { namespace gpu { namespace cudev
{
void loadHueCSC(float hueCSC[9]);
void NV12_to_RGB(const PtrStepb decodedFrame, PtrStepSz<uint> interopFrame, cudaStream_t stream = 0);
}}}
namespace
......@@ -187,7 +175,7 @@ namespace
void cudaPostProcessFrame(const cv::gpu::GpuMat& decodedFrame, cv::gpu::GpuMat& interopFrame, int width, int height)
{
using namespace cv::gpu::cudev::video_decoding;
using namespace cv::gpu::cudev;
static bool updateCSC = true;
static float hueColorSpaceMat[9];
......@@ -210,7 +198,7 @@ namespace
loadHueCSC(hueColorSpaceMat);
NV12ToARGB_gpu(decodedFrame, interopFrame);
NV12_to_RGB(decodedFrame, interopFrame);
}
}
......@@ -329,17 +317,17 @@ void cv::gpu::VideoReader_GPU::open(const cv::Ptr<VideoSource>& source)
{
CV_Assert( !source.empty() );
close();
impl_.reset(new Impl(source));
impl_ = new Impl(source);
}
bool cv::gpu::VideoReader_GPU::isOpened() const
{
return impl_.get() != 0;
return !impl_.empty();
}
void cv::gpu::VideoReader_GPU::close()
{
impl_.reset();
impl_.release();
}
bool cv::gpu::VideoReader_GPU::read(GpuMat& image)
......@@ -396,4 +384,9 @@ void cv::gpu::VideoReader_GPU::dumpFormat(std::ostream& st)
st << "Chroma Format : " << chromas[_format.chromaFormat] << std::endl;
}
#endif // HAVE_CUDA
#endif // HAVE_NVCUVID
template <> void cv::Ptr<cv::gpu::VideoReader_GPU::Impl>::delete_obj()
{
if (obj) delete obj;
}
modules/gpu/src/video_writer.cpp → modules/gpucodec/src/video_writer.cpp
View file @ 1b00a3ed
......@@ -42,7 +42,7 @@
#include "precomp.hpp"
#if !defined(HAVE_CUDA) || defined(CUDA_DISABLER) || !defined(HAVE_NVCUVID) || !defined(WIN32)
#if !defined(HAVE_NVCUVID) || !defined(WIN32)
class cv::gpu::VideoWriter_GPU::Impl
{
......@@ -70,13 +70,6 @@ void cv::gpu::VideoWriter_GPU::EncoderParams::save(const String&) const { throw_
#else // !defined HAVE_CUDA || !defined WIN32
#ifdef HAVE_FFMPEG
#include "../src/cap_ffmpeg_impl.hpp"
#else
#include "../src/cap_ffmpeg_api.hpp"
#endif
///////////////////////////////////////////////////////////////////////////
// VideoWriter_GPU::Impl
......@@ -91,7 +84,7 @@ namespace
err = NVGetHWEncodeCaps();
if (err)
CV_Error(CV_GpuNotSupported, "No CUDA capability present");
CV_Error(cv::Error::GpuNotSupported, "No CUDA capability present");
// Create the Encoder API Interface
err = NVCreateEncoder(&encoder_);
......@@ -212,7 +205,7 @@ void cv::gpu::VideoWriter_GPU::Impl::initEncoder(double fps)
};
err = NVSetCodec(encoder_, codecs_id[codec_]);
if (err)
CV_Error(CV_StsNotImplemented, "Codec format is not supported");
CV_Error(cv::Error::StsNotImplemented, "Codec format is not supported");
// Set default params
......@@ -501,14 +494,6 @@ void cv::gpu::VideoWriter_GPU::Impl::createHWEncoder()
CV_Assert( err == 0 );
}
namespace cv { namespace gpu { namespace cudev
{
namespace video_encoding
{
void YV12_gpu(const PtrStepSzb src, int cn, PtrStepSzb dst);
}
}}}
namespace
{
// UYVY/YUY2 are both 4:2:2 formats (16bpc)
......@@ -644,6 +629,11 @@ namespace
}
}
namespace cv { namespace gpu { namespace cudev
{
void RGB_to_YV12(const PtrStepSzb src, int cn, PtrStepSzb dst, cudaStream_t stream = 0);
}}}
void cv::gpu::VideoWriter_GPU::Impl::write(const cv::gpu::GpuMat& frame, bool lastFrame)
{
if (inputFormat_ == SF_BGR)
......@@ -674,7 +664,7 @@ void cv::gpu::VideoWriter_GPU::Impl::write(const cv::gpu::GpuMat& frame, bool la
CV_Assert( res == CUDA_SUCCESS );
if (inputFormat_ == SF_BGR)
cv::gpu::cudev::video_encoding::YV12_gpu(frame, frame.channels(), videoFrame_);
cv::gpu::cudev::RGB_to_YV12(frame, frame.channels(), videoFrame_);
else
{
switch (surfaceFormat_)
......@@ -829,11 +819,14 @@ void EncoderCallBackFFMPEG::releaseBitStream(unsigned char* data, int size)
void EncoderCallBackFFMPEG::onBeginFrame(int frameNumber, PicType picType)
{
(void) frameNumber;
isKeyFrame_ = picType == IFRAME;
}
void EncoderCallBackFFMPEG::onEndFrame(int frameNumber, PicType picType)
{
(void) frameNumber;
(void) picType;
}
///////////////////////////////////////////////////////////////////////////
......@@ -885,23 +878,23 @@ void cv::gpu::VideoWriter_GPU::open(const String& fileName, cv::Size frameSize,
void cv::gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, SurfaceFormat format)
{
close();
impl_.reset(new Impl(encoderCallback, frameSize, fps, format));
impl_ = new Impl(encoderCallback, frameSize, fps, format);
}
void cv::gpu::VideoWriter_GPU::open(const cv::Ptr<EncoderCallBack>& encoderCallback, cv::Size frameSize, double fps, const EncoderParams& params, SurfaceFormat format)
{
close();
impl_.reset(new Impl(encoderCallback, frameSize, fps, params, format));
impl_ = new Impl(encoderCallback, frameSize, fps, params, format);
}
bool cv::gpu::VideoWriter_GPU::isOpened() const
{
return impl_.get() != 0;
return !impl_.empty();
}
void cv::gpu::VideoWriter_GPU::close()
{
impl_.reset();
impl_.release();
}
void cv::gpu::VideoWriter_GPU::write(const cv::gpu::GpuMat& image, bool lastFrame)
......@@ -1002,3 +995,8 @@ void cv::gpu::VideoWriter_GPU::EncoderParams::save(const String& configFile) con
}
#endif // !defined HAVE_CUDA || !defined WIN32
template <> void cv::Ptr<cv::gpu::VideoWriter_GPU::Impl>::delete_obj()
{
if (obj) delete obj;
}
modules/gpucodec/test/test_main.cpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "test_precomp.hpp"
CV_TEST_MAIN("gpu")
modules/gpucodec/test/test_precomp.cpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "test_precomp.hpp"
modules/gpucodec/test/test_precomp.hpp 0 → 100644
View file @ 1b00a3ed
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#ifdef __GNUC__
# pragma GCC diagnostic ignored "-Wmissing-declarations"
# if defined __clang__ || defined __APPLE__
# pragma GCC diagnostic ignored "-Wmissing-prototypes"
# pragma GCC diagnostic ignored "-Wextra"
# endif
#endif
#ifndef __OPENCV_TEST_PRECOMP_HPP__
#define __OPENCV_TEST_PRECOMP_HPP__
#include "opencv2/ts.hpp"
#include "opencv2/ts/gpu_test.hpp"
#include "opencv2/gpucodec.hpp"
#include "opencv2/highgui.hpp"
#endif
modules/gpu/test/test_video.cpp → modules/gpucodec/test/test_video.cpp
View file @ 1b00a3ed
......@@ -42,29 +42,21 @@
#include "test_precomp.hpp"
#if defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#ifdef HAVE_NVCUVID
PARAM_TEST_CASE(Video, cv::gpu::DeviceInfo, std::string)
{
};
//////////////////////////////////////////////////////
// VideoReader
PARAM_TEST_CASE(VideoReader, cv::gpu::DeviceInfo, std::string)
GPU_TEST_P(Video, Reader)
{
cv::gpu::DeviceInfo devInfo;
std::string inputFile;
cv::gpu::setDevice(GET_PARAM(0).deviceID());
virtual void SetUp()
{
devInfo = GET_PARAM(0);
inputFile = GET_PARAM(1);
cv::gpu::setDevice(devInfo.deviceID());
const std::string inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + GET_PARAM(1);
inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + inputFile;
}
};
GPU_TEST_P(VideoReader, Regression)
{
cv::gpu::VideoReader_GPU reader(inputFile);
ASSERT_TRUE(reader.isOpened());
......@@ -80,33 +72,17 @@ GPU_TEST_P(VideoReader, Regression)
ASSERT_FALSE(reader.isOpened());
}
INSTANTIATE_TEST_CASE_P(GPU_Video, VideoReader, testing::Combine(
ALL_DEVICES,
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi"))));
//////////////////////////////////////////////////////
// VideoWriter
#ifdef WIN32
PARAM_TEST_CASE(VideoWriter, cv::gpu::DeviceInfo, std::string)
GPU_TEST_P(Video, Writer)
{
cv::gpu::DeviceInfo devInfo;
std::string inputFile;
virtual void SetUp()
{
devInfo = GET_PARAM(0);
inputFile = GET_PARAM(1);
cv::gpu::setDevice(GET_PARAM(0).deviceID());
cv::gpu::setDevice(devInfo.deviceID());
inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + std::string("video/") + inputFile;
}
};
const std::string inputFile = std::string(cvtest::TS::ptr()->get_data_path()) + "video/" + GET_PARAM(1);
GPU_TEST_P(VideoWriter, Regression)
{
std::string outputFile = cv::tempfile(".avi");
const double FPS = 25.0;
......@@ -144,10 +120,10 @@ GPU_TEST_P(VideoWriter, Regression)
}
}
INSTANTIATE_TEST_CASE_P(GPU_Video, VideoWriter, testing::Combine(
#endif // WIN32
INSTANTIATE_TEST_CASE_P(GPU_Codec, Video, testing::Combine(
ALL_DEVICES,
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi"))));
#endif // WIN32
#endif // defined(HAVE_CUDA) && defined(HAVE_NVCUVID)
#endif // HAVE_NVCUVID
modules/superres/CMakeLists.txt
View file @ 1b00a3ed
......@@ -4,4 +4,4 @@ endif()
set(the_description "Super Resolution")
ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 -Wundef)
ocv_define_module(superres opencv_imgproc opencv_video OPTIONAL opencv_gpu opencv_highgui)
ocv_define_module(superres opencv_imgproc opencv_video OPTIONAL opencv_gpu opencv_highgui opencv_gpucodec)
modules/superres/src/frame_source.cpp
View file @ 1b00a3ed
......@@ -187,7 +187,7 @@ Ptr<FrameSource> cv::superres::createFrameSource_Camera(int deviceId)
//////////////////////////////////////////////////////
// VideoFrameSource_GPU
#ifndef HAVE_OPENCV_GPU
#ifndef HAVE_OPENCV_GPUCODEC
Ptr<FrameSource> cv::superres::createFrameSource_Video_GPU(const String& fileName)
{
......@@ -196,7 +196,7 @@ Ptr<FrameSource> cv::superres::createFrameSource_Video_GPU(const String& fileNam
return Ptr<FrameSource>();
}
#else // HAVE_OPENCV_GPU
#else // HAVE_OPENCV_GPUCODEC
namespace
{
......@@ -250,4 +250,4 @@ Ptr<FrameSource> cv::superres::createFrameSource_Video_GPU(const String& fileNam
return new VideoFrameSource(fileName);
}
#endif // HAVE_OPENCV_GPU
#endif // HAVE_OPENCV_GPUCODEC
modules/superres/src/precomp.hpp
View file @ 1b00a3ed
......@@ -60,6 +60,10 @@
# include "opencv2/core/gpu_private.hpp"
#endif
#ifdef HAVE_OPENCV_GPUCODEC
# include "opencv2/gpucodec.hpp"
#endif
#ifdef HAVE_OPENCV_HIGHGUI
#include "opencv2/highgui.hpp"
#endif
......
samples/gpu/CMakeLists.txt
View file @ 1b00a3ed
SET(OPENCV_GPU_SAMPLES_REQUIRED_DEPS opencv_core opencv_flann opencv_imgproc opencv_highgui
opencv_ml opencv_video opencv_objdetect opencv_features2d
opencv_calib3d opencv_legacy opencv_contrib opencv_gpu
opencv_nonfree opencv_softcascade opencv_superres)
opencv_nonfree opencv_softcascade opencv_superres
opencv_gpucodec)
ocv_check_dependencies(${OPENCV_GPU_SAMPLES_REQUIRED_DEPS})
......
samples/gpu/video_reader.cpp
View file @ 1b00a3ed
......@@ -4,11 +4,11 @@
#include <algorithm>
#include <numeric>
#include <opencv2/core/core.hpp>
#include <opencv2/core.hpp>
#include <opencv2/core/opengl.hpp>
#include <opencv2/gpu/gpu.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/contrib/contrib.hpp>
#include <opencv2/gpucodec.hpp>
#include <opencv2/highgui.hpp>
#include <opencv2/contrib.hpp>
int main(int argc, const char* argv[])
{
......
samples/gpu/video_writer.cpp
View file @ 1b00a3ed
......@@ -2,10 +2,10 @@
#include <vector>
#include <numeric>
#include "opencv2/core/core.hpp"
#include "opencv2/gpu/gpu.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/contrib/contrib.hpp"
#include "opencv2/core.hpp"
#include "opencv2/gpucodec.hpp"
#include "opencv2/highgui.hpp"
#include "opencv2/contrib.hpp"
int main(int argc, const char* argv[])
{
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment