Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
42 changes: 40 additions & 2 deletions modules/cudacodec/include/opencv2/cudacodec.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -385,6 +385,8 @@ struct CV_EXPORTS_W_SIMPLE FormatInfo
*/
class CV_EXPORTS_W NVSurfaceToColorConverter {
public:
virtual ~NVSurfaceToColorConverter() {}

/** @brief Performs the conversion from the raw YUV Surface output from VideoReader to the requested color format. Use this function when you want to convert the raw YUV Surface output from VideoReader to more than one color format or you want both the raw Surface output in addition to a color frame.
* @param yuv The raw YUV Surface output from VideoReader see @ref SurfaceFormat.
* @param color The converted frame.
Expand Down Expand Up @@ -551,8 +553,7 @@ class CV_EXPORTS_W VideoReader
- Out: Value of the property.
@return `true` unless the property is not supported.
*/
virtual bool get(const VideoReaderProps propertyId, double& propertyVal) const = 0;
CV_WRAP virtual bool getVideoReaderProps(const VideoReaderProps propertyId, CV_OUT double& propertyValOut, double propertyValIn = 0) const = 0;
CV_WRAP_AS(getVideoReaderProps) virtual bool get(const VideoReaderProps propertyId, CV_OUT size_t& propertyVal) const = 0;

/** @brief Retrieves the specified property used by the VideoSource.

Expand All @@ -563,6 +564,43 @@ class CV_EXPORTS_W VideoReader
@return `true` unless the property is unset set or not supported.
*/
CV_WRAP virtual bool get(const int propertyId, CV_OUT double& propertyVal) const = 0;

/** @brief Determine whether the raw package at \a idx contains encoded data for a key frame.

@param idx Index of the encoded package to check.

@returns `true` if the raw package at \a idx contains encoded data for a key frame and `false` otherwise.

@note A typical use case is deciding to archive live video after streaming has been initialized. In this scenario you would not want to write any encoded frames before a key frame. This is simulated in the example below where VideoReader is initialized without enabling raw mode
```
VideoReaderInitParams params;
params.rawMode = false;
Ptr<VideoReader> reader = createVideoReader(rtspUrl, {}, params);
```
and then at some point in the future raw mode is enabled to enable the footage to be archived
```
reader->set(VideoReaderProps::PROP_RAW_MODE, true);
```
Because raw mode has been enabled mid stream the first raw package retrieved now is not guaranteed to contain a key frame. To locate the next raw package containing a key frame rawPackageHasKeyFrame() can then be used as shown below.
```
double iRawPacketBase = -1;
reader->get(VideoReaderProps::PROP_RAW_PACKAGES_BASE_INDEX, iRawPacketBase);
GpuMat frame;
while (reader->nextFrame(frame)) {
double nRawPackets = -1;
reader->get(VideoReaderProps::PROP_NUMBER_OF_RAW_PACKAGES_SINCE_LAST_GRAB, nRawPackets);
for (int iRawPacketToWrite = static_cast<int>(iRawPacketBase); iRawPacketToWrite < static_cast<int>(iRawPacketBase + nRawPackets); iRawPacketToWrite++) {
if (reader->rawPackageHasKeyFrame(iRawPacketToWrite)) {
Mat packageToWrite;
reader->retrieve(packageToWrite, iRawPacketToWrite);
...
}
}
}
```
\sa retrieve
*/
CV_WRAP virtual bool rawPackageHasKeyFrame(const size_t idx) const = 0;
};

/** @brief Interface for video demultiplexing. :
Expand Down
46 changes: 23 additions & 23 deletions modules/cudacodec/misc/python/test/test_cudacodec.py
Original file line number Diff line number Diff line change
Expand Up @@ -48,37 +48,37 @@ def test_reader(self):
ret, raw_mode = reader.getVideoReaderProps(cv.cudacodec.VideoReaderProps_PROP_RAW_MODE)
self.assertTrue(ret and raw_mode)

# Retrieve image histogram. Not all GPUs support histogram. Just check the method is called correctly
ret, gpu_mat, hist = reader.nextFrameWithHist()
self.assertTrue(ret and not gpu_mat.empty())
ret, gpu_mat_, hist_ = reader.nextFrameWithHist(gpu_mat, hist)
self.assertTrue(ret and not gpu_mat.empty())
self.assertTrue(gpu_mat_.cudaPtr() == gpu_mat.cudaPtr())
# Read raw encoded bitstream
ret, i_base = reader.getVideoReaderProps(cv.cudacodec.VideoReaderProps_PROP_RAW_PACKAGES_BASE_INDEX)
self.assertTrue(ret and i_base == 2.0)
ret, gpu_mat_2 = reader.nextFrame()
self.assertTrue(ret and isinstance(gpu_mat_2,cv.cuda.GpuMat) and not gpu_mat_2.empty())
ret = reader.retrieve(gpu_mat_2)
self.assertTrue(ret and isinstance(gpu_mat_2,cv.cuda.GpuMat) and not gpu_mat_2.empty())
ret, n_raw_packages_since_last_grab = reader.getVideoReaderProps(cv.cudacodec.VideoReaderProps_PROP_NUMBER_OF_RAW_PACKAGES_SINCE_LAST_GRAB)
self.assertTrue(ret and n_raw_packages_since_last_grab > 0)
self.assertTrue(reader.rawPackageHasKeyFrame(int(i_base)))
ret, raw_data = reader.retrieve(int(i_base))
self.assertTrue(ret and isinstance(raw_data,np.ndarray) and np.any(raw_data))

# Check post processing applied
self.assertTrue(gpu_mat.size() == post_processed_sz)
self.assertTrue(gpu_mat_2.size() == post_processed_sz)

# Retrieve image histogram. Not all GPUs support histogram. Just check the method is called correctly
ret, gpu_mat_3, hist = reader.nextFrameWithHist()
self.assertTrue(ret and not gpu_mat_3.empty())
ret, gpu_mat_3_, hist_ = reader.nextFrameWithHist(gpu_mat_3, hist)
self.assertTrue(ret and not gpu_mat_3.empty())
self.assertTrue(gpu_mat_3_.cudaPtr() == gpu_mat_3.cudaPtr())

# Change color format
ret, colour_code = reader.getVideoReaderProps(cv.cudacodec.VideoReaderProps_PROP_COLOR_FORMAT)
self.assertTrue(ret and colour_code == cv.cudacodec.ColorFormat_BGRA)
colour_code_gs = cv.cudacodec.ColorFormat_GRAY
self.assertTrue(ret and colour_code == cv.cudacodec.BGRA)
colour_code_gs = cv.cudacodec.GRAY
reader.set(colour_code_gs)
ret, colour_code = reader.getVideoReaderProps(cv.cudacodec.VideoReaderProps_PROP_COLOR_FORMAT)
self.assertTrue(ret and colour_code == colour_code_gs)

# Read raw encoded bitstream
ret, i_base = reader.getVideoReaderProps(cv.cudacodec.VideoReaderProps_PROP_RAW_PACKAGES_BASE_INDEX)
self.assertTrue(ret and i_base == 2.0)
self.assertTrue(reader.grab())
ret, gpu_mat3 = reader.retrieve()
self.assertTrue(ret and isinstance(gpu_mat3,cv.cuda.GpuMat) and not gpu_mat3.empty())
ret = reader.retrieve(gpu_mat3)
self.assertTrue(ret and isinstance(gpu_mat3,cv.cuda.GpuMat) and not gpu_mat3.empty())
ret, n_raw_packages_since_last_grab = reader.getVideoReaderProps(cv.cudacodec.VideoReaderProps_PROP_NUMBER_OF_RAW_PACKAGES_SINCE_LAST_GRAB)
self.assertTrue(ret and n_raw_packages_since_last_grab > 0)
ret, raw_data = reader.retrieve(int(i_base))
self.assertTrue(ret and isinstance(raw_data,np.ndarray) and np.any(raw_data))

except cv.error as e:
notSupported = (e.code == cv.Error.StsNotImplemented or e.code == cv.Error.StsUnsupportedFormat or e.code == cv.Error.GPU_API_CALL_ERROR)
self.assertTrue(notSupported)
Expand Down Expand Up @@ -107,7 +107,7 @@ def test_writer(self):
encoder_params_in.gopLength = 10
stream = cv.cuda.Stream()
sz = (1920,1080)
writer = cv.cudacodec.createVideoWriter(fname, sz, cv.cudacodec.H264, 30, cv.cudacodec.ColorFormat_BGR,
writer = cv.cudacodec.createVideoWriter(fname, sz, cv.cudacodec.H264, 30, cv.cudacodec.BGR,
encoder_params_in, stream=stream)
blankFrameIn = cv.cuda.GpuMat(sz,cv.CV_8UC3)
writer.write(blankFrameIn)
Expand Down
18 changes: 8 additions & 10 deletions modules/cudacodec/src/NvEncoder.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -101,11 +101,9 @@ void NvEncoder::CreateDefaultEncoderParams(NV_ENC_INITIALIZE_PARAMS* pIntializeP
#endif
pIntializeParams->tuningInfo = tuningInfo;
pIntializeParams->encodeConfig->rcParams.rateControlMode = NV_ENC_PARAMS_RC_CONSTQP;
#if ((NVENCAPI_MAJOR_VERSION == 12 && NVENCAPI_MINOR_VERSION >= 2) || NVENCAPI_MAJOR_VERSION > 12)
NV_ENC_PRESET_CONFIG presetConfig = { NV_ENC_PRESET_CONFIG_VER, 0, { NV_ENC_CONFIG_VER } };
#else
NV_ENC_PRESET_CONFIG presetConfig = { NV_ENC_PRESET_CONFIG_VER, { NV_ENC_CONFIG_VER } };
#endif
NV_ENC_PRESET_CONFIG presetConfig = {};
presetConfig.version = NV_ENC_PRESET_CONFIG_VER;
presetConfig.presetCfg.version = NV_ENC_CONFIG_VER;
m_nvenc.nvEncGetEncodePresetConfigEx(m_hEncoder, codecGuid, presetGuid, tuningInfo, &presetConfig);
memcpy(pIntializeParams->encodeConfig, &presetConfig.presetCfg, sizeof(NV_ENC_CONFIG));

Expand Down Expand Up @@ -205,11 +203,9 @@ void NvEncoder::CreateEncoder(const NV_ENC_INITIALIZE_PARAMS* pEncoderParams)
}
else
{
#if ((NVENCAPI_MAJOR_VERSION == 12 && NVENCAPI_MINOR_VERSION >= 2) || NVENCAPI_MAJOR_VERSION > 12)
NV_ENC_PRESET_CONFIG presetConfig = { NV_ENC_PRESET_CONFIG_VER, 0, { NV_ENC_CONFIG_VER } };
#else
NV_ENC_PRESET_CONFIG presetConfig = { NV_ENC_PRESET_CONFIG_VER, { NV_ENC_CONFIG_VER } };
#endif
NV_ENC_PRESET_CONFIG presetConfig = {};
presetConfig.version = NV_ENC_PRESET_CONFIG_VER;
presetConfig.presetCfg.version = NV_ENC_CONFIG_VER;
m_nvenc.nvEncGetEncodePresetConfigEx(m_hEncoder, pEncoderParams->encodeGUID, pEncoderParams->presetGUID, pEncoderParams->tuningInfo, &presetConfig);
memcpy(&m_encodeConfig, &presetConfig.presetCfg, sizeof(NV_ENC_CONFIG));
}
Expand Down Expand Up @@ -570,6 +566,8 @@ void NvEncoder::WaitForCompletionEvent(int iEvent)
NVENC_THROW_ERROR("Failed to encode frame", NV_ENC_ERR_GENERIC);
}
#endif
#else
CV_UNUSED(iEvent);
#endif
}

Expand Down
5 changes: 4 additions & 1 deletion modules/cudacodec/src/NvEncoder.h
Original file line number Diff line number Diff line change
Expand Up @@ -263,7 +263,10 @@ class NvEncoder
/**
* @brief This function returns the completion event.
*/
void* GetCompletionEvent(uint32_t eventIdx) { return (m_vpCompletionEvent.size() == m_nEncoderBuffer) ? m_vpCompletionEvent[eventIdx] : nullptr; }
void* GetCompletionEvent(uint32_t eventIdx) {
CV_Assert(m_nEncoderBuffer >= 0);
return (m_vpCompletionEvent.size() == static_cast<size_t>(m_nEncoderBuffer)) ? m_vpCompletionEvent[eventIdx] : nullptr;
}

/**
* @brief This function returns the current pixel format.
Expand Down
23 changes: 14 additions & 9 deletions modules/cudacodec/src/cuda/ColorSpace.cu
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,11 @@

namespace cv { namespace cuda { namespace device {

void SetMatYuv2Rgb(int iMatrix, bool fullRange = false);
void Y8ToGray8(uint8_t* dpY8, int nY8Pitch, uint8_t* dpGray, int nGrayPitch, int nWidth, int nHeight, bool videoFullRangeFlag, const cudaStream_t stream);
void Y8ToGray16(uint8_t* dpY8, int nY8Pitch, uint8_t* dpGray, int nGrayPitch, int nWidth, int nHeight, bool videoFullRangeFlag, const cudaStream_t stream);
void Y16ToGray8(uint8_t* dpY16, int nY16Pitch, uint8_t* dpGray, int nGrayPitch, int nWidth, int nHeight, bool videoFullRangeFlag, const cudaStream_t stream);
void Y16ToGray16(uint8_t* dpY16, int nY16Pitch, uint8_t* dpGray, int nGrayPitch, int nWidth, int nHeight, bool videoFullRangeFlag, const cudaStream_t stream);
__constant__ float matYuv2Color[3][3];

void inline GetConstants(int iMatrix, float& wr, float& wb, int& black, int& white, int& uvWhite, int& max, bool fullRange = false) {
Expand Down Expand Up @@ -49,7 +54,7 @@ void inline GetConstants(int iMatrix, float& wr, float& wb, int& black, int& whi
}
}

void SetMatYuv2Rgb(int iMatrix, bool fullRange = false) {
void SetMatYuv2Rgb(int iMatrix, bool fullRange) {
float wr, wb;
int black, white, max, uvWhite;
GetConstants(iMatrix, wr, wb, black, white, uvWhite, max, fullRange);
Expand Down Expand Up @@ -160,15 +165,15 @@ __global__ static void YuvToColorKernel(uint8_t* pYuv, int nYuvPitch, uint8_t* p

union ColorOutx2 {
Colorx2 d;
Color Color[2];
Color color[2];
};
ColorOutx2 l1Out;
l1Out.Color[0] = YuvToColorForPixel<Color>(l0.x, ch.x, ch.y, videoFullRangeFlag);
l1Out.Color[1] = YuvToColorForPixel<Color>(l0.y, ch.x, ch.y, videoFullRangeFlag);
l1Out.color[0] = YuvToColorForPixel<Color>(l0.x, ch.x, ch.y, videoFullRangeFlag);
l1Out.color[1] = YuvToColorForPixel<Color>(l0.y, ch.x, ch.y, videoFullRangeFlag);
*(Colorx2*)pDst = l1Out.d;
ColorOutx2 l2Out;
l2Out.Color[0] = YuvToColorForPixel<Color>(l1.x, ch.x, ch.y, videoFullRangeFlag);
l2Out.Color[1] = YuvToColorForPixel<Color>(l1.y, ch.x, ch.y, videoFullRangeFlag);
l2Out.color[0] = YuvToColorForPixel<Color>(l1.x, ch.x, ch.y, videoFullRangeFlag);
l2Out.color[1] = YuvToColorForPixel<Color>(l1.y, ch.x, ch.y, videoFullRangeFlag);
*(Colorx2*)(pDst + nColorPitch) = l2Out.d;
}

Expand Down Expand Up @@ -214,11 +219,11 @@ __global__ static void Yuv444ToColorKernel(uint8_t* pYuv, int nYuvPitch, uint8_t

union ColorOutx2 {
Colorx2 d;
Color Color[2];
Color color[2];
};
ColorOutx2 out;
out.Color[0] = YuvToColorForPixel<Color>(l0.x, ch1.x, ch2.x, videoFullRangeFlag);
out.Color[1] = YuvToColorForPixel<Color>(l0.y, ch1.y, ch2.y, videoFullRangeFlag);
out.color[0] = YuvToColorForPixel<Color>(l0.x, ch1.x, ch2.x, videoFullRangeFlag);
out.color[1] = YuvToColorForPixel<Color>(l0.y, ch1.y, ch2.y, videoFullRangeFlag);
*(Colorx2*)pDst = out.d;
}

Expand Down
2 changes: 1 addition & 1 deletion modules/cudacodec/src/cuvid_video_source.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,7 @@ cv::cudacodec::detail::CuvidVideoSource::CuvidVideoSource(const String& fname)
CUVIDEOFORMAT vidfmt;
cuSafeCall( cuvidGetSourceVideoFormat(videoSource_, &vidfmt, 0) );

CV_Assert(Codec::NumCodecs == cudaVideoCodec::cudaVideoCodec_NumCodecs);
CV_Assert(static_cast<int>(Codec::NumCodecs) == static_cast<int>(cudaVideoCodec::cudaVideoCodec_NumCodecs));
format_.codec = static_cast<Codec>(vidfmt.codec);
format_.chromaFormat = static_cast<ChromaFormat>(vidfmt.chroma_format);
format_.nBitDepthMinus8 = vidfmt.bit_depth_luma_minus8;
Expand Down
1 change: 1 addition & 0 deletions modules/cudacodec/src/cuvid_video_source.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -56,6 +56,7 @@ class CuvidVideoSource : public VideoSource

FormatInfo format() const CV_OVERRIDE;
void updateFormat(const FormatInfo& videoFormat) CV_OVERRIDE;
bool get(const int, double&) const { return false; }
void start() CV_OVERRIDE;
void stop() CV_OVERRIDE;
bool isStarted() const CV_OVERRIDE;
Expand Down
12 changes: 7 additions & 5 deletions modules/cudacodec/src/ffmpeg_video_source.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -90,7 +90,7 @@ Codec FourccToCodec(int codec)
}

static
int StartCodeLen(unsigned char* data, const int sz) {
int StartCodeLen(unsigned char* data, const size_t sz) {
if (sz >= 3 && data[0] == 0 && data[1] == 0 && data[2] == 1)
return 3;
else if (sz >= 4 && data[0] == 0 && data[1] == 0 && data[2] == 0 && data[3] == 1)
Expand All @@ -99,7 +99,8 @@ int StartCodeLen(unsigned char* data, const int sz) {
return 0;
}

bool ParamSetsExist(unsigned char* parameterSets, const int szParameterSets, unsigned char* data, const int szData) {
static
bool ParamSetsExist(unsigned char* parameterSets, const size_t szParameterSets, unsigned char* data, const size_t szData) {
const int paramSetStartCodeLen = StartCodeLen(parameterSets, szParameterSets);
const int packetStartCodeLen = StartCodeLen(data, szData);
// weak test to see if the parameter set has already been included in the RTP stream
Expand Down Expand Up @@ -129,8 +130,8 @@ cv::cudacodec::detail::FFmpegVideoSource::FFmpegVideoSource(const String& fname,

int codec = (int)cap.get(CAP_PROP_FOURCC);
format_.codec = FourccToCodec(codec);
format_.height = cap.get(CAP_PROP_FRAME_HEIGHT);
format_.width = cap.get(CAP_PROP_FRAME_WIDTH);
format_.height = static_cast<int>(cap.get(CAP_PROP_FRAME_HEIGHT));
format_.width = static_cast<int>(cap.get(CAP_PROP_FRAME_WIDTH));
format_.displayArea = Rect(0, 0, format_.width, format_.height);
format_.valid = false;
format_.fps = cap.get(CAP_PROP_FPS);
Expand Down Expand Up @@ -181,7 +182,8 @@ bool cv::cudacodec::detail::FFmpegVideoSource::getNextPacket(unsigned char** dat
{
const size_t nBytesToTrimFromData = format_.codec == Codec::MPEG4 ? 3 : 0;
const size_t newSz = extraData.total() + *size - nBytesToTrimFromData;
dataWithHeader = Mat(1, newSz, CV_8UC1);
CV_Assert(newSz <= std::numeric_limits<int>::max());
dataWithHeader = Mat(1, static_cast<int>(newSz), CV_8UC1);
memcpy(dataWithHeader.data, extraData.data, extraData.total());
memcpy(dataWithHeader.data + extraData.total(), (*data) + nBytesToTrimFromData, *size - nBytesToTrimFromData);
*data = dataWithHeader.data;
Expand Down
Loading
Loading