AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Add 'projects/rocdecode/' from commit 'b0bab079403eda171f9056409fa96b0908f61073'

Przeglądaj źródła 
git-subtree-dir: projects/rocdecode
git-subtree-mainline: 5d609c1e57
git-subtree-split: b0bab07940
This commit is contained in:
Ameya Keshava Mallya
2026-01-30 20:33:26 +00:00
rodzic 5d609c1e57 b0bab07940
commit d0396f30b3
208 zmienionych plików z 43533 dodań i 0 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/rocdecode/utils/colorspace_kernels.cpp
+593
Wyświetl plik
@@ -0,0 +1,593 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "colorspace_kernels.h"
#include "roc_video_dec.h"
__constant__ float yuv_to_rgb_mat[3][3];
__constant__ float rgb_to_yuv_mat[3][3];
void inline GetColMatCoefficients(int col_standard, float &wr, float &wb, int &black, int &white, int &max) {
black = 16; white = 235;
max = 255;
switch (col_standard)
{
case ColorSpaceStandard_BT709:
default:
wr = 0.2126f; wb = 0.0722f;
break;
case ColorSpaceStandard_FCC:
wr = 0.30f; wb = 0.11f;
break;
case ColorSpaceStandard_BT470:
case ColorSpaceStandard_BT601:
wr = 0.2990f; wb = 0.1140f;
break;
case ColorSpaceStandard_SMPTE240M:
wr = 0.212f; wb = 0.087f;
break;
case ColorSpaceStandard_BT2020:
case ColorSpaceStandard_BT2020C:
wr = 0.2627f; wb = 0.0593f;
// 10-bit only
black = 64 << 6; white = 940 << 6;
max = (1 << 16) - 1;
break;
}
}
void SetMatYuv2Rgb(int col_standard) {
float wr, wb;
int black, white, max;
GetColMatCoefficients(col_standard, wr, wb, black, white, max);
float mat[3][3] = {
1.0f, 0.0f, (1.0f - wr) / 0.5f,
1.0f, -wb * (1.0f - wb) / 0.5f / (1 - wb - wr), -wr * (1 - wr) / 0.5f / (1 - wb - wr),
1.0f, (1.0f - wb) / 0.5f, 0.0f,
};
for (int i = 0; i < 3; i++) {
for (int j = 0; j < 3; j++) {
mat[i][j] = (float)(1.0 * max / (white - black) * mat[i][j]);
}
}
HIP_API_CALL(hipMemcpyToSymbol(yuv_to_rgb_mat, mat, sizeof(mat)));
}
void SetMatRgb2Yuv(int col_standard) {
float wr, wb;
int black, white, max;
GetColMatCoefficients(col_standard, wr, wb, black, white, max);
float mat[3][3] = {
wr, 1.0f - wb - wr, wb,
-0.5f * wr / (1.0f - wb), -0.5f * (1 - wb - wr) / (1.0f - wb), 0.5f,
0.5f, -0.5f * (1.0f - wb - wr) / (1.0f - wr), -0.5f * wb / (1.0f - wr),
};
for (int i = 0; i < 3; i++) {
for (int j = 0; j < 3; j++) {
mat[i][j] = (float)(1.0 * (white - black) / max * mat[i][j]);
}
}
HIP_API_CALL(hipMemcpyToSymbol(rgb_to_yuv_mat, mat, sizeof(mat)));
}
template<class T>
__device__ static T Clamp(T x, T lower, T upper) {
return x < lower ? lower : (x > upper ? upper : x);
}
template<class Rgb, class YuvUnit>
__device__ inline Rgb YuvToRgbForPixel(YuvUnit y, YuvUnit u, YuvUnit v) {
const int
low = 1 << (sizeof(YuvUnit) * 8 - 4),
mid = 1 << (sizeof(YuvUnit) * 8 - 1);
float fy = (int)y - low, fu = (int)u - mid, fv = (int)v - mid;
const float maxf = (1 << sizeof(YuvUnit) * 8) - 1.0f;
YuvUnit
r = (YuvUnit)Clamp(yuv_to_rgb_mat[0][0] * fy + yuv_to_rgb_mat[0][1] * fu + yuv_to_rgb_mat[0][2] * fv, 0.0f, maxf),
g = (YuvUnit)Clamp(yuv_to_rgb_mat[1][0] * fy + yuv_to_rgb_mat[1][1] * fu + yuv_to_rgb_mat[1][2] * fv, 0.0f, maxf),
b = (YuvUnit)Clamp(yuv_to_rgb_mat[2][0] * fy + yuv_to_rgb_mat[2][1] * fu + yuv_to_rgb_mat[2][2] * fv, 0.0f, maxf);
Rgb rgb{};
const int nShift = abs((int)sizeof(YuvUnit) - (int)sizeof(rgb.c.r)) * 8;
if (sizeof(YuvUnit) >= sizeof(rgb.c.r)) {
rgb.c.r = r >> nShift;
rgb.c.g = g >> nShift;
rgb.c.b = b >> nShift;
} else {
rgb.c.r = r << nShift;
rgb.c.g = g << nShift;
rgb.c.b = b << nShift;
}
return rgb;
}
// yuv to RGBA (32/64 bit)
template<class YuvUnitx2, class Rgb, class RgbIntx2>
__global__ static void YuvToRgbaKernel(uint8_t *dp_yuv, int yuv_pitch, uint8_t *dp_rgb, int rgb_pitch, int width, int height, int v_pitch) {
int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y) * 2;
if (x + 1 >= width || y + 1 >= height) {
return;
}
uint8_t *p_src = dp_yuv + x * sizeof(YuvUnitx2) / 2 + y * yuv_pitch;
uint8_t *p_dst = dp_rgb + x * sizeof(Rgb) + y * rgb_pitch;
YuvUnitx2 l0 = *(YuvUnitx2 *)p_src;
YuvUnitx2 l1 = *(YuvUnitx2 *)(p_src + yuv_pitch);
YuvUnitx2 ch = *(YuvUnitx2 *)(p_src + (v_pitch - y / 2) * yuv_pitch);
*(RgbIntx2 *)p_dst = RgbIntx2 {
YuvToRgbForPixel<Rgb>(l0.x, ch.x, ch.y).d,
YuvToRgbForPixel<Rgb>(l0.y, ch.x, ch.y).d,
};
*(RgbIntx2 *)(p_dst + rgb_pitch) = RgbIntx2 {
YuvToRgbForPixel<Rgb>(l1.x, ch.x, ch.y).d,
YuvToRgbForPixel<Rgb>(l1.y, ch.x, ch.y).d,
};
}
// yuv to RGB (24/48 bit)
template<class YuvUnitx2, class Rgb, class RgbInt1, class RgbInt2>
__global__ static void YuvToRgbKernel(uint8_t *dp_yuv, int yuv_pitch, uint8_t *dp_rgb, int rgb_pitch, int width, int height, int v_pitch) {
int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y) * 2;
if (x + 1 >= width || y + 1 >= height) {
return;
}
uint8_t *p_src = dp_yuv + x * sizeof(YuvUnitx2) / 2 + y * yuv_pitch;
uint8_t *p_dst = dp_rgb + x * sizeof(Rgb) + y * rgb_pitch;
uint8_t *p_dst1 = p_dst + rgb_pitch;
YuvUnitx2 l0 = *(YuvUnitx2 *)p_src;
YuvUnitx2 l1 = *(YuvUnitx2 *)(p_src + yuv_pitch);
YuvUnitx2 ch = *(YuvUnitx2 *)(p_src + (v_pitch - y / 2) * yuv_pitch);
Rgb rgb0 = YuvToRgbForPixel<Rgb>(l0.x, ch.x, ch.y),
rgb1 = YuvToRgbForPixel<Rgb>(l0.y, ch.x, ch.y),
rgb2 = YuvToRgbForPixel<Rgb>(l1.x, ch.x, ch.y),
rgb3 = YuvToRgbForPixel<Rgb>(l1.y, ch.x, ch.y);
*(RgbInt1 *)p_dst = RgbInt1 { rgb0.v.x, rgb0.v.y, rgb0.v.z, rgb1.v.x };
*(RgbInt2 *)(p_dst + sizeof(RgbInt1)) = RgbInt2 { rgb1.v.y, rgb1.v.z };
*(RgbInt1 *)(p_dst1) = RgbInt1 { rgb2.v.x, rgb2.v.y, rgb2.v.z, rgb3.v.x };
*(RgbInt2 *)(p_dst1 + sizeof(RgbInt1)) = RgbInt2 { rgb3.v.y, rgb3.v.z };
}
// yuv444 to RGBA (32/64 bit)
template<class YuvUnitx2, class Rgb, class RgbIntx2>
__global__ static void Yuv444ToRgbaKernel(uint8_t *dp_yuv, int yuv_pitch, uint8_t *dp_rgb, int rgb_pitch, int width, int height, int v_pitch) {
int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y);
if (x + 1 >= width || y >= height) {
return;
}
uint8_t *p_src = dp_yuv + x * sizeof(YuvUnitx2) / 2 + y * yuv_pitch;
uint8_t *p_dst = dp_rgb + x * sizeof(Rgb) + y * rgb_pitch;
YuvUnitx2 l0 = *(YuvUnitx2 *)p_src;
YuvUnitx2 ch1 = *(YuvUnitx2 *)(p_src + (v_pitch * yuv_pitch));
YuvUnitx2 ch2 = *(YuvUnitx2 *)(p_src + (2 * v_pitch * yuv_pitch));
*(RgbIntx2 *)p_dst = RgbIntx2{
YuvToRgbForPixel<Rgb>(l0.x, ch1.x, ch2.x).d,
YuvToRgbForPixel<Rgb>(l0.y, ch1.y, ch2.y).d,
};
}
// yuv444 to RGB (24/48 bit)
template<class YuvUnitx2, class Rgb, class RgbInt1, class RgbInt2>
__global__ static void Yuv444ToRgbKernel(uint8_t *dp_yuv, int yuv_pitch, uint8_t *dp_rgb, int rgb_pitch, int width, int height, int v_pitch) {
int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y);
if (x + 1 >= width || y >= height) {
return;
}
uint8_t *p_src = dp_yuv + x * sizeof(YuvUnitx2) / 2 + y * yuv_pitch;
uint8_t *p_dst = dp_rgb + x * sizeof(Rgb) + y * rgb_pitch;
YuvUnitx2 l0 = *(YuvUnitx2 *)p_src;
YuvUnitx2 ch1 = *(YuvUnitx2 *)(p_src + (v_pitch * yuv_pitch));
YuvUnitx2 ch2 = *(YuvUnitx2 *)(p_src + (2 * v_pitch * yuv_pitch));
Rgb rgb0 = YuvToRgbForPixel<Rgb>(l0.x, ch1.x, ch2.x),
rgb1 = YuvToRgbForPixel<Rgb>(l0.y, ch1.y, ch2.y);
*(RgbInt1 *)p_dst = RgbInt1 { rgb0.v.x, rgb0.v.y, rgb0.v.z, rgb1.v.x };
*(RgbInt2 *)(p_dst + sizeof(RgbInt1)) = RgbInt2 { rgb1.v.y, rgb1.v.z };
}
template<class YuvUnitx2, class Rgb, class RgbUnitx2>
__global__ static void YuvToRgbPlanarKernel(uint8_t *dp_yuv, int yuv_pitch, uint8_t *dp_rgbp, int nRgbpPitch, int width, int height, int v_pitch) {
int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y) * 2;
if (x + 1 >= width || y + 1 >= height) {
return;
}
uint8_t *p_src = dp_yuv + x * sizeof(YuvUnitx2) / 2 + y * yuv_pitch;
YuvUnitx2 l0 = *(YuvUnitx2 *)p_src;
YuvUnitx2 l1 = *(YuvUnitx2 *)(p_src + yuv_pitch);
YuvUnitx2 ch = *(YuvUnitx2 *)(p_src + (v_pitch - y / 2) * yuv_pitch);
Rgb rgb0 = YuvToRgbForPixel<Rgb>(l0.x, ch.x, ch.y),
rgb1 = YuvToRgbForPixel<Rgb>(l0.y, ch.x, ch.y),
rgb2 = YuvToRgbForPixel<Rgb>(l1.x, ch.x, ch.y),
rgb3 = YuvToRgbForPixel<Rgb>(l1.y, ch.x, ch.y);
uint8_t *p_dst = dp_rgbp + x * sizeof(RgbUnitx2) / 2 + y * nRgbpPitch;
*(RgbUnitx2 *)p_dst = RgbUnitx2 {rgb0.v.x, rgb1.v.x};
*(RgbUnitx2 *)(p_dst + nRgbpPitch) = RgbUnitx2 {rgb2.v.x, rgb3.v.x};
p_dst += nRgbpPitch * height;
*(RgbUnitx2 *)p_dst = RgbUnitx2 {rgb0.v.y, rgb1.v.y};
*(RgbUnitx2 *)(p_dst + nRgbpPitch) = RgbUnitx2 {rgb2.v.y, rgb3.v.y};
p_dst += nRgbpPitch * height;
*(RgbUnitx2 *)p_dst = RgbUnitx2 {rgb0.v.z, rgb1.v.z};
*(RgbUnitx2 *)(p_dst + nRgbpPitch) = RgbUnitx2 {rgb2.v.z, rgb3.v.z};
}
template<class YuvUnitx2, class Rgb, class RgbUnitx2>
__global__ static void Yuv444ToRgbPlanarKernel(uint8_t *dp_yuv, int yuv_pitch, uint8_t *dp_rgbp, int nRgbpPitch, int width, int height, int v_pitch) {
int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y);
if (x + 1 >= width || y >= height) {
return;
}
uint8_t *p_src = dp_yuv + x * sizeof(YuvUnitx2) / 2 + y * yuv_pitch;
YuvUnitx2 l0 = *(YuvUnitx2 *)p_src;
YuvUnitx2 ch1 = *(YuvUnitx2 *)(p_src + (v_pitch * yuv_pitch));
YuvUnitx2 ch2 = *(YuvUnitx2 *)(p_src + (2 * v_pitch * yuv_pitch));
Rgb rgb0 = YuvToRgbForPixel<Rgb>(l0.x, ch1.x, ch2.x),
rgb1 = YuvToRgbForPixel<Rgb>(l0.y, ch1.y, ch2.y);
uint8_t *p_dst = dp_rgbp + x * sizeof(RgbUnitx2) / 2 + y * nRgbpPitch;
*(RgbUnitx2 *)p_dst = RgbUnitx2{ rgb0.v.x, rgb1.v.x };
p_dst += nRgbpPitch * height;
*(RgbUnitx2 *)p_dst = RgbUnitx2{ rgb0.v.y, rgb1.v.y };
p_dst += nRgbpPitch * height;
*(RgbUnitx2 *)p_dst = RgbUnitx2{ rgb0.v.z, rgb1.v.z };
}
template <class COLOR32>
void Nv12ToColor32(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbaKernel<uchar2, COLOR32, uint2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_nv12, nv12_pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR64>
void Nv12ToColor64(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbaKernel<uchar2, COLOR64, ulonglong2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_nv12, nv12_pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR32>
void YUV444ToColor32(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbaKernel<uchar2, COLOR32, uint2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR64>
void YUV444ToColor64(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbaKernel<uchar2, COLOR64, ulonglong2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR32>
void P016ToColor32(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbaKernel<ushort2, COLOR32, uint2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_p016, p016_pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR64>
void P016ToColor64(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbaKernel<ushort2, COLOR64, ulonglong2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_p016, p016_pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR32>
void YUV444P16ToColor32(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbaKernel<ushort2, COLOR32, uint2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR64>
void YUV444P16ToColor64(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbaKernel<ushort2, COLOR64, ulonglong2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR32>
void Nv12ToColorPlanar(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbPlanarKernel<uchar2, COLOR32, uchar2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_nv12, nv12_pitch, dp_bgrp, nBgrpPitch, width, height, v_pitch);
}
template <class COLOR32>
void P016ToColorPlanar(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbPlanarKernel<ushort2, COLOR32, uchar2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_p016, p016_pitch, dp_bgrp, nBgrpPitch, width, height, v_pitch);
}
template <class COLOR32>
void YUV444ToColorPlanar(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbPlanarKernel<uchar2, COLOR32, uchar2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgrp, nBgrpPitch, width, height, v_pitch);
}
template <class COLOR32>
void YUV444P16ToColorPlanar(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbPlanarKernel<ushort2, COLOR32, uchar2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgrp, nBgrpPitch, width, height, v_pitch);
}
// Explicit Instantiation: for RGB32/BGR32 and RGB64/BGR64 formats
template void Nv12ToColor32<BGRA32>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void Nv12ToColor32<RGBA32>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void Nv12ToColor64<BGRA64>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void Nv12ToColor64<RGBA64>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColor32<BGRA32>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColor32<RGBA32>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColor64<BGRA64>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColor64<RGBA64>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColor32<BGRA32>(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColor32<RGBA32>(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColor64<BGRA64>(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColor64<RGBA64>(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColor32<BGRA32>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColor32<RGBA32>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColor64<BGRA64>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColor64<RGBA64>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void Nv12ToColorPlanar<BGRA32>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void Nv12ToColorPlanar<RGBA32>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColorPlanar<BGRA32>(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColorPlanar<RGBA32>(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColorPlanar<BGRA32>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColorPlanar<RGBA32>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColorPlanar<BGRA32>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColorPlanar<RGBA32>(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgrp, int nBgrpPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR24>
void Nv12ToColor24(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbKernel<uchar2, COLOR24, uchar4, uchar2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_nv12, nv12_pitch, dp_bgr, bgr_pitch, width, height, v_pitch);
}
template <class COLOR48>
void Nv12ToColor48(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbKernel<uchar2, COLOR48, ushort4, ushort2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_nv12, nv12_pitch, dp_bgr, bgr_pitch, width, height, v_pitch);
}
template <class COLOR24>
void YUV444ToColor24(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbKernel<uchar2, COLOR24, uchar4, uchar2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgr, bgr_pitch, width, height, v_pitch);
}
template <class COLOR48>
void YUV444ToColor48(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbKernel<uchar2, COLOR48, ushort4, ushort2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgr, bgr_pitch, width, height, v_pitch);
}
template <class COLOR24>
void P016ToColor24(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbKernel<ushort2, COLOR24, uchar4, uchar2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_p016, p016_pitch, dp_bgr, bgr_pitch, width, height, v_pitch);
}
template <class COLOR48>
void P016ToColor48(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
YuvToRgbKernel<ushort2, COLOR48, ushort4, ushort2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_p016, p016_pitch, dp_bgr, bgr_pitch, width, height, v_pitch);
}
template <class COLOR24>
void YUV444P16ToColor24(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbKernel<ushort2, COLOR24, uchar4, uchar2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch);
}
template <class COLOR48>
void YUV444P16ToColor48(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream) {
SetMatYuv2Rgb(col_standard);
Yuv444ToRgbKernel<ushort2, COLOR48, ushort4, ushort2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2), dim3(32, 2), 0, hip_stream>>>
(dp_yuv_444, pitch, dp_bgr, bgr_pitch, width, height, v_pitch);
}
// Explicit Instantiation: for RGB24/BGR24 and RGB48/BGR48 formats
template void Nv12ToColor24<BGR24>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void Nv12ToColor24<RGB24>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void Nv12ToColor48<BGR48>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void Nv12ToColor48<RGB48>(uint8_t *dp_nv12, int nv12_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColor24<BGR24>(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColor24<RGB24>(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColor48<BGR48>(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444ToColor48<RGB48>(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColor24<BGR24>(uint8_t *dp_p016, int p016_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColor24<RGB24>(uint8_t *dp_p016, int p016_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColor48<BGR48>(uint8_t *dp_p016, int p016_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void P016ToColor48<RGB48>(uint8_t *dp_p016, int p016_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColor24<BGR24>(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColor24<RGB24>(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColor48<BGR48>(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template void YUV444P16ToColor48<RGB48>(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template<class YuvUnit, class RgbUnit>
__device__ inline YuvUnit RgbToY(RgbUnit r, RgbUnit g, RgbUnit b) {
const YuvUnit low = 1 << (sizeof(YuvUnit) * 8 - 4);
return rgb_to_yuv_mat[0][0] * r + rgb_to_yuv_mat[0][1] * g + rgb_to_yuv_mat[0][2] * b + low;
}
template<class YuvUnit, class RgbUnit>
__device__ inline YuvUnit RgbToU(RgbUnit r, RgbUnit g, RgbUnit b) {
const YuvUnit mid = 1 << (sizeof(YuvUnit) * 8 - 1);
return rgb_to_yuv_mat[1][0] * r + rgb_to_yuv_mat[1][1] * g + rgb_to_yuv_mat[1][2] * b + mid;
}
template<class YuvUnit, class RgbUnit>
__device__ inline YuvUnit RgbToV(RgbUnit r, RgbUnit g, RgbUnit b) {
const YuvUnit mid = 1 << (sizeof(YuvUnit) * 8 - 1);
return rgb_to_yuv_mat[2][0] * r + rgb_to_yuv_mat[2][1] * g + rgb_to_yuv_mat[2][2] * b + mid;
}
template<class YuvUnitx2, class Rgb, class RgbIntx2>
__global__ static void RgbaToYuvKernel(uint8_t *dp_rgb, int rgba_pitch, uint8_t *dp_yuv, int yuv_pitch, int width, int height) {
int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y) * 2;
if (x + 1 >= width || y + 1 >= height) {
return;
}
uint8_t *p_src = dp_rgb + x * sizeof(Rgb) + y * rgba_pitch;
RgbIntx2 int2a = *(RgbIntx2 *)p_src;
RgbIntx2 int2b = *(RgbIntx2 *)(p_src + rgba_pitch);
Rgb rgb[4] = {int2a.x, int2a.y, int2b.x, int2b.y};
decltype(Rgb::c.r)
r = (rgb[0].c.r + rgb[1].c.r + rgb[2].c.r + rgb[3].c.r) / 4,
g = (rgb[0].c.g + rgb[1].c.g + rgb[2].c.g + rgb[3].c.g) / 4,
b = (rgb[0].c.b + rgb[1].c.b + rgb[2].c.b + rgb[3].c.b) / 4;
uint8_t *p_dst = dp_yuv + x * sizeof(YuvUnitx2) / 2 + y * yuv_pitch;
*(YuvUnitx2 *)p_dst = YuvUnitx2 {
RgbToY<decltype(YuvUnitx2::x)>(rgb[0].c.r, rgb[0].c.g, rgb[0].c.b),
RgbToY<decltype(YuvUnitx2::x)>(rgb[1].c.r, rgb[1].c.g, rgb[1].c.b),
};
*(YuvUnitx2 *)(p_dst + yuv_pitch) = YuvUnitx2 {
RgbToY<decltype(YuvUnitx2::x)>(rgb[2].c.r, rgb[2].c.g, rgb[2].c.b),
RgbToY<decltype(YuvUnitx2::x)>(rgb[3].c.r, rgb[3].c.g, rgb[3].c.b),
};
*(YuvUnitx2 *)(p_dst + (height - y / 2) * yuv_pitch) = YuvUnitx2 {
RgbToU<decltype(YuvUnitx2::x)>(r, g, b),
RgbToV<decltype(YuvUnitx2::x)>(r, g, b),
};
}
void Bgra64ToP016(uint8_t *dp_bgra, int bgra_pitch, uint8_t *dp_p016, int p016_pitch, int width, int height, int col_standard, hipStream_t hip_stream) {
SetMatRgb2Yuv(col_standard);
RgbaToYuvKernel<ushort2, BGRA64, ulonglong2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(dp_bgra, bgra_pitch, dp_p016, p016_pitch, width, height);
}
template<class YuvUnitx2, class Rgb, class RgbInt1, class RgbInt2>
__global__ static void RgbToYuvKernel(uint8_t *dp_rgb, int rgb_pitch, uint8_t *dp_yuv, int yuv_pitch, int width, int height) {
int x = (threadIdx.x + blockIdx.x * blockDim.x) * 2;
int y = (threadIdx.y + blockIdx.y * blockDim.y) * 2;
if (x + 1 >= width || y + 1 >= height) {
return;
}
uint8_t *p_src = dp_rgb + x * sizeof(Rgb) + y * rgb_pitch;
RgbInt1 int1a = *(RgbInt1 *)p_src;
RgbInt2 int2a = *(RgbInt2 *)(p_src + sizeof(RgbInt1));
RgbInt1 int1b = *(RgbInt1 *)(p_src + rgb_pitch);
RgbInt2 int2b = *(RgbInt2 *)(p_src + rgb_pitch + sizeof(RgbInt1));
Rgb rgb[4];
rgb[0].v = {int1a.x, int1a.y, int1a.z},
rgb[1].v = {int1a.w, int2a.x, int2a.y},
rgb[2].v = {int1b.x, int1b.y, int1b.z},
rgb[3].v = {int1b.w, int2b.x, int2b.y};
decltype(Rgb::c.r)
r = (rgb[0].c.r + rgb[1].c.r + rgb[2].c.r + rgb[3].c.r) / 4,
g = (rgb[0].c.g + rgb[1].c.g + rgb[2].c.g + rgb[3].c.g) / 4,
b = (rgb[0].c.b + rgb[1].c.b + rgb[2].c.b + rgb[3].c.b) / 4;
uint8_t *p_dst = dp_yuv + x * sizeof(YuvUnitx2) / 2 + y * yuv_pitch;
*(YuvUnitx2 *)p_dst = YuvUnitx2 {
RgbToY<decltype(YuvUnitx2::x)>(rgb[0].c.r, rgb[0].c.g, rgb[0].c.b),
RgbToY<decltype(YuvUnitx2::x)>(rgb[1].c.r, rgb[1].c.g, rgb[1].c.b),
};
*(YuvUnitx2 *)(p_dst + yuv_pitch) = YuvUnitx2 {
RgbToY<decltype(YuvUnitx2::x)>(rgb[2].c.r, rgb[2].c.g, rgb[2].c.b),
RgbToY<decltype(YuvUnitx2::x)>(rgb[3].c.r, rgb[3].c.g, rgb[3].c.b),
};
*(YuvUnitx2 *)(p_dst + (height - y / 2) * yuv_pitch) = YuvUnitx2 {
RgbToU<decltype(YuvUnitx2::x)>(r, g, b),
RgbToV<decltype(YuvUnitx2::x)>(r, g, b),
};
}
void Bgr48ToP016(uint8_t *p_bgr, int bgr_pitch, uint8_t *dp_p016, int p016_pitch, int width, int height, int col_standard, hipStream_t hip_stream) {
SetMatRgb2Yuv(col_standard);
RgbToYuvKernel<ushort2, BGR48, ushort4, ushort2>
<<<dim3((width + 63) / 32 / 2, (height + 3) / 2 / 2), dim3(32, 2), 0, hip_stream>>>
(p_bgr, bgr_pitch, dp_p016, p016_pitch, width, height);
}
projects/rocdecode/utils/colorspace_kernels.h
+143
Wyświetl plik
@@ -0,0 +1,143 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include <stdint.h>
#include <hip/hip_runtime.h>
/*!
* \file
* \brief The AMD Color Space Standards for VCN Decode Library.
*
* \defgroup group_amd_vcn_colorspace colorSpace: AMD VCN Color Space API
* \brief AMD The vcnDECODE Color Space API.
*/
typedef enum ColorSpaceStandard_ {
ColorSpaceStandard_BT709 = 1,
ColorSpaceStandard_Unspecified = 2,
ColorSpaceStandard_Reserved = 3,
ColorSpaceStandard_FCC = 4,
ColorSpaceStandard_BT470 = 5,
ColorSpaceStandard_BT601 = 6,
ColorSpaceStandard_SMPTE240M = 7,
ColorSpaceStandard_YCgCo = 8,
ColorSpaceStandard_BT2020 = 9,
ColorSpaceStandard_BT2020C = 10
} ColorSpaceStandard;
union BGR24 {
uchar3 v;
struct {
uint8_t b, g, r;
} c;
};
union RGB24 {
uchar3 v;
struct {
uint8_t r, g, b;
} c;
};
union BGR48 {
ushort3 v;
struct {
uint16_t b, g, r;
} c;
};
union RGB48 {
ushort3 v;
struct {
uint16_t r, g, b;
} c;
};
union BGRA32 {
uint32_t d;
uchar4 v;
struct {
uint8_t b, g, r, a;
} c;
};
union RGBA32 {
uint32_t d;
uchar4 v;
struct {
uint8_t r, g, b, a;
} c;
};
union BGRA64 {
uint64_t d;
ushort4 v;
struct {
uint16_t b, g, r, a;
} c;
};
union RGBA64 {
uint64_t d;
ushort4 v;
struct {
uint16_t r, g, b, a;
} c;
};
// color-convert hip kernel function definitions
template <class COLOR32>
void YUV444ToColor32(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR64>
void YUV444ToColor64(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR24>
void YUV444ToColor24(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR48>
void YUV444ToColor48(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR24>
void Nv12ToColor24(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR32>
void Nv12ToColor32(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR48>
void Nv12ToColor48(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR64>
void Nv12ToColor64(uint8_t *dp_nv12, int nv12_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR24>
void YUV444P16ToColor24(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR48>
void YUV444P16ToColor48(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR32>
void YUV444P16ToColor32(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR64>
void YUV444P16ToColor64(uint8_t *dp_yuv_444, int pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR32>
void P016ToColor32(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR64>
void P016ToColor64(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgra, int bgra_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR24>
void P016ToColor24(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
template <class COLOR48>
void P016ToColor48(uint8_t *dp_p016, int p016_pitch, uint8_t *dp_bgr, int bgr_pitch, int width, int height, int v_pitch, int col_standard, hipStream_t hip_stream);
projects/rocdecode/utils/ffmpegvideodecode/ffmpeg_video_dec.cpp
+628
Wyświetl plik
@@ -0,0 +1,628 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "ffmpeg_video_dec.h"
//helper function
static inline float GetChromaWidthFactor(rocDecVideoSurfaceFormat surface_format) {
float factor = 0.5;
switch (surface_format) {
case rocDecVideoSurfaceFormat_NV12:
case rocDecVideoSurfaceFormat_P016:
case rocDecVideoSurfaceFormat_YUV444:
case rocDecVideoSurfaceFormat_YUV444_16Bit:
factor = 1.0;
break;
case rocDecVideoSurfaceFormat_YUV420:
case rocDecVideoSurfaceFormat_YUV420_16Bit:
case rocDecVideoSurfaceFormat_YUV422:
case rocDecVideoSurfaceFormat_YUV422_16Bit:
factor = 0.5;
break;
}
return factor;
};
FFMpegVideoDecoder::FFMpegVideoDecoder(int device_id, OutputSurfaceMemoryType out_mem_type, rocDecVideoCodec codec, bool force_zero_latency,
const Rect *p_crop_rect, bool extract_user_sei_Message, uint32_t disp_delay, int max_width, int max_height, uint32_t clk_rate) :
RocVideoDecoder(device_id, out_mem_type, codec, force_zero_latency, p_crop_rect, extract_user_sei_Message, disp_delay, max_width, max_height, clk_rate, true) {
if ((out_mem_type_ == OUT_SURFACE_MEM_DEV_INTERNAL) || (out_mem_type_ == OUT_SURFACE_MEM_NOT_MAPPED)) {
ROCDEC_THROW("Unsupported output memory type", ROCDEC_INVALID_PARAMETER);
}
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
if (!InitHIP(device_id_)) {
ROCDEC_THROW("Failed to initilize the HIP", ROCDEC_DEVICE_INVALID);
}
}
// many of the decoder parameters are hardcoded below for just creating the decoder.
// In the handlevideosequence callback, the decoder will get reconfigured to the actual parameters in the sequence header
RocDecoderHostCreateInfo create_info = {};
create_info.codec_type = codec;
create_info.num_decode_threads = 0; // default
create_info.max_width = max_width;
create_info.max_height = max_height;
create_info.width = max_width;
create_info.height = max_height;
create_info.target_width = max_width;
create_info.target_height = max_height;
create_info.display_rect.left = 0;
create_info.display_rect.right = static_cast<short>(max_width);
create_info.display_rect.top = 0;
create_info.display_rect.bottom = static_cast<short>(max_height);
create_info.chroma_format = rocDecVideoChromaFormat_420;
create_info.output_format = rocDecVideoSurfaceFormat_P016;
create_info.bit_depth_minus_8 = 2;
create_info.num_output_surfaces = 1;
create_info.user_data = this;
create_info.pfn_sequence_callback = FFMpegHandleVideoSequenceProc;
create_info.pfn_display_picture = FFMpegHandlePictureDisplayProc;
create_info.pfn_get_sei_msg = nullptr; // tobe supported in future
ROCDEC_API_CALL(rocDecCreateDecoderHost(&roc_decoder_, &create_info));
// set disp_width and height to non_zero values for it doesn't trigger decoding error before actual start of decoding
disp_width_ = max_width;
disp_height_ = max_height;
// fill output_surface_info_
output_surface_info_.output_width = max_width;
output_surface_info_.output_height = max_height;
output_surface_info_.output_pitch = max_width * 2; // bytes_per_pixel 2
output_surface_info_.output_vstride = max_height;
output_surface_info_.bit_depth = bitdepth_minus_8_ + 8;
output_surface_info_.bytes_per_pixel = 2;
output_surface_info_.surface_format = rocDecVideoSurfaceFormat_P016;
output_surface_info_.num_chroma_planes = 2;
output_surface_info_.mem_type = OUT_SURFACE_MEM_HOST_COPIED;
}
FFMpegVideoDecoder::~FFMpegVideoDecoder() {
std::lock_guard<std::mutex> lock(mtx_vp_frame_);
for (auto &p_frame : vp_frames_) {
if (p_frame.frame_ptr) {
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
hipError_t hip_status = hipFree(p_frame.frame_ptr);
if (hip_status != hipSuccess) {
std::cerr << "ERROR: hipFree failed! (" << hip_status << ")" << std::endl;
}
}
}
}
}
/* Return value from HandleVideoSequence() are interpreted as :
* 0: fail, 1: succeeded, > 1: override dpb size of parser (set by RocdecParserParams::max_num_decode_surfaces while creating parser)
*/
int FFMpegVideoDecoder::HandleVideoSequence(RocdecVideoFormatHost *format_host) {
if (format_host == nullptr) {
ROCDEC_THROW("Rocdec:: Invalid video format in HandleVideoSequence: ", ROCDEC_INVALID_PARAMETER);
return 0;
}
auto start_time = StartTimer();
RocdecVideoFormat *p_video_format = &format_host->video_format;
input_video_info_str_.str("");
input_video_info_str_.clear();
input_video_info_str_ << "Input Video Information" << std::endl
<< "\tCodec : " << GetCodecFmtName(p_video_format->codec) << std::endl;
if (p_video_format->frame_rate.numerator && p_video_format->frame_rate.denominator) {
input_video_info_str_ << "\tFrame rate : " << p_video_format->frame_rate.numerator << "/" << p_video_format->frame_rate.denominator << " = " << 1.0 * p_video_format->frame_rate.numerator / p_video_format->frame_rate.denominator << " fps" << std::endl;
}
input_video_info_str_ << "\tSequence : " << (p_video_format->progressive_sequence ? "Progressive" : "Interlaced") << std::endl
<< "\tCoded size : [" << p_video_format->coded_width << ", " << p_video_format->coded_height << "]" << std::endl
<< "\tDisplay area : [" << p_video_format->display_area.left << ", " << p_video_format->display_area.top << ", "
<< p_video_format->display_area.right << ", " << p_video_format->display_area.bottom << "]" << std::endl
<< "\tBit depth : " << p_video_format->bit_depth_luma_minus8 + 8
;
input_video_info_str_ << std::endl;
int num_decode_surfaces = p_video_format->min_num_decode_surfaces;
if (curr_video_format_ptr_ == nullptr) {
curr_video_format_ptr_ = new RocdecVideoFormat();
}
// store current video format: this is required to call reconfigure from application in case of random seek
if (curr_video_format_ptr_) memcpy(curr_video_format_ptr_, p_video_format, sizeof(RocdecVideoFormat));
if (coded_width_ && coded_height_) {
// rocdecCreateDecoder() has been called before, and now there's possible config change
return ReconfigureDecoder(p_video_format);
}
// e_codec has been set in the constructor (for parser). Here it's set again for potential correction
codec_id_ = p_video_format->codec;
video_chroma_format_ = p_video_format->chroma_format;
bitdepth_minus_8_ = p_video_format->bit_depth_luma_minus8;
byte_per_pixel_ = bitdepth_minus_8_ > 0 ? 2 : 1;
// convert AVPixelFormat to rocDecVideoChromaFormat
video_surface_format_ = format_host->video_surface_format;
coded_width_ = p_video_format->coded_width;
coded_height_ = p_video_format->coded_height;
disp_rect_.top = p_video_format->display_area.top;
disp_rect_.bottom = p_video_format->display_area.bottom;
disp_rect_.left = p_video_format->display_area.left;
disp_rect_.right = p_video_format->display_area.right;
disp_width_ = p_video_format->display_area.right - p_video_format->display_area.left;
disp_height_ = p_video_format->display_area.bottom - p_video_format->display_area.top;
// AV1 has max width/height of sequence in sequence header
if (codec_id_ == rocDecVideoCodec_AV1 && p_video_format->seqhdr_data_length > 0) {
// dont overwrite if it is already set from cmdline or reconfig.txt
if (!(max_width_ > p_video_format->coded_width || max_height_ > p_video_format->coded_height)) {
RocdecVideoFormatEx *vidFormatEx = (RocdecVideoFormatEx *)p_video_format;
max_width_ = vidFormatEx->max_width;
max_height_ = vidFormatEx->max_height;
}
}
if (max_width_ < static_cast<int>(p_video_format->coded_width))
max_width_ = p_video_format->coded_width;
if (max_height_ < static_cast<int>(p_video_format->coded_height))
max_height_ = p_video_format->coded_height;
if (!(crop_rect_.right && crop_rect_.bottom)) {
target_width_ = (disp_width_ + 1) & ~1;
target_height_ = (disp_height_ + 1) & ~1;
} else {
target_width_ = (crop_rect_.right - crop_rect_.left + 1) & ~1;
target_height_ = (crop_rect_.bottom - crop_rect_.top + 1) & ~1;
}
chroma_height_ = static_cast<int>(ceil(target_height_ * GetChromaHeightFactor(video_surface_format_)));
chroma_width_ = static_cast<int>(ceil(target_width_ * GetChromaWidthFactor(video_surface_format_)));
num_chroma_planes_ = GetChromaPlaneCount(video_surface_format_);
if (video_chroma_format_ == rocDecVideoChromaFormat_Monochrome) num_chroma_planes_ = 0;
surface_stride_ = target_width_ * byte_per_pixel_;
// fill output_surface_info_
output_surface_info_.output_width = target_width_;
output_surface_info_.output_height = target_height_;
output_surface_info_.output_pitch = surface_stride_;
output_surface_info_.output_vstride = target_height_;
output_surface_info_.bit_depth = bitdepth_minus_8_ + 8;
output_surface_info_.bytes_per_pixel = byte_per_pixel_;
output_surface_info_.surface_format = video_surface_format_;
output_surface_info_.num_chroma_planes = num_chroma_planes_;
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
output_surface_info_.output_surface_size_in_bytes = GetFrameSize();
output_surface_info_.mem_type = OUT_SURFACE_MEM_DEV_COPIED;
} else if (out_mem_type_ == OUT_SURFACE_MEM_HOST_COPIED){
output_surface_info_.output_surface_size_in_bytes = GetFrameSize();
output_surface_info_.mem_type = OUT_SURFACE_MEM_HOST_COPIED;
}
input_video_info_str_ << "Video Decoding Params:" << std::endl
<< "\tNum Surfaces : " << num_decode_surfaces << std::endl
<< "\tCrop : [" << disp_rect_.left << ", " << disp_rect_.top << ", "
<< disp_rect_.right << ", " << disp_rect_.bottom << "]" << std::endl
<< "\tResize : " << target_width_ << "x" << target_height_ << std::endl
;
input_video_info_str_ << std::endl;
std::cout << input_video_info_str_.str();
double elapsed_time = StopTimer(start_time);
AddDecoderSessionOverHead(std::this_thread::get_id(), elapsed_time);
return num_decode_surfaces;
}
bool FFMpegVideoDecoder::GetOutputSurfaceInfo(OutputSurfaceInfo **surface_info) {
if (!disp_width_ || !disp_height_) {
std::cerr << "ERROR: FFMpegVideo is not initialized" << std::endl;
return false;
}
*surface_info = &output_surface_info_;
return true;
}
/**
* @brief function to reconfigure decoder if there is a change in sequence params.
*
* @param p_video_format
* @return int 1: success 0: fail
*/
int FFMpegVideoDecoder::ReconfigureDecoder(RocdecVideoFormat *p_video_format) {
if (p_video_format->codec != codec_id_) {
ROCDEC_THROW("Reconfigure Not supported for codec change", ROCDEC_NOT_SUPPORTED);
return 0;
}
if (p_video_format->chroma_format != video_chroma_format_) {
ROCDEC_THROW("Reconfigure Not supported for chroma format change", ROCDEC_NOT_SUPPORTED);
return 0;
}
if (p_video_format->bit_depth_luma_minus8 != bitdepth_minus_8_){
ROCDEC_THROW("Reconfigure Not supported for bit depth change", ROCDEC_NOT_SUPPORTED);
return 0;
}
bool is_decode_res_changed = !(p_video_format->coded_width == coded_width_ && p_video_format->coded_height == coded_height_);
bool is_display_rect_changed = !(p_video_format->display_area.bottom == disp_rect_.bottom &&
p_video_format->display_area.top == disp_rect_.top &&
p_video_format->display_area.left == disp_rect_.left &&
p_video_format->display_area.right == disp_rect_.right);
if (!is_decode_res_changed && !is_display_rect_changed && !b_force_recofig_flush_) {
return 1;
}
// Flush and clear internal frame store to reconfigure when either coded size or display size has changed.
if (p_reconfig_params_ && p_reconfig_params_->p_fn_reconfigure_flush)
num_frames_flushed_during_reconfig_ += p_reconfig_params_->p_fn_reconfigure_flush(this, p_reconfig_params_->reconfig_flush_mode, static_cast<void *>(p_reconfig_params_->p_reconfig_user_struct));
// clear the existing output buffers of different size
// note that app lose the remaining frames in the vp_frames in case application didn't set p_fn_reconfigure_flush_ callback
std::lock_guard<std::mutex> lock(mtx_vp_frame_);
while(!vp_frames_.empty()) {
DecFrameBuffer *p_frame = &vp_frames_.back();
// pop decoded frame
vp_frames_.pop_back();
if (p_frame->frame_ptr) {
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
hipError_t hip_status = hipFree(p_frame->frame_ptr);
if (hip_status != hipSuccess) std::cerr << "ERROR: hipFree failed! (" << hip_status << ")" << std::endl;
}
}
}
output_frame_cnt_ = 0; // reset frame_count
if (is_decode_res_changed) {
coded_width_ = p_video_format->coded_width;
coded_height_ = p_video_format->coded_height;
}
if (is_display_rect_changed) {
disp_rect_.left = p_video_format->display_area.left;
disp_rect_.right = p_video_format->display_area.right;
disp_rect_.top = p_video_format->display_area.top;
disp_rect_.bottom = p_video_format->display_area.bottom;
disp_width_ = p_video_format->display_area.right - p_video_format->display_area.left;
disp_height_ = p_video_format->display_area.bottom - p_video_format->display_area.top;
if (!(crop_rect_.right && crop_rect_.bottom)) {
target_width_ = (disp_width_ + 1) & ~1;
target_height_ = (disp_height_ + 1) & ~1;
} else {
target_width_ = (crop_rect_.right - crop_rect_.left + 1) & ~1;
target_height_ = (crop_rect_.bottom - crop_rect_.top + 1) & ~1;
}
is_output_surface_changed_ = true;
}
surface_stride_ = target_width_ * byte_per_pixel_;
chroma_height_ = static_cast<int>(std::ceil(target_height_ * GetChromaHeightFactor(video_surface_format_)));
chroma_width_ = static_cast<int>(ceil(target_width_ * GetChromaWidthFactor(video_surface_format_)));
num_chroma_planes_ = GetChromaPlaneCount(video_surface_format_);
if (p_video_format->chroma_format == rocDecVideoChromaFormat_Monochrome) num_chroma_planes_ = 0;
// Fill output_surface_info_
output_surface_info_.output_width = target_width_;
output_surface_info_.output_height = target_height_;
output_surface_info_.output_pitch = surface_stride_;
output_surface_info_.output_vstride = (out_mem_type_ == OUT_SURFACE_MEM_DEV_INTERNAL) ? surface_vstride_ : target_height_;
output_surface_info_.bit_depth = bitdepth_minus_8_ + 8;
output_surface_info_.bytes_per_pixel = byte_per_pixel_;
output_surface_info_.surface_format = video_surface_format_;
output_surface_info_.num_chroma_planes = num_chroma_planes_;
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
output_surface_info_.output_surface_size_in_bytes = GetFrameSize();
output_surface_info_.mem_type = OUT_SURFACE_MEM_DEV_COPIED;
} else if (out_mem_type_ == OUT_SURFACE_MEM_HOST_COPIED) {
output_surface_info_.output_surface_size_in_bytes = GetFrameSize();
output_surface_info_.mem_type = OUT_SURFACE_MEM_HOST_COPIED;
}
// If the coded_width or coded_height hasn't changed but display resolution has changed, then need to update width and height for
// correct output with cropping. There is no need to reconfigure the decoder.
if (!is_decode_res_changed && is_display_rect_changed) {
return 1;
}
input_video_info_str_.str("");
input_video_info_str_.clear();
input_video_info_str_ << "Input Video Resolution Changed:" << std::endl
<< "\tCoded size : [" << p_video_format->coded_width << ", " << p_video_format->coded_height << "]" << std::endl
<< "\tDisplay area : [" << p_video_format->display_area.left << ", " << p_video_format->display_area.top << ", "
<< p_video_format->display_area.right << ", " << p_video_format->display_area.bottom << "]" << std::endl;
input_video_info_str_ << std::endl;
return 1;
}
/**
* @brief function to handle display picture
*
* @param pDispInfo
* @return int 0:fail 1: success
*/
int FFMpegVideoDecoder::HandlePictureDisplay(RocdecParserDispInfo *pDispInfo) {
if (b_extract_sei_message_) {
if (sei_message_display_q_[pDispInfo->picture_index].sei_data) {
// Write SEI Message
uint8_t *sei_buffer = static_cast<uint8_t *>(sei_message_display_q_[pDispInfo->picture_index].sei_data);
uint32_t sei_num_messages = sei_message_display_q_[pDispInfo->picture_index].sei_message_count;
RocdecSeiMessage *sei_message = sei_message_display_q_[pDispInfo->picture_index].sei_message;
if (fp_sei_) {
for (uint32_t i = 0; i < sei_num_messages; i++) {
if (codec_id_ == rocDecVideoCodec_AVC || rocDecVideoCodec_HEVC) {
switch (sei_message[i].sei_message_type) {
case SEI_TYPE_TIME_CODE: {
//todo:: check if we need to write timecode
}
break;
case SEI_TYPE_USER_DATA_UNREGISTERED: {
fwrite(sei_buffer, sei_message[i].sei_message_size, 1, fp_sei_);
}
break;
}
}
if (codec_id_ == rocDecVideoCodec_AV1) {
fwrite(sei_buffer, sei_message[i].sei_message_size, 1, fp_sei_);
}
sei_buffer += sei_message[i].sei_message_size;
}
}
free(sei_message_display_q_[pDispInfo->picture_index].sei_data);
sei_message_display_q_[pDispInfo->picture_index].sei_data = NULL; // to avoid double free
free(sei_message_display_q_[pDispInfo->picture_index].sei_message);
sei_message_display_q_[pDispInfo->picture_index].sei_message = NULL; // to avoid double free
}
}
RocdecParserDispInfo *p_disp_info = static_cast<RocdecParserDispInfo *>(pDispInfo);
RocdecProcParams video_proc_params = {};
video_proc_params.progressive_frame = p_disp_info->progressive_frame;
video_proc_params.top_field_first = p_disp_info->top_field_first;
void * src_ptr[3] = { 0 };
uint32_t src_pitch[3] = { 0 };
ROCDEC_API_CALL(rocDecGetVideoFrameHost(roc_decoder_, pDispInfo->picture_index, src_ptr, src_pitch, &video_proc_params));
// copy the decoded surface info device or host
uint8_t *p_dec_frame = nullptr;
{
std::lock_guard<std::mutex> lock(mtx_vp_frame_);
// if not enough frames in stock, allocate
if (++output_frame_cnt_ > vp_frames_.size()) {
num_alloced_frames_++;
DecFrameBuffer dec_frame = {0};
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
// allocate device memory
HIP_API_CALL(hipMalloc((void **)&dec_frame.frame_ptr, GetFrameSize()));
} else {
dec_frame.frame_ptr = new uint8_t[GetFrameSize()];
}
dec_frame.pts = pDispInfo->pts;
dec_frame.picture_index = pDispInfo->picture_index; //picture_index is not used here since it is handled within FFMpeg decoder
vp_frames_.push_back(dec_frame);
}
p_dec_frame = vp_frames_[output_frame_cnt_ - 1].frame_ptr;
}
// Copy luma data
int dst_pitch = disp_width_ * byte_per_pixel_;
uint8_t *p_src_ptr_y = static_cast<uint8_t *>(src_ptr[0]) + (disp_rect_.top + crop_rect_.top) * src_pitch[0] + (disp_rect_.left + crop_rect_.left) * byte_per_pixel_;
uint8_t *p_frame_y = p_dec_frame;
if (!p_frame_y && !p_src_ptr_y) {
std::cerr << "HandlePictureDisplay: Invalid Memory address for src/dst" << std::endl;
return 0;
}
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
if (src_pitch[0] == dst_pitch) {
int luma_size = src_pitch[0] * disp_height_;
HIP_API_CALL(hipMemcpyHtoDAsync(p_frame_y, p_src_ptr_y, luma_size, hip_stream_));
} else {
// use 2d copy to copy an ROI
HIP_API_CALL(hipMemcpy2DAsync(p_frame_y, dst_pitch, p_src_ptr_y, src_pitch[0], dst_pitch, disp_height_, hipMemcpyHostToDevice, hip_stream_));
}
} else {
if (src_pitch[0] == dst_pitch) {
int luma_size = src_pitch[0] * disp_height_;
memcpy(p_frame_y, p_src_ptr_y, luma_size);
} else {
for (int i = 0; i < disp_height_; i++) {
memcpy(p_dec_frame, p_src_ptr_y, dst_pitch);
p_frame_y += dst_pitch;
p_src_ptr_y += src_pitch[0];
}
}
}
// Copy chroma plane/s
// rocDec output gives pointer to luma and chroma pointers seperated for the decoded frame
uint8_t *p_frame_uv = p_dec_frame + dst_pitch * disp_height_;
uint8_t *p_src_ptr_uv = static_cast<uint8_t *>(src_ptr[1]) + ((disp_rect_.top + crop_rect_.top) >> 1) * src_pitch[1] + ((disp_rect_.left + crop_rect_.left)>>1) * byte_per_pixel_ ;
dst_pitch = chroma_width_ * byte_per_pixel_;
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
if (src_pitch[1] == dst_pitch) {
int chroma_size = chroma_height_ * dst_pitch;
HIP_API_CALL(hipMemcpyHtoDAsync(p_frame_uv, p_src_ptr_uv, chroma_size, hip_stream_));
} else {
// use 2d copy to copy an ROI
HIP_API_CALL(hipMemcpy2DAsync(p_frame_uv, dst_pitch, p_src_ptr_uv, src_pitch[1], dst_pitch, chroma_height_, hipMemcpyHostToDevice, hip_stream_));
}
} else {
if (src_pitch[1] == dst_pitch) {
int chroma_size = chroma_height_ * dst_pitch;
memcpy(p_frame_uv, p_src_ptr_uv, chroma_size);
}
else {
for (int i = 0; i < chroma_height_; i++) {
memcpy(p_frame_uv, p_src_ptr_uv, dst_pitch);
p_frame_uv += dst_pitch;
p_src_ptr_uv += src_pitch[1];
}
}
}
if (num_chroma_planes_ == 2) {
uint8_t *p_frame_v = p_frame_uv + dst_pitch * chroma_height_;
uint8_t *p_src_ptr_v = static_cast<uint8_t *>(src_ptr[2]) + (disp_rect_.top + crop_rect_.top) * src_pitch[2] + ((disp_rect_.left + crop_rect_.left) >> 1) * byte_per_pixel_;
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) {
if (src_pitch[2] == dst_pitch) {
int chroma_size = chroma_height_ * dst_pitch;
HIP_API_CALL(hipMemcpyDtoDAsync(p_frame_v, p_src_ptr_v, chroma_size, hip_stream_));
} else {
// use 2d copy to copy an ROI
HIP_API_CALL(hipMemcpy2DAsync(p_frame_v, dst_pitch, p_src_ptr_v, src_pitch[2], dst_pitch, chroma_height_, hipMemcpyHostToDevice, hip_stream_));
}
}
else {
if (src_pitch[2] == dst_pitch) {
int chroma_size = chroma_height_ * dst_pitch;
memcpy(p_frame_v, p_src_ptr_v, chroma_size);
}
else {
for (int i = 0; i < chroma_height_; i++) {
memcpy(p_frame_v, p_src_ptr_v, dst_pitch);
p_frame_v += dst_pitch;
p_src_ptr_v += src_pitch[1];
}
}
}
}
if (out_mem_type_ == OUT_SURFACE_MEM_DEV_COPIED) HIP_API_CALL(hipStreamSynchronize(hip_stream_));
return 1;
}
int FFMpegVideoDecoder::DecodeFrame(const uint8_t *data, size_t size, int pkt_flags, int64_t pts, int *num_decoded_pics) {
output_frame_cnt_ = 0, output_frame_cnt_ret_ = 0;
decoded_pic_cnt_ = 0;
RocdecPicParamsHost pic_params = {};
pic_params.bitstream_data_len = size;
pic_params.bitstream_data = data;
if (!data || size == 0) {
pic_params.flags = ROCDEC_PKT_ENDOFPICTURE; // mark end_of_picture flag for last frame
}
ROCDEC_API_CALL(rocDecDecodeFrameHost(roc_decoder_, &pic_params));
if (num_decoded_pics) {
*num_decoded_pics = output_frame_cnt_;
}
return output_frame_cnt_;
}
uint8_t* FFMpegVideoDecoder::GetFrame(int64_t *pts) {
if (output_frame_cnt_ > 0) {
std::lock_guard<std::mutex> lock(mtx_vp_frame_);
if (vp_frames_.size() > 0){
output_frame_cnt_--;
if (pts) *pts = vp_frames_[output_frame_cnt_ret_].pts;
return vp_frames_[output_frame_cnt_ret_++].frame_ptr;
}
}
return nullptr;
}
bool FFMpegVideoDecoder::ReleaseFrame(int64_t pTimestamp, bool b_flushing) {
// if not flushing the buffers are re-used, so keep them
if (!b_flushing)
return true;
else {
// remove frames in flushing mode
std::lock_guard<std::mutex> lock(mtx_vp_frame_);
DecFrameBuffer *fb = &vp_frames_[0];
if (pTimestamp != fb->pts) {
std::cerr << "Decoded Frame is released out of order" << std::endl;
return false;
}
vp_frames_.erase(vp_frames_.begin()); // get rid of the frames from the framestore
}
return true;
}
void FFMpegVideoDecoder::SaveFrameToFile(std::string output_file_name, void *surf_mem, OutputSurfaceInfo *surf_info, size_t rgb_image_size) {
uint8_t *hst_ptr = nullptr;
bool is_rgb = (rgb_image_size != 0);
uint64_t output_image_size = is_rgb ? rgb_image_size : surf_info->output_surface_size_in_bytes;
if (surf_info->mem_type == OUT_SURFACE_MEM_DEV_COPIED) {
if (hst_ptr == nullptr) {
hst_ptr = new uint8_t [output_image_size];
}
hipError_t hip_status = hipSuccess;
hip_status = hipMemcpyDtoH((void *)hst_ptr, surf_mem, output_image_size);
if (hip_status != hipSuccess) {
std::cerr << "ERROR: hipMemcpyDtoH failed! (" << hipGetErrorName(hip_status) << ")" << std::endl;
delete [] hst_ptr;
return;
}
} else
hst_ptr = static_cast<uint8_t *> (surf_mem);
if (current_output_filename.empty()) {
current_output_filename = output_file_name;
}
// don't overwrite to the same file if reconfigure is detected for a resolution changes.
if (is_output_surface_changed_) {
if (fp_out_) {
fclose(fp_out_);
fp_out_ = nullptr;
}
// Append the width and height of the new stream to the old file name to create a file name to save the new frames
// do this only if resolution changes within a stream (e.g., decoding a multi-resolution stream using the videoDecode app)
// don't append to the output_file_name if multiple output file name is provided (e.g., decoding multi-files using the videDecodeMultiFiles)
if (!current_output_filename.compare(output_file_name)) {
std::string::size_type const pos(output_file_name.find_last_of('.'));
extra_output_file_count_++;
std::string to_append = "_" + std::to_string(surf_info->output_width) + "_" + std::to_string(surf_info->output_height) + "_" + std::to_string(extra_output_file_count_);
if (pos != std::string::npos) {
output_file_name.insert(pos, to_append);
} else {
output_file_name += to_append;
}
}
is_output_surface_changed_ = false;
}
if (fp_out_ == nullptr) {
fp_out_ = fopen(output_file_name.c_str(), "wb");
}
if (fp_out_) {
if (!is_rgb) {
uint8_t *tmp_hst_ptr = hst_ptr;
int img_width = surf_info->output_width;
int img_height = surf_info->output_height;
int output_stride = surf_info->output_pitch;
if (img_width * surf_info->bytes_per_pixel == output_stride && img_height == surf_info->output_vstride) {
fwrite(hst_ptr, 1, output_image_size, fp_out_);
} else {
uint32_t width = surf_info->output_width * surf_info->bytes_per_pixel;
if (surf_info->bit_depth <= 16) {
for (int i = 0; i < surf_info->output_height; i++) {
fwrite(tmp_hst_ptr, 1, width, fp_out_);
tmp_hst_ptr += output_stride;
}
// dump chroma
uint32_t chroma_stride = (output_stride >> 1);
uint8_t *u_hst_ptr = hst_ptr + output_stride * surf_info->output_height;
uint8_t *v_hst_ptr = u_hst_ptr + chroma_stride * chroma_height_;
for (int i = 0; i < chroma_height_; i++) {
fwrite(u_hst_ptr, 1, chroma_width_, fp_out_);
u_hst_ptr += chroma_stride;
}
if (num_chroma_planes_ == 2) {
for (int i = 0; i < chroma_height_; i++) {
fwrite(v_hst_ptr, 1, chroma_width_, fp_out_);
v_hst_ptr += chroma_stride;
}
}
}
}
} else {
fwrite(hst_ptr, 1, rgb_image_size, fp_out_);
}
}
if (hst_ptr && (surf_info->mem_type != OUT_SURFACE_MEM_HOST_COPIED)) {
delete [] hst_ptr;
}
}
projects/rocdecode/utils/ffmpegvideodecode/ffmpeg_video_dec.h
+155
Wyświetl plik
@@ -0,0 +1,155 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
extern "C" {
#include <libavcodec/avcodec.h>
#include <libavformat/avformat.h>
#include <libavutil/pixdesc.h>
#if USE_AVCODEC_GREATER_THAN_58_134
#include <libavcodec/bsf.h>
#endif
}
#include "rocvideodecode/roc_video_dec.h" // for derived class
#include "rocdecode/rocdecode_host.h"
/**
* FFMpegVideoDecoder: Derived class for FFMpeg based host decoder
*/
class FFMpegVideoDecoder: public RocVideoDecoder {
public:
/**
* @brief Construct a new FFMpegVideoDecoder object
*
* @param num_threads : number of cpu threads for the decoder
* @param out_mem_type : out_mem_type for the decoded surface
* @param codec : codec type
* @param force_zero_latency : no support in FFMpeg decoding (false)
* @param p_crop_rect : to crop output
* @param extract_user_SEI_Message : enable to extract SEI
* @param disp_delay : output delayed by #disp_delay surfaces
* @param max_width : Max. width for the output surface
* @param max_height : Max. height for the output surface
* @param clk_rate : FPS clock-rate
* @param no_multithreading : run FFMpeg decoder in the main thread (no multithreading)
*/
FFMpegVideoDecoder(int num_threads, OutputSurfaceMemoryType out_mem_type, rocDecVideoCodec codec, bool force_zero_latency = false,
const Rect *p_crop_rect = nullptr, bool extract_user_SEI_Message = false, uint32_t disp_delay = 0, int max_width = 0, int max_height = 0,
uint32_t clk_rate = 1000);
/**
* @brief destructor
*
*/
~FFMpegVideoDecoder();
/**
* @brief this function decodes a frame and returns the number of frames avalable for display
*
* @param data - pointer to the compressed data buffer that is to be decoded
* @param size - size of the data buffer in bytes
* @param pts - presentation timestamp
* @param flags - video packet flags
* @param num_decoded_pics - nummber of pictures decoded in this call
* @return int - num of frames to display
*/
int DecodeFrame(const uint8_t *data, size_t size, int pkt_flags, int64_t pts = 0, int *num_decoded_pics = nullptr) override;
/**
* @brief Get the pointer to the Output Image Info
*
* @param surface_info ptr to output surface info
* @return true
* @return false
*/
bool GetOutputSurfaceInfo(OutputSurfaceInfo **surface_info) override;
/**
* @brief This function returns a decoded frame and timestamp. This should be called in a loop fetching all the available frames
*
*/
uint8_t* GetFrame(int64_t *pts) override;
/**
* @brief function to release frame after use by the application: Only used with "OUT_SURFACE_MEM_DEV_INTERNAL"
*
* @param pTimestamp - timestamp of the frame to be released (unmapped)
* @param b_flushing - true when flushing
* @return true - success
* @return false - falied
*/
bool ReleaseFrame(int64_t pTimestamp, bool b_flushing = false) override;
/**
* @brief Helper function to dump decoded output surface to file
*
* @param output_file_name - Output file name
* @param dev_mem - pointer to surface memory
* @param surf_info - surface info
* @param rgb_image_size - image size for rgb (optional). A non_zero value indicates the surf_mem holds an rgb interleaved image and the entire size will be dumped to file
*/
void SaveFrameToFile(std::string output_file_name, void *surf_mem, OutputSurfaceInfo *surf_info, size_t rgb_image_size = 0) override;
/**
* @brief This function is used to get the current frame size based on pixel format.
*/
virtual int GetFrameSize() override {CHECK_ZERO("Display width", disp_width_); return ((disp_width_ * disp_height_) + ((chroma_height_ * chroma_width_) * num_chroma_planes_)) * byte_per_pixel_; }
/**
* @brief This function reconfigure decoder if there is a change in sequence params.
*/
int ReconfigureDecoder(RocdecVideoFormat *p_video_format) override;
private:
/**
* @brief Callback function to be registered for getting a callback when decoding of sequence starts
*/
static int ROCDECAPI FFMpegHandleVideoSequenceProc(void *p_user_data, RocdecVideoFormatHost *p_video_format) { return ((FFMpegVideoDecoder *)p_user_data)->HandleVideoSequence(p_video_format); }
/**
* @brief Callback function to be registered for getting a callback when a decoded frame is available for display
*/
static int ROCDECAPI FFMpegHandlePictureDisplayProc(void *p_user_data, RocdecParserDispInfo *p_disp_info) { return ((FFMpegVideoDecoder *)p_user_data)->HandlePictureDisplay(p_disp_info); }
/**
* @brief Callback function to be registered for getting a callback when all the unregistered user SEI Messages are parsed for a frame.
*/
static int ROCDECAPI FFMpegHandleSEIMessagesProc(void *p_user_data, RocdecSeiMessageInfo *p_sei_message_info) { return ((FFMpegVideoDecoder *)p_user_data)->GetSEIMessage(p_sei_message_info); }
/**
* @brief This function gets called when a sequence is ready to be decoded. The function also gets called
when there is format change
*/
int HandleVideoSequence(RocdecVideoFormatHost *p_video_format);
/**
* @brief This function gets called after a picture is decoded and available for display. Frames are fetched and stored in
internal buffer
*/
int HandlePictureDisplay(RocdecParserDispInfo *p_disp_info);
/**
* @brief This function gets called when all unregistered user SEI messages are parsed for a frame
*/
int GetSEIMessage(RocdecSeiMessageInfo *p_sei_message_info) { return RocVideoDecoder::GetSEIMessage(p_sei_message_info);};
};
projects/rocdecode/utils/md5.h
+166
Wyświetl plik
@@ -0,0 +1,166 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
extern "C" {
#include "libavutil/md5.h"
#include "libavutil/mem.h"
}
#include "roc_video_dec.h"
/*!
* \file
* \brief The MD5 message digest generation utility.
*/
class MD5Generator {
public:
MD5Generator() {};
~MD5Generator() {};
/*! \brief Function to start MD5 calculation
*/
void InitMd5() {
md5_ctx_ = av_md5_alloc();
av_md5_init(md5_ctx_);
}
/*! \brief Function to update MD5 digest for a device data buffer
* \param [in] data_buf Pointer to the data buffer
* \param [in] buf_size Buffer info
*/
void UpdateMd5ForDataBuffer(void *data_buf, int buf_size) {
uint8_t *hstPtr = nullptr;
hstPtr = new uint8_t[buf_size];
hipError_t hip_status = hipSuccess;
hip_status = hipMemcpyDtoH((void *)hstPtr, data_buf, buf_size);
if (hip_status != hipSuccess) {
std::cerr << "ERROR: hipMemcpyDtoH failed! (" << hip_status << ")" << std::endl;
delete [] hstPtr;
return;
}
av_md5_update(md5_ctx_, hstPtr, buf_size);
if (hstPtr) {
delete [] hstPtr;
}
}
/*! \brief Function to update MD5 digest for a decoded frame
* \param [in] surf_mem Pointer to surface memory
* \param [in] surf_info Surface info
*/
void UpdateMd5ForFrame(void *surf_mem, OutputSurfaceInfo *surf_info) {
int i;
uint8_t *hst_ptr = nullptr;
uint64_t output_image_size = surf_info->output_surface_size_in_bytes;
if (surf_info->mem_type == OUT_SURFACE_MEM_DEV_INTERNAL || surf_info->mem_type == OUT_SURFACE_MEM_DEV_COPIED) {
if (hst_ptr == nullptr) {
hst_ptr = new uint8_t [output_image_size];
}
hipError_t hip_status = hipSuccess;
hip_status = hipMemcpyDtoH((void *)hst_ptr, surf_mem, output_image_size);
if (hip_status != hipSuccess) {
std::cerr << "ERROR: hipMemcpyDtoH failed! (" << hip_status << ")" << std::endl;
delete [] hst_ptr;
return;
}
} else
hst_ptr = static_cast<uint8_t *> (surf_mem);
if (hst_ptr == nullptr) {
ROCDEC_ERR("Null surface pointer.");
return;
}
// Need to covert interleaved planar to stacked planar, assuming 4:2:0 chroma sampling.
uint8_t *stacked_ptr = new uint8_t [output_image_size];
uint8_t *tmp_hst_ptr = hst_ptr;
int output_stride = surf_info->output_pitch;
tmp_hst_ptr += (surf_info->disp_rect.top * output_stride) + surf_info->disp_rect.left * surf_info->bytes_per_pixel;
uint8_t *tmp_stacked_ptr = stacked_ptr;
int img_width = surf_info->output_width;
int img_height = surf_info->output_height;
// Luma
if (img_width * surf_info->bytes_per_pixel == output_stride && img_height == surf_info->output_vstride) {
memcpy(stacked_ptr, hst_ptr, img_width * surf_info->bytes_per_pixel * img_height);
} else {
for (i = 0; i < img_height; i++) {
memcpy(tmp_stacked_ptr, tmp_hst_ptr, img_width * surf_info->bytes_per_pixel);
tmp_hst_ptr += output_stride;
tmp_stacked_ptr += img_width * surf_info->bytes_per_pixel;
}
}
// Chroma
int img_width_chroma = img_width >> 1;
tmp_hst_ptr = hst_ptr + output_stride * surf_info->output_vstride;
if (surf_info->mem_type == OUT_SURFACE_MEM_DEV_INTERNAL) {
tmp_hst_ptr += ((surf_info->disp_rect.top >> 1) * output_stride) + (surf_info->disp_rect.left * surf_info->bytes_per_pixel);
}
tmp_stacked_ptr = stacked_ptr + img_width * surf_info->bytes_per_pixel * img_height; // Cb
uint8_t *tmp_stacked_ptr_v = tmp_stacked_ptr + img_width_chroma * surf_info->bytes_per_pixel * surf_info->chroma_height; // Cr
for (i = 0; i < surf_info->chroma_height; i++) {
for ( int j = 0; j < img_width_chroma; j++) {
uint8_t *src_ptr, *dst_ptr;
// Cb
src_ptr = &tmp_hst_ptr[j * surf_info->bytes_per_pixel * 2];
dst_ptr = &tmp_stacked_ptr[j * surf_info->bytes_per_pixel];
memcpy(dst_ptr, src_ptr, surf_info->bytes_per_pixel);
// Cr
src_ptr += surf_info->bytes_per_pixel;
dst_ptr = &tmp_stacked_ptr_v[j * surf_info->bytes_per_pixel];
memcpy(dst_ptr, src_ptr, surf_info->bytes_per_pixel);
}
tmp_hst_ptr += output_stride;
tmp_stacked_ptr += img_width_chroma * surf_info->bytes_per_pixel;
tmp_stacked_ptr_v += img_width_chroma * surf_info->bytes_per_pixel;
}
int img_size = img_width * surf_info->bytes_per_pixel * (img_height + surf_info->chroma_height);
// For 10/12 bit, convert from P010/P012 to LSB to match reference decoder output
if (surf_info->bytes_per_pixel == 2) {
uint16_t *ptr = reinterpret_cast<uint16_t *> (stacked_ptr);
for (i = 0; i < img_size / 2; i++) {
ptr[i] = ptr[i] >> (16 - surf_info->bit_depth);
}
}
av_md5_update(md5_ctx_, stacked_ptr, img_size);
if (hst_ptr && (surf_info->mem_type != OUT_SURFACE_MEM_HOST_COPIED)) {
delete [] hst_ptr;
}
delete [] stacked_ptr;
}
/*! \brief Function to complete MD5 calculation
* \param [out] digest Pointer to the 16 byte message digest
*/
void FinalizeMd5(uint8_t **digest) {
av_md5_final(md5_ctx_, md5_digest_);
av_freep(&md5_ctx_);
*digest = md5_digest_;
}
private:
struct AVMD5 *md5_ctx_;
uint8_t md5_digest_[16];
};
projects/rocdecode/utils/resize_kernels.cpp
+186
Wyświetl plik
@@ -0,0 +1,186 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include "resize_kernels.h"
#include "roc_video_dec.h"
/**
* @brief low level HIP kernel for Resize using nearest neighbor interpolation
*
* @tparam YuvUnitx2
* @param p_src - src Y pointer
* @param p_src_uv - src UV pointer
* @param src_pitch - src pitch
* @param p_dst - dst Y pointer
* @param p_dst_uv - dst UV pointer
* @param pitch - dst pitch
* @param width - dst width
* @param height - dst height
* @param fx_scale - xscale
* @param fy_scale - yscale
* @return
*/
template<typename YuvUnitx2>
static __global__ void ResizeHip(uint8_t *p_src, uint8_t *p_src_uv, int src_pitch,
uint8_t *p_dst, uint8_t *p_dst_uv, int pitch, int width, int height, float fx_scale, float fy_scale) {
int ix = blockIdx.x * blockDim.x + threadIdx.x,
iy = blockIdx.y * blockDim.y + threadIdx.y;
if (ix >= width / 2 || iy >= height / 2) {
return;
}
int x = ix * 2, y = iy * 2;
typedef decltype(YuvUnitx2::x) YuvUnit;
uint8_t *p_src_y = p_src + src_pitch * static_cast<uint32_t>(fmaf(y, fy_scale, 0.5 * fy_scale));
*(YuvUnitx2 *)(p_dst + y * pitch + x * sizeof(YuvUnit)) = YuvUnitx2 {
*(YuvUnit *)(p_src_y + static_cast<uint>(fmaf(x, fx_scale, 0.5 * fx_scale)) * sizeof(YuvUnit)),
*(YuvUnit *)(p_src_y + static_cast<uint>(fmaf(x + 1, fx_scale, 0.5 * fx_scale) * sizeof(YuvUnit)))
};
y++;
p_src_y = p_src + src_pitch * static_cast<uint32_t>(fmaf(y, fy_scale, 0.5 * fy_scale));
*(YuvUnitx2 *)(p_dst + y * pitch + x * sizeof(YuvUnit)) = YuvUnitx2 {
*(YuvUnit *)(p_src_y + static_cast<uint>(fmaf(x, fx_scale, 0.5 * fx_scale)) * sizeof(YuvUnit)),
*(YuvUnit *)(p_src_y + static_cast<uint>(fmaf(x + 1, fx_scale, 0.5 * fx_scale)) * sizeof(YuvUnit))
};
YuvUnit *p_uv = (YuvUnit *) (p_src_uv + static_cast<uint>(fmaf(ix, fx_scale, fx_scale * 0.5)) * sizeof(YuvUnit) * 2 +
src_pitch * static_cast<uint>(fmaf(iy, fy_scale, 0.5 * fy_scale)));
*(YuvUnitx2 *)(p_dst_uv + iy * pitch + ix * 2 * sizeof(YuvUnit)) = YuvUnitx2{ (YuvUnit)p_uv[0], (YuvUnit)p_uv[1] };
}
template <typename YuvUnitx2>
static void Resize(unsigned char *p_dst, unsigned char* p_dst_uv, int dst_pitch, int dst_width, int dst_height,
unsigned char *p_src, unsigned char *p_src_uv, int src_pitch, int src_width, int src_height, hipStream_t hip_stream) {
ResizeHip<YuvUnitx2> <<<dim3((dst_width + 31) / 32, (dst_height + 31) / 32), dim3(16, 16), 0, hip_stream >>>(p_src, p_src_uv, src_pitch, p_dst, p_dst_uv,
dst_pitch, dst_width, dst_height, 1.0f * src_width / dst_width, 1.0f * src_height / dst_height);
}
void ResizeNv12(unsigned char *p_dst_nv12, int dst_pitch, int dst_width, int dst_height, unsigned char *p_src_nv12,
int src_pitch, int src_width, int src_height, unsigned char* p_src_nv12_uv, unsigned char* p_dst_nv12_uv, hipStream_t hip_stream)
{
unsigned char* p_src_uv = p_src_nv12_uv ? p_src_nv12_uv : p_src_nv12 + (src_pitch*src_height);
unsigned char* p_dst_uv = p_dst_nv12_uv ? p_dst_nv12_uv : p_dst_nv12 + (dst_pitch*dst_height);
return Resize<uchar2>(p_dst_nv12, p_dst_uv, dst_pitch, dst_width, dst_height, p_src_nv12, p_src_uv, src_pitch, src_width, src_height, hip_stream);
}
void ResizeP016(unsigned char *p_dst_p016, int dst_pitch, int dst_width, int dst_height, unsigned char *p_src_p016,
int src_pitch, int src_width, int src_height, unsigned char* p_src_p016_uv, unsigned char* p_dst_p016_uv, hipStream_t hip_stream)
{
unsigned char* p_src_uv = p_src_p016_uv ? p_src_p016_uv : p_src_p016 + (src_pitch*src_height);
unsigned char* p_dst_uv = p_dst_p016_uv ? p_dst_p016_uv : p_dst_p016 + (dst_pitch*dst_height);
return Resize<ushort2>(p_dst_p016, p_dst_uv, dst_pitch, dst_width, dst_height, p_src_p016, p_src_uv, src_pitch, src_width, src_height, hip_stream);
}
static __global__ void Scale(uint8_t *p_src, int src_pitch, uint8_t *p_dst, int pitch, int width,
int height, float fx_scale, float fy_scale) {
int x = blockIdx.x * blockDim.x + threadIdx.x,
y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= width || y >= height){
return;
}
// do nearest neighbor interpolation
uint8_t *p_src_xy = p_src + src_pitch * static_cast<uint>(fmaf(y, fy_scale, 0.5 * fy_scale)) + static_cast<uint>(fmaf(x, fx_scale, 0.5*fx_scale));
*(uint8_t*)(p_dst + (y * pitch) + x) = *p_src_xy;
}
static __global__ void Scale_UV(uint8_t *p_src, int src_pitch, uint8_t *p_dst, int pitch, int width,
int height, float fx_scale, float fy_scale) {
int x = blockIdx.x * blockDim.x + threadIdx.x,
y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= width || y >= height) {
return;
}
// do nearest neighbor interpolation
uint8_t *p_src_uv = p_src + src_pitch * static_cast<uint>(fmaf(y , fy_scale, 0.5 * fy_scale)) + static_cast<uint>(fmaf(x, fx_scale, 0.5 * fx_scale)) * 2;
uchar2 dst_uv = uchar2{ p_src_uv[0], p_src_uv[1] };
*(uchar2*)(p_dst + (y * pitch) + 2 * x) = dst_uv;
}
/**
* @brief Resize a single plane of Y/U/V or UV interleaved (reserved for future)
*
* @param dp_dst - dest pointer
* @param dst_pitch - Pitch of the dst plane
* @param dst_width - Width of the dst plane
* @param dst_height - Height of the dst plane
* @param dp_src - source pointer
* @param src_pitch - source pitch
* @param src_width - source width
* @param src_height - source height
* @param b_resize_uv - to resize UV plance or not
* @param hip_stream - Stream for launching the kernel
*/
void ResizeYUVHipLaunchKernel(uint8_t *dp_dst, int dst_pitch, int dst_width, int dst_height, uint8_t *dp_src, int src_pitch,
int src_width, int src_height, bool b_resize_uv, hipStream_t hip_stream) {
dim3 blockSize(16, 16, 1);
dim3 gridSize(((uint32_t)dst_width + blockSize.x - 1) / blockSize.x, ((uint32_t)dst_height + blockSize.y - 1) / blockSize.y, 1);
if (b_resize_uv) {
Scale_UV <<<gridSize, blockSize, 0, hip_stream >>>(dp_src, src_pitch, dp_dst,
dst_pitch, dst_width, dst_height, 1.0f * src_width / dst_width, 1.0f * src_height / dst_height);
}
else {
Scale <<<gridSize, blockSize, 0, hip_stream >>>(dp_src, src_pitch, dp_dst,
dst_pitch, dst_width, dst_height, 1.0f * src_width / dst_width, 1.0f * src_height / dst_height);
}
}
void ResizeYUV420(uint8_t *p_dst_y,
uint8_t* p_dst_u,
uint8_t* p_dst_v,
int dst_pitch_y,
int dst_pitch_uv,
int dst_width,
int dst_height,
uint8_t *p_src_y,
uint8_t* p_src_u,
uint8_t* p_src_v,
int src_pitch_y,
int src_pitch_uv,
int src_width,
int src_height,
bool b_nv12,
hipStream_t hip_stream) {
int uv_width_dst = (dst_width + 1) >> 1;
int uv_height_dst = (dst_width + 1) >> 1;
int uv_width_src = (src_width + 1) >> 1;
int uv_height_src = (src_height + 1) >> 1;
// Scale Y plane
ResizeYUVHipLaunchKernel(p_dst_y, dst_pitch_y, dst_width, dst_height, p_src_y, src_pitch_y, src_width, src_height, 0, hip_stream);
if (b_nv12) {
ResizeYUVHipLaunchKernel(p_dst_u, dst_pitch_uv, uv_width_dst, uv_height_dst, p_src_u, src_pitch_uv, uv_width_src, uv_height_src, b_nv12, hip_stream);
} else {
ResizeYUVHipLaunchKernel(p_dst_u, dst_pitch_uv, uv_width_dst, uv_height_dst, p_src_u, src_pitch_uv, uv_width_src, uv_height_src, b_nv12, hip_stream);
ResizeYUVHipLaunchKernel(p_dst_v, dst_pitch_uv, uv_width_dst, uv_height_dst, p_src_v, src_pitch_uv, uv_width_src, uv_height_src, b_nv12, hip_stream);
}
}
projects/rocdecode/utils/resize_kernels.h
+103
Wyświetl plik
@@ -0,0 +1,103 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include <stdint.h>
#include <hip/hip_runtime.h>
/**
* @brief Function to resize both planes of an NV12 image
*
*
* @param p_dst_nv12 - destination pointer Y plane
* @param dst_pitch - destination pitch
* @param dst_width - destination width
* @param dst_height - destination height
* @param p_src_nv12 - source pointer
* @param src_pitch - source pitch
* @param src_width - source width
* @param src_height - source height
* @param p_src_nv12_uv - source pointer of UV plane
* @param hip_stream - Stream for launching the kernel
*/
void ResizeNv12(uint8_t *p_dst_nv12, int dst_pitch, int dst_width, int dst_height, uint8_t *p_src_nv12,
int src_pitch, int src_width, int src_height, unsigned char* p_src_nv12_uv, unsigned char* p_dst_nv12_uv, hipStream_t hip_stream);
/**
* @brief
*
* @param p_dst_p016
* @param dst_pitch
* @param dst_width
* @param dst_height
* @param p_src_p016
* @param src_pitch
* @param src_width
* @param src_height
* @param p_src_p016_uv
* @param p_dst_p016_uv
* @param hip_stream - Stream for launching the kernel
*/
void ResizeP016(uint8_t *p_dst_p016, int dst_pitch, int dst_width, int dst_height, uint8_t *p_src_p016, int src_pitch,
int src_width, int src_height, unsigned char* p_src_p016_uv, unsigned char* p_dst_p016_uv, hipStream_t hip_stream);
/**
* @brief Function to resize 420 YUV image
*
* @param p_dst_y - Destination Y plane pointer
* @param p_dst_u - Destination U plane pointer
* @param p_dst_v - Destination V plane pointer
* @param dst_pitch_y - Destination Pitch Y
* @param dst_pitch_uv - Destination Pitch UV
* @param dst_width - Destination Width
* @param dst_height - Destination Height
* @param p_src_y - Src Y plane pointer
* @param p_src_u - Src U plane pointer
* @param p_src_v - Src V plane pointer
* @param src_pitch_y - Src Pitch Y
* @param src_pitch_uv - Src Pitch UV
* @param src_width - Src Width
* @param src_height - Src Height
* @param b_nv12 - Is uv interleaved?
* @param hip_stream - Stream for launching the kernel
*/
void ResizeYUV420(uint8_t *p_dst_y, uint8_t* p_dst_u, uint8_t* p_dst_v, int dst_pitch_y, int dst_pitch_uv,
int dst_width, int dst_height, uint8_t *p_src_y, uint8_t* p_src_u, uint8_t* p_src_v,
int src_pitch_y, int src_pitch_uv, int src_width, int src_height, bool b_nv12 = false, hipStream_t hip_stream = nullptr);
/**
* @brief The function to launch ResizeYUV HIP kernel
*
* @param dp_dst - dest pointer
* @param dst_pitch - Pitch of the dst plane
* @param dst_width - Width of the dst plane
* @param dst_height - Height of the dst plane
* @param dp_src - source pointer
* @param src_pitch - source pitch
* @param src_width - source width
* @param src_height - source height
* @param b_resize_uv - to resize UV plance or not
* @param hip_stream - Stream for launching the kernel
*/
void ResizeYUVHipLaunchKernel(uint8_t *dp_dst, int dst_pitch, int dst_width, int dst_height, uint8_t *dp_src, int src_pitch,
int src_width, int src_height, bool b_resize_uv = false, hipStream_t hip_stream = nullptr);
projects/rocdecode/utils/rocvideodecode/roc_video_dec.cpp
+1122
Wyświetl plik
Plik diff jest za duży Load Diff
projects/rocdecode/utils/rocvideodecode/roc_video_dec.h
+531
Wyświetl plik
@@ -0,0 +1,531 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include <stdint.h>
#include <mutex>
#include <vector>
#include <string>
#include <iostream>
#include <sstream>
#include <string.h>
#include <queue>
#include <stdexcept>
#include <exception>
#include <cstring>
#include <unordered_map>
#include <chrono>
#include <thread>
#include <hip/hip_runtime.h>
#include "rocdecode/rocdecode.h"
#include "rocdecode/rocparser.h"
/*!
* \file
* \brief The AMD Video Decode Library.
*
* \defgroup group_amd_roc_video_dec rocDecode Video Decode: AMD Video Decode API
* \brief AMD The rocDecode video decoder for AMD’s GPUs.
*/
#define MAX_FRAME_NUM 16
typedef int (ROCDECAPI *PFNRECONFIGUEFLUSHCALLBACK)(void *, uint32_t, void *);
typedef enum SeiAvcHevcPayloadType_enum {
SEI_TYPE_TIME_CODE = 136,
SEI_TYPE_USER_DATA_UNREGISTERED = 5
} SeiAvcHevcPayloadType;
typedef enum OutputSurfaceMemoryType_enum {
OUT_SURFACE_MEM_DEV_INTERNAL = 0, /**< Internal interopped decoded surface memory(original mapped decoded surface) */
OUT_SURFACE_MEM_DEV_COPIED = 1, /**< decoded output will be copied to a separate device memory (the user doesn't need to call release) **/
OUT_SURFACE_MEM_HOST_COPIED = 2, /**< decoded output will be copied to a separate host memory (the user doesn't need to call release) **/
OUT_SURFACE_MEM_NOT_MAPPED = 3 /**< < decoded output is not available (interop won't be used): useful for decode only performance app*/
} OutputSurfaceMemoryType;
#define TOSTR(X) std::to_string(static_cast<int>(X))
#define STR(X) std::string(X)
#if DBGINFO
#define ROCDEC_INFO(X) std::clog << "[INF] " << " {" << __func__ <<"} " << " " << X << std::endl;
#else
#define ROCDEC_INFO(X) ;
#endif
#define ROCDEC_ERR(X) std::cerr << "[ERR] " << " {" << __func__ <<"} " << " " << X << std::endl;
inline int GetChromaPlaneCount(rocDecVideoSurfaceFormat surface_format) {
int num_planes = 1;
switch (surface_format) {
case rocDecVideoSurfaceFormat_NV12:
case rocDecVideoSurfaceFormat_P016:
num_planes = 1;
break;
case rocDecVideoSurfaceFormat_YUV444:
case rocDecVideoSurfaceFormat_YUV444_16Bit:
case rocDecVideoSurfaceFormat_YUV420:
case rocDecVideoSurfaceFormat_YUV420_16Bit:
case rocDecVideoSurfaceFormat_YUV422:
case rocDecVideoSurfaceFormat_YUV422_16Bit:
num_planes = 2;
break;
}
return num_planes;
};
inline float GetChromaHeightFactor(rocDecVideoSurfaceFormat surface_format) {
float factor = 0.5;
switch (surface_format) {
case rocDecVideoSurfaceFormat_NV12:
case rocDecVideoSurfaceFormat_P016:
case rocDecVideoSurfaceFormat_YUV420:
case rocDecVideoSurfaceFormat_YUV420_16Bit:
factor = 0.5;
break;
case rocDecVideoSurfaceFormat_YUV422:
case rocDecVideoSurfaceFormat_YUV422_16Bit:
case rocDecVideoSurfaceFormat_YUV444:
case rocDecVideoSurfaceFormat_YUV444_16Bit:
factor = 1.0;
break;
}
return factor;
};
class RocVideoDecodeException : public std::exception {
public:
explicit RocVideoDecodeException(const std::string& message, const int err_code):_message(message), _err_code(err_code) {}
explicit RocVideoDecodeException(const std::string& message):_message(message), _err_code(-1) {}
virtual const char* what() const throw() override {
return _message.c_str();
}
int Geterror_code() const { return _err_code; }
private:
std::string _message;
int _err_code;
};
#define ROCDEC_THROW(X, CODE) throw RocVideoDecodeException(" { " + std::string(__func__) + " } " + X , CODE);
#define ROCDEC_API_CALL( rocDecAPI ) \
do { \
rocDecStatus error_code = rocDecAPI; \
if( error_code != ROCDEC_SUCCESS) { \
std::ostringstream error_log; \
error_log << #rocDecAPI << " returned " << rocDecGetErrorName(error_code) << " at " <<__FILE__ <<":" << __LINE__;\
ROCDEC_THROW(error_log.str(), error_code); \
} \
} while (0)
#define HIP_API_CALL( call ) \
do { \
hipError_t hip_status = call; \
if (hip_status != hipSuccess) { \
const char *sz_err_name = NULL; \
sz_err_name = hipGetErrorName(hip_status); \
std::ostringstream error_log; \
error_log << "hip API error " << sz_err_name ; \
ROCDEC_THROW(error_log.str(), hip_status); \
} \
} \
while (0)
#define CHECK_ZERO(str, value) \
if (value == 0) { \
ROCDEC_ERR(STR(str) + " is 0."); \
}
struct Rect {
int left;
int top;
int right;
int bottom;
};
struct Dim {
int w, h;
};
static inline int align(int value, int alignment) {
return (value + alignment - 1) & ~(alignment - 1);
}
typedef struct DecFrameBuffer_ {
uint8_t *frame_ptr; /**< device memory pointer for the decoded frame */
int64_t pts; /**< timestamp for the decoded frame */
int picture_index; /**< surface index for the decoded frame */
} DecFrameBuffer;
typedef struct OutputSurfaceInfoType {
uint32_t output_width; /**< Output width of decoded surface*/
uint32_t output_height; /**< Output height of decoded surface*/
uint32_t output_pitch; /**< Output pitch in bytes of luma plane, chroma pitch can be inferred based on chromaFormat*/
uint32_t output_vstride; /**< Output vertical stride in case of using internal mem pointer **/
uint32_t chroma_height; /**< Chroma plane height **/
Rect disp_rect; /**< Display area **/
uint32_t bytes_per_pixel; /**< Output BytesPerPixel of decoded image*/
uint32_t bit_depth; /**< Output BitDepth of the image*/
uint32_t num_chroma_planes; /**< Output Chroma number of planes*/
uint64_t output_surface_size_in_bytes; /**< Output Image Size in Bytes; including both luma and chroma planes*/
rocDecVideoSurfaceFormat surface_format; /**< Chroma format of the decoded image*/
OutputSurfaceMemoryType mem_type; /**< Output mem_type of the surface*/
} OutputSurfaceInfo;
typedef struct ReconfigParams_t {
PFNRECONFIGUEFLUSHCALLBACK p_fn_reconfigure_flush;
void *p_reconfig_user_struct;
uint32_t reconfig_flush_mode;
} ReconfigParams;
class RocVideoDecoder {
public:
/**
* @brief Construct a new Roc Video Decoder object
*
* @param device_id : device_id to initialize HIP and VCN
* @param out_mem_type : out_mem_type for the decoded surface
* @param codec : codec type
* @param force_zero_latency : to force zero latency (output in decoding orde)
* @param p_crop_rect : to crop output
* @param extract_user_SEI_Message : enable to extract SEI
* @param disp_delay : output delayed by #disp_delay surfaces
* @param max_width : Max. width for the output surface
* @param max_height : Max. height for the output surface
* @param clk_rate : FPS clock-rate
*/
RocVideoDecoder(int device_id, OutputSurfaceMemoryType out_mem_type, rocDecVideoCodec codec, bool force_zero_latency = false,
const Rect *p_crop_rect = nullptr, bool extract_user_SEI_Message = false, uint32_t disp_delay = 0, int max_width = 0, int max_height = 0,
uint32_t clk_rate = 1000, bool skip_init = false);
virtual ~RocVideoDecoder();
rocDecVideoCodec GetCodecId() { return codec_id_; }
/**
* @brief Get the output frame width
*/
uint32_t GetWidth() {CHECK_ZERO("Display width", disp_width_); return disp_width_;}
/**
* @brief This function is used to get the actual decode width
*/
int GetDecodeWidth() {CHECK_ZERO("Coded width", coded_width_); return coded_width_; }
/**
* @brief Get the output frame height
*/
uint32_t GetHeight() {CHECK_ZERO("Display height", disp_height_); return disp_height_; }
/**
* @brief This function is used to get the current chroma height.
*/
int GetChromaHeight() {CHECK_ZERO("Chroma height", chroma_height_); return chroma_height_; }
/**
* @brief This function is used to get the number of chroma planes.
*/
int GetNumChromaPlanes() {return num_chroma_planes_; }
/**
* @brief This function is used to get the current frame size based on pixel format.
*/
virtual int GetFrameSize() {CHECK_ZERO("Display width", disp_width_); return disp_width_ * (disp_height_ + (chroma_height_ * num_chroma_planes_)) * byte_per_pixel_; }
/**
* @brief Get the Bit Depth and BytesPerPixel associated with the pixel format
*
* @return uint32_t
*/
uint32_t GetBitDepth() {return (bitdepth_minus_8_ + 8); }
uint32_t GetBytePerPixel() {CHECK_ZERO("Bytes per pixel", byte_per_pixel_); return byte_per_pixel_; }
/**
* @brief Functions to get the output surface attributes
*/
size_t GetSurfaceSize() {CHECK_ZERO("Surface size", surface_size_); return surface_size_; }
uint32_t GetSurfaceStride() {CHECK_ZERO("Surface stride", surface_stride_); return surface_stride_; }
//RocDecImageFormat GetSubsampling() { return subsampling_; }
/**
* @brief Get the name of the output format
*
* @param codec_id
* @return std::string
*/
const char *GetCodecFmtName(rocDecVideoCodec codec_id);
/**
* @brief function to return the name from surface_format_id
*
* @param surface_format_id - enum for surface format
* @return const char*
*/
const char *GetSurfaceFmtName(rocDecVideoSurfaceFormat surface_format_id);
/**
* @brief Get the pointer to the Output Image Info
*
* @param surface_info ptr to output surface info
* @return true
* @return false
*/
virtual bool GetOutputSurfaceInfo(OutputSurfaceInfo **surface_info);
/**
* @brief Function to set the Reconfig Params object
*
* @param p_reconfig_params: pointer to reconfig params struct
* @return true : success
* @return false : fail
*/
bool SetReconfigParams(ReconfigParams *p_reconfig_params, bool b_force_reconfig_flush = false);
/**
* @brief Function to force Reconfigure Flush: needed for random seeking to key frames
*
* @return int 1: Success 0: Fail
*/
int FlushAndReconfigure();
/**
* @brief this function decodes a frame and returns the number of frames avalable for display
*
* @param data - pointer to the data buffer that is to be decode
* @param size - size of the data buffer in bytes
* @param pts - presentation timestamp
* @param flags - video packet flags
* @param num_decoded_pics - nummber of pictures decoded in this call
* @return int - num of frames to display
*/
virtual int DecodeFrame(const uint8_t *data, size_t size, int pkt_flags, int64_t pts = 0, int *num_decoded_pics = nullptr);
/**
* @brief This function returns a decoded frame and timestamp. This should be called in a loop fetching all the available frames
*
*/
virtual uint8_t* GetFrame(int64_t *pts);
/**
* @brief function to release frame after use by the application: Only used with "OUT_SURFACE_MEM_DEV_INTERNAL"
*
* @param pTimestamp - timestamp of the frame to be released (unmapped)
* @param b_flushing - true when flushing
* @return true - success
* @return false - falied
*/
virtual bool ReleaseFrame(int64_t pTimestamp, bool b_flushing = false);
/**
* @brief utility function to save image to a file
*
* @param output_file_name - file to write
* @param dev_mem - dev_memory pointer of the frame
* @param image_info - output image info
* @param is_output_RGB - to write in RGB
*/
//void SaveImage(std::string output_file_name, void* dev_mem, OutputImageInfo* image_info, bool is_output_RGB = 0);
/**
* @brief Get the Device info for the current device
*
* @param device_name
* @param gcn_arch_name
* @param pci_bus_id
* @param pci_domain_id
* @param pci_device_id
*/
void GetDeviceinfo(std::string &device_name, std::string &gcn_arch_name, int &pci_bus_id, int &pci_domain_id, int &pci_device_id);
/**
* @brief Helper function to dump decoded output surface to file
*
* @param output_file_name - Output file name
* @param dev_mem - pointer to surface memory
* @param surf_info - surface info
* @param rgb_image_size - image size for rgb (optional). A non_zero value indicates the surf_mem holds an rgb interleaved image and the entire size will be dumped to file
*/
virtual void SaveFrameToFile(std::string output_file_name, void *surf_mem, OutputSurfaceInfo *surf_info, size_t rgb_image_size = 0);
/**
* @brief Helper funtion to close a existing file and dump to new file in case of multiple files using same decoder
*/
virtual void ResetSaveFrameToFile();
/**
* @brief Get the Num Of Flushed Frames from video decoder object
*
* @return int32_t
*/
int32_t GetNumOfFlushedFrames() { return num_frames_flushed_during_reconfig_;}
/*! \brief Function to wait for the decode completion of the last submitted picture
*/
void WaitForDecodeCompletion();
// Session overhead refers to decoder initialization and deinitialization time
void AddDecoderSessionOverHead(std::thread::id session_id, double duration) { session_overhead_[session_id] += duration; }
double GetDecoderSessionOverHead(std::thread::id session_id) {
if (session_overhead_.find(session_id) != session_overhead_.end()) {
return session_overhead_[session_id];
} else {
return 0;
}
}
/**
* @brief Check if the given Video Codec is supported on the given GPU
*
* @return rocDecStatus
*/
bool CodecSupported(int device_id, rocDecVideoCodec codec_id, uint32_t bit_depth);
/**
* @brief This function reconfigure decoder if there is a change in sequence params.
*/
virtual int ReconfigureDecoder(RocdecVideoFormat *p_video_format);
protected:
/**
* @brief Callback function to be registered for getting a callback when decoding of sequence starts
*/
static int ROCDECAPI HandleVideoSequenceProc(void *p_user_data, RocdecVideoFormat *p_video_format) { return ((RocVideoDecoder *)p_user_data)->HandleVideoSequence(p_video_format); }
/**
* @brief Callback function to be registered for getting a callback when a decoded frame is ready to be decoded
*/
static int ROCDECAPI HandlePictureDecodeProc(void *p_user_data, RocdecPicParams *p_pic_params) { return ((RocVideoDecoder *)p_user_data)->HandlePictureDecode(p_pic_params); }
/**
* @brief Callback function to be registered for getting a callback when a decoded frame is available for display
*/
static int ROCDECAPI HandlePictureDisplayProc(void *p_user_data, RocdecParserDispInfo *p_disp_info) { return ((RocVideoDecoder *)p_user_data)->HandlePictureDisplay(p_disp_info); }
/**
* @brief Callback function to be registered for getting a callback when all the unregistered user SEI Messages are parsed for a frame.
*/
static int ROCDECAPI HandleSEIMessagesProc(void *p_user_data, RocdecSeiMessageInfo *p_sei_message_info) { return ((RocVideoDecoder *)p_user_data)->GetSEIMessage(p_sei_message_info); }
/**
* @brief This function gets called when a sequence is ready to be decoded. The function also gets called
when there is format change
*/
int HandleVideoSequence(RocdecVideoFormat *p_video_format);
/**
* @brief This function gets called when a picture is ready to be decoded. rocDecDecodeFrame is called from this function
* to decode the picture
*/
int HandlePictureDecode(RocdecPicParams *p_pic_params);
/**
* @brief This function gets called after a picture is decoded and available for display. Frames are fetched and stored in
internal buffer
*/
int HandlePictureDisplay(RocdecParserDispInfo *p_disp_info);
/**
* @brief This function gets called when all unregistered user SEI messages are parsed for a frame
*/
int GetSEIMessage(RocdecSeiMessageInfo *p_sei_message_info);
/**
* @brief function to release all internal frames and clear the vp_frames_q_ (used with reconfigure): Only used with "OUT_SURFACE_MEM_DEV_INTERNAL"
*
* @return true - success
* @return false - falied
*/
bool ReleaseInternalFrames();
/**
* @brief Function to Initialize GPU-HIP
*
*/
bool InitHIP(int device_id);
/**
* @brief Function to get start time
*
*/
std::chrono::_V2::system_clock::time_point StartTimer();
/**
* @brief Function to get elapsed time
*
*/
double StopTimer(const std::chrono::_V2::system_clock::time_point &start_time);
int num_devices_;
int device_id_;
RocdecVideoParser rocdec_parser_ = nullptr;
rocDecDecoderHandle roc_decoder_ = nullptr;
OutputSurfaceMemoryType out_mem_type_ = OUT_SURFACE_MEM_DEV_INTERNAL;
rocDecVideoCodec codec_id_ = rocDecVideoCodec_NumCodecs;
bool b_force_zero_latency_ = false;
bool b_extract_sei_message_ = false;
uint32_t disp_delay_;
ReconfigParams *p_reconfig_params_ = nullptr;
bool b_force_recofig_flush_ = false;
int32_t num_frames_flushed_during_reconfig_ = 0;
hipDeviceProp_t hip_dev_prop_;
hipStream_t hip_stream_ = nullptr;
rocDecVideoChromaFormat video_chroma_format_ = rocDecVideoChromaFormat_420;
rocDecVideoSurfaceFormat video_surface_format_ = rocDecVideoSurfaceFormat_NV12;
RocdecSeiMessageInfo *curr_sei_message_ptr_ = nullptr;
RocdecSeiMessageInfo sei_message_display_q_[MAX_FRAME_NUM];
RocdecVideoFormat *curr_video_format_ptr_ = nullptr;
int output_frame_cnt_ = 0, output_frame_cnt_ret_ = 0;
int decoded_pic_cnt_ = 0;
int decode_poc_ = 0, pic_num_in_dec_order_[MAX_FRAME_NUM];
int num_alloced_frames_ = 0;
int last_decode_surf_idx_ = 0;
std::ostringstream input_video_info_str_;
int bitdepth_minus_8_ = 0;
uint32_t byte_per_pixel_ = 1;
uint32_t coded_width_ = 0;
uint32_t disp_width_ = 0;
uint32_t coded_height_ = 0;
uint32_t disp_height_ = 0;
uint32_t target_width_ = 0;
uint32_t target_height_ = 0;
int max_width_ = 0, max_height_ = 0;
uint32_t chroma_height_ = 0, chroma_width_ = 0;
uint32_t num_decode_surfaces_ = 0;
uint32_t num_chroma_planes_ = 0;
uint32_t num_components_ = 0;
uint32_t surface_stride_ = 0;
uint32_t surface_vstride_ = 0, chroma_vstride_ = 0; // vertical stride between planes: used when using internal dev memory
size_t surface_size_ = 0;
OutputSurfaceInfo output_surface_info_ = {};
std::mutex mtx_vp_frame_;
std::vector<DecFrameBuffer> vp_frames_; // vector of decoded frames
std::queue<DecFrameBuffer> vp_frames_q_;
Rect disp_rect_ = {}; // displayable area specified in the bitstream
Rect crop_rect_ = {}; // user specified region of interest within diplayable area disp_rect_
FILE *fp_sei_ = NULL;
FILE *fp_out_ = NULL;
bool is_output_surface_changed_ = false;
std::string current_output_filename = "";
uint32_t extra_output_file_count_ = 0;
std::thread::id decoder_session_id_; // Decoder session identifier. Used to gather session level stats.
std::unordered_map<std::thread::id, double> session_overhead_; // Records session overhead of initialization+deinitialization time. Format is (thread id, duration)
};
projects/rocdecode/utils/video_demuxer.h
+589
Wyświetl plik
@@ -0,0 +1,589 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include <iostream>
extern "C" {
#include <libavcodec/avcodec.h>
#include <libavformat/avformat.h>
#if USE_AVCODEC_GREATER_THAN_58_134
#include <libavcodec/bsf.h>
#endif
}
#include "rocdecode/rocdecode.h"
/*!
* \file
* \brief The AMD Video Demuxer for rocDecode Library.
*
* \defgroup group_amd_rocdecode_videodemuxer videoDemuxer: AMD rocDecode Video Demuxer API
* \brief AMD The rocDecode video demuxer API.
*/
/**
* @brief Enum for Seek mode
*
*/
typedef enum SeekModeEnum {
SEEK_MODE_EXACT_FRAME = 0,
SEEK_MODE_PREV_KEY_FRAME = 1,
SEEK_MODE_NUM,
} SeekMode;
/**
* @brief Enum for Seek Criteria
*
*/
typedef enum SeekCriteriaEnum {
SEEK_CRITERIA_FRAME_NUM = 0,
SEEK_CRITERIA_TIME_STAMP = 1,
SEEK_CRITERIA_NUM,
} SeekCriteria;
struct PacketData {
int32_t key;
int64_t pts;
int64_t dts;
uint64_t pos;
uintptr_t bsl_data;
uint64_t bsl;
uint64_t duration;
};
class VideoSeekContext {
public:
VideoSeekContext()
: use_seek_(false), seek_frame_(0), seek_mode_(SEEK_MODE_PREV_KEY_FRAME), seek_crit_(SEEK_CRITERIA_FRAME_NUM),
out_frame_pts_(0), out_frame_duration_(0), num_frames_decoded_(0U) {}
VideoSeekContext(uint64_t frame_id)
: use_seek_(true), seek_frame_(frame_id), seek_mode_(SEEK_MODE_PREV_KEY_FRAME),
seek_crit_(SEEK_CRITERIA_FRAME_NUM), out_frame_pts_(0), out_frame_duration_(0), num_frames_decoded_(0U) {}
VideoSeekContext& operator=(const VideoSeekContext& other) {
use_seek_ = other.use_seek_;
seek_frame_ = other.seek_frame_;
seek_mode_ = other.seek_mode_;
seek_crit_ = other.seek_crit_;
out_frame_pts_ = other.out_frame_pts_;
out_frame_duration_ = other.out_frame_duration_;
num_frames_decoded_ = other.num_frames_decoded_;
return *this;
}
/* Will be set to false when not seeking, true otherwise;
*/
bool use_seek_;
/* Frame we want to get. Set by user.
* Shall be set to frame timestamp in case seek is done by time.
*/
uint64_t seek_frame_;
/* Mode in which we seek. */
SeekMode seek_mode_;
/* Criteria by which we seek. */
SeekCriteria seek_crit_;
/* PTS of frame found after seek. */
int64_t out_frame_pts_;
/* Duration of frame found after seek. */
int64_t out_frame_duration_;
/* Number of frames that were decoded during seek. */
uint64_t num_frames_decoded_;
/* PTS of frame to seek as set by the user in seek_frame_. */
int64_t requested_frame_pts_;
};
// Video Demuxer Interface class
class VideoDemuxer {
public:
class StreamProvider {
public:
virtual ~StreamProvider() {}
virtual int GetData(uint8_t *buf, int buf_size) = 0;
virtual size_t GetBufferSize() = 0;
};
AVCodecID GetCodecID() { return av_video_codec_id_; };
VideoDemuxer(const char *input_file_path) : VideoDemuxer(CreateFmtContextUtil(input_file_path)) {}
VideoDemuxer(StreamProvider *stream_provider) : VideoDemuxer(CreateFmtContextUtil(stream_provider)) {av_io_ctx_ = av_fmt_input_ctx_->pb;}
~VideoDemuxer() {
if (!av_fmt_input_ctx_) {
return;
}
if (packet_) {
av_packet_free(&packet_);
}
if (packet_filtered_) {
av_packet_free(&packet_filtered_);
}
if (av_bsf_ctx_) {
av_bsf_free(&av_bsf_ctx_);
}
avformat_close_input(&av_fmt_input_ctx_);
if (av_io_ctx_) {
av_freep(&av_io_ctx_->buffer);
av_freep(&av_io_ctx_);
}
if (data_with_header_) {
av_free(data_with_header_);
}
}
bool Demux(uint8_t **video, int *video_size, int64_t *pts = nullptr) {
if (!av_fmt_input_ctx_) {
return false;
}
*video_size = 0;
if (packet_->data) {
av_packet_unref(packet_);
}
int ret = 0;
while ((ret = av_read_frame(av_fmt_input_ctx_, packet_)) >= 0 && packet_->stream_index != av_stream_) {
av_packet_unref(packet_);
}
if (ret < 0) {
return false;
}
if (is_h264_ || is_hevc_) {
if (packet_filtered_->data) {
av_packet_unref(packet_filtered_);
}
if (av_bsf_send_packet(av_bsf_ctx_, packet_) != 0) {
std::cerr << "ERROR: av_bsf_send_packet failed!" << std::endl;
return false;
}
if (av_bsf_receive_packet(av_bsf_ctx_, packet_filtered_) != 0) {
std::cerr << "ERROR: av_bsf_receive_packet failed!" << std::endl;
return false;
}
*video = packet_filtered_->data;
*video_size = packet_filtered_->size;
if (packet_filtered_->dts != AV_NOPTS_VALUE) {
pkt_dts_ = packet_filtered_->dts;
} else {
pkt_dts_ = packet_filtered_->pts;
}
if (pts) {
*pts = (int64_t) (packet_filtered_->pts * default_time_scale_ * time_base_);
pkt_duration_ = packet_filtered_->duration;
}
} else {
if (is_mpeg4_ && (frame_count_ == 0)) {
int ext_data_size = av_fmt_input_ctx_->streams[av_stream_]->codecpar->extradata_size;
if (ext_data_size > 0) {
data_with_header_ = (uint8_t *)av_malloc(ext_data_size + packet_->size - 3 * sizeof(uint8_t));
if (!data_with_header_) {
std::cerr << "ERROR: av_malloc failed!" << std::endl;
return false;
}
memcpy(data_with_header_, av_fmt_input_ctx_->streams[av_stream_]->codecpar->extradata, ext_data_size);
memcpy(data_with_header_ + ext_data_size, packet_->data + 3, packet_->size - 3 * sizeof(uint8_t));
*video = data_with_header_;
*video_size = ext_data_size + packet_->size - 3 * sizeof(uint8_t);
}
} else {
*video = packet_->data;
*video_size = packet_->size;
}
if (packet_->dts != AV_NOPTS_VALUE) {
pkt_dts_ = packet_->dts;
} else {
pkt_dts_ = packet_->pts;
}
if (pts) {
*pts = (int64_t)(packet_->pts * default_time_scale_ * time_base_);
pkt_duration_ = packet_->duration;
}
}
frame_count_++;
return true;
}
bool Seek(VideoSeekContext& seek_ctx, uint8_t** pp_video, int* video_size) {
/* !!! IMPORTANT !!!
* Across this function, packet decode timestamp (DTS) values are used to
* compare given timestamp against. This is done because DTS values shall
* monotonically increase during the course of decoding unlike PTS values
* which may be affected by frame reordering due to B frames.
*/
if (!is_seekable_) {
std::cerr << "ERROR: Seek isn't supported for this input." << std::endl;
return false;
}
if (IsVFR() && (SEEK_CRITERIA_FRAME_NUM == seek_ctx.seek_crit_)) {
std::cerr << "ERROR: Can't seek by frame number in VFR sequences. Seek by timestamp instead." << std::endl;
return false;
}
int64_t timestamp = 0;
// Seek for single frame;
auto seek_frame = [&](VideoSeekContext const& seek_ctx, int flags) {
bool seek_backward = true;
int ret = 0;
switch (seek_ctx.seek_crit_) {
case SEEK_CRITERIA_FRAME_NUM:
timestamp = TsFromFrameNumber(seek_ctx.seek_frame_);
ret = av_seek_frame(av_fmt_input_ctx_, av_stream_, timestamp, seek_backward ? AVSEEK_FLAG_BACKWARD | flags : flags);
break;
case SEEK_CRITERIA_TIME_STAMP:
timestamp = TsFromTime(seek_ctx.seek_frame_);
ret = av_seek_frame(av_fmt_input_ctx_, av_stream_, timestamp, seek_backward ? AVSEEK_FLAG_BACKWARD | flags : flags);
break;
default:
std::cerr << "ERROR: Invalid seek mode" << std::endl;
ret = -1;
}
if (ret < 0) {
throw std::runtime_error("ERROR: seeking for frame");
}
};
// Check if frame satisfies seek conditions;
auto is_seek_done = [&](PacketData& pkt_data, VideoSeekContext const& seek_ctx) {
int64_t target_ts = 0;
switch (seek_ctx.seek_crit_) {
case SEEK_CRITERIA_FRAME_NUM:
target_ts = TsFromFrameNumber(seek_ctx.seek_frame_);
break;
case SEEK_CRITERIA_TIME_STAMP:
target_ts = TsFromTime(seek_ctx.seek_frame_);
break;
default:
std::cerr << "ERROR::Invalid seek criteria" << std::endl;
return -1;
}
if (pkt_dts_ == target_ts) {
return 0;
} else if (pkt_dts_ > target_ts) {
return 1;
} else {
return -1;
};
};
/* This will seek for exact frame number;
* Note that decoder may not be able to decode such frame; */
auto seek_for_exact_frame = [&](PacketData& pkt_data, VideoSeekContext& seek_ctx) {
// Repetititive seek until seek condition is satisfied;
VideoSeekContext tmp_ctx(seek_ctx.seek_frame_);
seek_frame(tmp_ctx, AVSEEK_FLAG_ANY);
int seek_done = 0;
do {
if (!Demux(pp_video, video_size, &pkt_data.pts)) {
throw std::runtime_error("ERROR: Demux failed trying to seek for specified frame number/timestamp");
}
seek_done = is_seek_done(pkt_data, seek_ctx);
//TODO: one last condition, check for a target too high than available for timestamp
if (seek_done > 0) { // We've gone too far and need to seek backwards;
if ((tmp_ctx.seek_frame_--) >= 0) {
seek_frame(tmp_ctx, AVSEEK_FLAG_ANY);
}
} else if (seek_done < 0) { // Need to read more frames until we reach requested number;
tmp_ctx.seek_frame_++;
seek_frame(tmp_ctx, AVSEEK_FLAG_ANY);
}
if (tmp_ctx.seek_frame_ == seek_ctx.seek_frame_) // if frame 'N' is too far and frame 'N-1' is too less from target. Avoids infinite loop between N & N-1
break;
} while (seek_done != 0);
seek_ctx.out_frame_pts_ = pkt_data.pts;
seek_ctx.out_frame_duration_ = pkt_data.duration = pkt_duration_;
seek_ctx.requested_frame_pts_ = (int64_t) (timestamp * default_time_scale_ * time_base_);
};
// Seek for closest key frame in the past;
auto seek_for_prev_key_frame = [&](PacketData& pkt_data, VideoSeekContext& seek_ctx) {
seek_frame(seek_ctx, AVSEEK_FLAG_BACKWARD);
Demux(pp_video, video_size, &pkt_data.pts);
seek_ctx.num_frames_decoded_ = static_cast<uint64_t>(pkt_data.pts / 1000 * frame_rate_);
seek_ctx.out_frame_pts_ = pkt_data.pts;
seek_ctx.out_frame_duration_ = pkt_data.duration = pkt_duration_;
seek_ctx.requested_frame_pts_ = (int64_t) (timestamp * default_time_scale_ * time_base_);
};
PacketData pktData;
pktData.bsl_data = size_t(*pp_video);
pktData.bsl = *video_size;
switch (seek_ctx.seek_mode_) {
case SEEK_MODE_EXACT_FRAME:
seek_for_exact_frame(pktData, seek_ctx);
break;
case SEEK_MODE_PREV_KEY_FRAME:
seek_for_prev_key_frame(pktData, seek_ctx);
break;
default:
throw std::runtime_error("ERROR::Unsupported seek mode");
break;
}
return true;
}
const uint32_t GetWidth() const { return width_;}
const uint32_t GetHeight() const { return height_;}
const uint32_t GetChromaHeight() const { return chroma_height_;}
const uint32_t GetBitDepth() const { return bit_depth_;}
const uint32_t GetBytePerPixel() const { return byte_per_pixel_;}
const uint32_t GetBitRate() const { return bit_rate_;}
const double GetFrameRate() const {return frame_rate_;};
bool IsVFR() const { return frame_rate_ != avg_frame_rate_; };
int64_t TsFromTime(double ts_sec) {
// Convert integer timestamp representation to AV_TIME_BASE and switch to fixed_point
auto const ts_tbu = llround(ts_sec * AV_TIME_BASE);
// Rescale the timestamp to value represented in stream base units;
AVRational time_factor = {1, AV_TIME_BASE};
return av_rescale_q(ts_tbu, time_factor, av_fmt_input_ctx_->streams[av_stream_]->time_base);
}
int64_t TsFromFrameNumber(int64_t frame_num) {
auto const ts_sec = static_cast<double>(frame_num) / frame_rate_;
return TsFromTime(ts_sec);
}
private:
VideoDemuxer(AVFormatContext *av_fmt_input_ctx) : av_fmt_input_ctx_(av_fmt_input_ctx) {
av_log_set_level(AV_LOG_QUIET);
if (!av_fmt_input_ctx_) {
std::cerr << "ERROR: av_fmt_input_ctx_ is not vaild!" << std::endl;
return;
}
packet_ = av_packet_alloc();
packet_filtered_ = av_packet_alloc();
if (!packet_ || !packet_filtered_) {
std::cerr << "ERROR: av_packet_alloc failed!" << std::endl;
return;
}
if (avformat_find_stream_info(av_fmt_input_ctx_, nullptr) < 0) {
std::cerr << "ERROR: avformat_find_stream_info failed!" << std::endl;
return;
}
av_stream_ = av_find_best_stream(av_fmt_input_ctx_, AVMEDIA_TYPE_VIDEO, -1, -1, nullptr, 0);
if (av_stream_ < 0) {
std::cerr << "ERROR: av_find_best_stream failed!" << std::endl;
av_packet_free(&packet_);
av_packet_free(&packet_filtered_);
return;
}
av_video_codec_id_ = av_fmt_input_ctx_->streams[av_stream_]->codecpar->codec_id;
width_ = av_fmt_input_ctx_->streams[av_stream_]->codecpar->width;
height_ = av_fmt_input_ctx_->streams[av_stream_]->codecpar->height;
chroma_format_ = (AVPixelFormat)av_fmt_input_ctx_->streams[av_stream_]->codecpar->format;
bit_rate_ = av_fmt_input_ctx_->streams[av_stream_]->codecpar->bit_rate;
if (av_fmt_input_ctx_->streams[av_stream_]->r_frame_rate.den != 0)
frame_rate_ = static_cast<double>(av_fmt_input_ctx_->streams[av_stream_]->r_frame_rate.num) / static_cast<double>(av_fmt_input_ctx_->streams[av_stream_]->r_frame_rate.den);
if (av_fmt_input_ctx_->streams[av_stream_]->avg_frame_rate.den != 0)
avg_frame_rate_ = static_cast<double>(av_fmt_input_ctx_->streams[av_stream_]->avg_frame_rate.num) / static_cast<double>(av_fmt_input_ctx_->streams[av_stream_]->avg_frame_rate.den);
switch (chroma_format_) {
case AV_PIX_FMT_YUV420P10LE:
case AV_PIX_FMT_GRAY10LE:
bit_depth_ = 10;
chroma_height_ = (height_ + 1) >> 1;
byte_per_pixel_ = 2;
break;
case AV_PIX_FMT_YUV420P12LE:
bit_depth_ = 12;
chroma_height_ = (height_ + 1) >> 1;
byte_per_pixel_ = 2;
break;
case AV_PIX_FMT_YUV444P10LE:
bit_depth_ = 10;
chroma_height_ = height_ << 1;
byte_per_pixel_ = 2;
break;
case AV_PIX_FMT_YUV444P12LE:
bit_depth_ = 12;
chroma_height_ = height_ << 1;
byte_per_pixel_ = 2;
break;
case AV_PIX_FMT_YUV444P:
bit_depth_ = 8;
chroma_height_ = height_ << 1;
byte_per_pixel_ = 1;
break;
case AV_PIX_FMT_YUV420P:
case AV_PIX_FMT_YUVJ420P:
case AV_PIX_FMT_YUVJ422P:
case AV_PIX_FMT_YUVJ444P:
case AV_PIX_FMT_GRAY8:
bit_depth_ = 8;
chroma_height_ = (height_ + 1) >> 1;
byte_per_pixel_ = 1;
break;
default:
chroma_format_ = AV_PIX_FMT_YUV420P;
bit_depth_ = 8;
chroma_height_ = (height_ + 1) >> 1;
byte_per_pixel_ = 1;
}
AVRational time_base = av_fmt_input_ctx_->streams[av_stream_]->time_base;
time_base_ = av_q2d(time_base);
is_h264_ = av_video_codec_id_ == AV_CODEC_ID_H264 && (!strcmp(av_fmt_input_ctx_->iformat->long_name, "QuickTime / MOV")
|| !strcmp(av_fmt_input_ctx_->iformat->long_name, "FLV (Flash Video)")
|| !strcmp(av_fmt_input_ctx_->iformat->long_name, "Matroska / WebM"));
is_hevc_ = av_video_codec_id_ == AV_CODEC_ID_HEVC && (!strcmp(av_fmt_input_ctx_->iformat->long_name, "QuickTime / MOV")
|| !strcmp(av_fmt_input_ctx_->iformat->long_name, "FLV (Flash Video)")
|| !strcmp(av_fmt_input_ctx_->iformat->long_name, "Matroska / WebM"));
is_mpeg4_ = av_video_codec_id_ == AV_CODEC_ID_MPEG4 && (!strcmp(av_fmt_input_ctx_->iformat->long_name, "QuickTime / MOV")
|| !strcmp(av_fmt_input_ctx_->iformat->long_name, "FLV (Flash Video)")
|| !strcmp(av_fmt_input_ctx_->iformat->long_name, "Matroska / WebM"));
// Check if the input file allow seek functionality.
#if USE_AVCODEC_GREATER_THAN_58_134
is_seekable_ = true; //for latest version of FFMPeg, read_seek and read_seek2 is not exposed in AVFormatContext
#else
is_seekable_ = av_fmt_input_ctx_->iformat->read_seek || av_fmt_input_ctx_->iformat->read_seek2;
#endif
if (is_h264_) {
const AVBitStreamFilter *bsf = av_bsf_get_by_name("h264_mp4toannexb");
if (!bsf) {
std::cerr << "ERROR: av_bsf_get_by_name() failed" << std::endl;
av_packet_free(&packet_);
av_packet_free(&packet_filtered_);
return;
}
if (av_bsf_alloc(bsf, &av_bsf_ctx_) != 0) {
std::cerr << "ERROR: av_bsf_alloc failed!" << std::endl;
return;
}
avcodec_parameters_copy(av_bsf_ctx_->par_in, av_fmt_input_ctx_->streams[av_stream_]->codecpar);
if (av_bsf_init(av_bsf_ctx_) < 0) {
std::cerr << "ERROR: av_bsf_init failed!" << std::endl;
return;
}
}
if (is_hevc_) {
const AVBitStreamFilter *bsf = av_bsf_get_by_name("hevc_mp4toannexb");
if (!bsf) {
std::cerr << "ERROR: av_bsf_get_by_name() failed" << std::endl;
av_packet_free(&packet_);
av_packet_free(&packet_filtered_);
return;
}
if (av_bsf_alloc(bsf, &av_bsf_ctx_) != 0 ) {
std::cerr << "ERROR: av_bsf_alloc failed!" << std::endl;
return;
}
avcodec_parameters_copy(av_bsf_ctx_->par_in, av_fmt_input_ctx_->streams[av_stream_]->codecpar);
if (av_bsf_init(av_bsf_ctx_) < 0) {
std::cerr << "ERROR: av_bsf_init failed!" << std::endl;
return;
}
}
}
AVFormatContext *CreateFmtContextUtil(StreamProvider *stream_provider) {
AVFormatContext *ctx = nullptr;
if (!(ctx = avformat_alloc_context())) {
std::cerr << "ERROR: avformat_alloc_context failed" << std::endl;
return nullptr;
}
uint8_t *avioc_buffer = nullptr;
int avioc_buffer_size = stream_provider->GetBufferSize();
avioc_buffer = (uint8_t *)av_malloc(avioc_buffer_size);
if (!avioc_buffer) {
std::cerr << "ERROR: av_malloc failed!" << std::endl;
return nullptr;
}
av_io_ctx_ = avio_alloc_context(avioc_buffer, avioc_buffer_size,
0, stream_provider, &ReadPacket, nullptr, nullptr);
if (!av_io_ctx_) {
std::cerr << "ERROR: avio_alloc_context failed!" << std::endl;
return nullptr;
}
ctx->pb = av_io_ctx_;
if (avformat_open_input(&ctx, nullptr, nullptr, nullptr) != 0) {
std::cerr << "ERROR: avformat_open_input failed!" << std::endl;
return nullptr;
}
return ctx;
}
AVFormatContext *CreateFmtContextUtil(const char *input_file_path) {
avformat_network_init();
AVFormatContext *ctx = nullptr;
if (avformat_open_input(&ctx, input_file_path, nullptr, nullptr) != 0 ) {
std::cerr << "ERROR: avformat_open_input failed!" << std::endl;
return nullptr;
}
return ctx;
}
static int ReadPacket(void *data, uint8_t *buf, int buf_size) {
return ((StreamProvider *)data)->GetData(buf, buf_size);
}
AVFormatContext *av_fmt_input_ctx_ = nullptr;
AVIOContext *av_io_ctx_ = nullptr;
AVPacket* packet_ = nullptr;
AVPacket* packet_filtered_ = nullptr;
AVBSFContext *av_bsf_ctx_ = nullptr;
AVCodecID av_video_codec_id_;
AVPixelFormat chroma_format_;
double frame_rate_ = 0.0;
double avg_frame_rate_ = 0.0;
uint8_t *data_with_header_ = nullptr;
int av_stream_ = 0;
bool is_h264_ = false;
bool is_hevc_ = false;
bool is_mpeg4_ = false;
bool is_seekable_ = false;
int64_t default_time_scale_ = 1000;
double time_base_ = 0.0;
uint32_t frame_count_ = 0;
uint32_t width_ = 0;
uint32_t height_ = 0;
uint32_t chroma_height_ = 0;
uint32_t bit_depth_ = 0;
uint32_t byte_per_pixel_ = 0;
uint32_t bit_rate_ = 0;
// used for Seek Exact frame
int64_t pkt_dts_ = 0;
int64_t pkt_duration_ = 0;
};
static inline rocDecVideoCodec AVCodec2RocDecVideoCodec(AVCodecID av_codec) {
switch (av_codec) {
case AV_CODEC_ID_MPEG1VIDEO : return rocDecVideoCodec_MPEG1;
case AV_CODEC_ID_MPEG2VIDEO : return rocDecVideoCodec_MPEG2;
case AV_CODEC_ID_MPEG4 : return rocDecVideoCodec_MPEG4;
case AV_CODEC_ID_H264 : return rocDecVideoCodec_AVC;
case AV_CODEC_ID_HEVC : return rocDecVideoCodec_HEVC;
case AV_CODEC_ID_VP8 : return rocDecVideoCodec_VP8;
case AV_CODEC_ID_VP9 : return rocDecVideoCodec_VP9;
case AV_CODEC_ID_MJPEG : return rocDecVideoCodec_JPEG;
case AV_CODEC_ID_AV1 : return rocDecVideoCodec_AV1;
default : return rocDecVideoCodec_NumCodecs;
}
}
projects/rocdecode/utils/video_post_process.h
+119
Wyświetl plik
@@ -0,0 +1,119 @@
/*
Copyright (c) 2023 - 2026 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include "colorspace_kernels.h"
#include "resize_kernels.h"
#include "rocvideodecode/roc_video_dec.h" //for OutputSurfaceInfo
enum OutputFormatEnum {
native = 0, bgr, bgr48, rgb, rgb48, bgra, bgra64, rgba, rgba64
};
class VideoPostProcess {
public:
VideoPostProcess(){};
~VideoPostProcess(){};
void ColorConvertYUV2RGB(uint8_t *p_src, OutputSurfaceInfo *surf_info, uint8_t *rgb_dev_mem_ptr, OutputFormatEnum e_output_format, hipStream_t hip_stream) {
int rgb_width = (surf_info->output_width + 1) & ~1; // has to be a multiple of 2 for hip colorconvert kernels
// todo:: get color standard from the decoder
if (surf_info->surface_format == rocDecVideoSurfaceFormat_YUV444) {
if (e_output_format == bgr)
YUV444ToColor24<BGR24>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == bgra)
YUV444ToColor32<BGRA32>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 4 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgb)
YUV444ToColor24<RGB24>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgba)
YUV444ToColor32<RGBA32>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 4 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
} else if (surf_info->surface_format == rocDecVideoSurfaceFormat_NV12) {
if (e_output_format == bgr)
Nv12ToColor24<BGR24>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == bgra)
Nv12ToColor32<BGRA32>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 4 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgb)
Nv12ToColor24<RGB24>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgba)
Nv12ToColor32<RGBA32>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 4 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
}
if (surf_info->surface_format == rocDecVideoSurfaceFormat_YUV444_16Bit) {
if (e_output_format == bgr)
YUV444P16ToColor24<BGR24>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgb)
YUV444P16ToColor24<RGB24>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == bgr48)
YUV444P16ToColor48<BGR48>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 6 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgb48)
YUV444P16ToColor48<RGB48>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 6 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == bgra64)
YUV444P16ToColor64<BGRA64>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 8 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgba64)
YUV444P16ToColor64<RGBA64>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 8 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
} else if (surf_info->surface_format == rocDecVideoSurfaceFormat_P016) {
if (e_output_format == bgr)
P016ToColor24<BGR24>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgb)
P016ToColor24<RGB24>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == bgr48)
P016ToColor48<BGR48>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 6 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgb48)
P016ToColor48<RGB48>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 6 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == bgra64)
P016ToColor64<BGRA64>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 8 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
else if (e_output_format == rgba64)
P016ToColor64<RGBA64>(p_src, surf_info->output_pitch, static_cast<uint8_t *>(rgb_dev_mem_ptr), 8 * rgb_width, surf_info->output_width,
surf_info->output_height, surf_info->output_vstride, 0, hip_stream);
}
};
uint32_t GetRgbStride(OutputFormatEnum e_output_format, OutputSurfaceInfo *surf_info) {
uint32_t rgb_stride;
uint32_t rgb_width = (surf_info->output_width + 1) & ~1; // has to be a multiple of 2 for hip colorconvert kernels
if (surf_info->bit_depth == 8) {
rgb_stride = ((e_output_format == bgr) || (e_output_format == rgb)) ? rgb_width * 3 : rgb_width * 4; // bgr/bgra/rgb/rgba
} else {
rgb_stride = ((e_output_format == bgr) || (e_output_format == rgb)) ? rgb_width * 3 :
((e_output_format == bgr48) || (e_output_format == rgb48)) ? rgb_width * 6 : rgb_width * 8;
}
return rgb_stride;
};
};