AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

added testcase for videodecode with resizing and colorconversion to rgb (#608)

Bladeren bron 
* added testcase for videodecode with resizing and colorconversion to rgb

* Update test/CMakeLists.txt

Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>

* removed texture scaling code as it doesn't work on MI3xx

---------

Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>

[ROCm/rocdecode commit: 26f9ec23a6]
This commit is contained in:
Rajy Rawther
2025-06-24 10:51:59 -07:00
gecommit door GitHub
bovenliggende 26e30ddc92
commit 72b5c717e1
2 gewijzigde bestanden met toevoegingen van 14 en 143 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/rocdecode/test/CMakeLists.txt
+13
Bestand weergeven
@@ -242,3 +242,16 @@ add_test(
--build-generator "${CMAKE_GENERATOR}"
--test-command "rocdecodenegativetest"
)
# 14 - videoDecodeRGBResize
add_test(
NAME
video_decodeRGB-Resize
COMMAND
"${CMAKE_CTEST_COMMAND}"
--build-and-test "${ROCM_PATH}/share/rocdecode/samples/videoDecodeRGB"
"${CMAKE_CURRENT_BINARY_DIR}/videoDecodeRGB"
--build-generator "${CMAKE_GENERATOR}"
--test-command "videodecodergb"
-i "${ROCM_PATH}/share/rocdecode/video/AMD_driving_virtual_20-H265.mp4" -resize 640x360 -of rgb
)
projects/rocdecode/utils/resize_kernels.cpp
+1 -143
Bestand weergeven
@@ -23,52 +23,6 @@ THE SOFTWARE.
#include "resize_kernels.h"
#include "roc_video_dec.h"
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
/**
* @brief low level HIP kernel for Resize using tex2d
*
* @tparam YuvUnitx2
* @param tex_y - text2D object Y pointer
* @param tex_uv - text2D object UV pointer
* @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(hipTextureObject_t tex_y, hipTextureObject_t tex_uv,
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;
const int max_yuv_value = (1 << (sizeof(YuvUnit) * 8)) - 1;
*(YuvUnitx2 *)(p_dst + y * pitch + x * sizeof(YuvUnit)) = YuvUnitx2 {
(YuvUnit)(tex2D<float>(tex_y, x * fx_scale, y * fy_scale) * max_yuv_value),
(YuvUnit)(tex2D<float>(tex_y, (x + 1) * fx_scale, y * fy_scale) * max_yuv_value)
};
y++;
*(YuvUnitx2 *)(p_dst + y * pitch + x * sizeof(YuvUnit)) = YuvUnitx2 {
(YuvUnit)(tex2D<float>(tex_y, x * fx_scale, y * fy_scale) * max_yuv_value),
(YuvUnit)(tex2D<float>(tex_y, (x + 1) * fx_scale, y * fy_scale) * max_yuv_value)
};
float2 uv = tex2D<float2>(tex_uv, ix * fx_scale, iy * fy_scale + 0.5f);
*(YuvUnitx2 *)(p_dst_uv + iy * pitch + ix * 2 * sizeof(YuvUnit)) = YuvUnitx2{ (YuvUnit)(uv.x * max_yuv_value), (YuvUnit)(uv.y * max_yuv_value) };
}
#endif
/**
* @brief low level HIP kernel for Resize using nearest neighbor interpolation
*
@@ -118,40 +72,9 @@ static __global__ void ResizeHip(uint8_t *p_src, uint8_t *p_src_uv, int src_pitc
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) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
hipResourceDesc res_desc = {};
res_desc.resType = hipResourceTypePitch2D;
res_desc.res.pitch2D.devPtr = p_src;
res_desc.res.pitch2D.desc = hipCreateChannelDesc<decltype(YuvUnitx2::x)>();
res_desc.res.pitch2D.width = src_width;
res_desc.res.pitch2D.height = src_height;
res_desc.res.pitch2D.pitchInBytes = src_pitch;
hipTextureDesc tex_desc = {};
tex_desc.filterMode = hipFilterModeLinear;
tex_desc.readMode = hipReadModeNormalizedFloat;
hipTextureObject_t tex_y=0;
HIP_API_CALL(hipCreateTextureObject(&tex_y, &res_desc, &tex_desc, NULL));
res_desc.res.pitch2D.devPtr = p_src_uv;
res_desc.res.pitch2D.desc = hipCreateChannelDesc<YuvUnitx2>();
res_desc.res.pitch2D.width = src_width >> 1;
res_desc.res.pitch2D.height = src_height / 2;
hipTextureObject_t tex_uv=0;
HIP_API_CALL(hipCreateTextureObject(&tex_uv, &res_desc, &tex_desc, NULL));
ResizeHip<YuvUnitx2> <<<dim3((dst_width + 31) / 32, (dst_height + 31) / 32), dim3(16, 16), 0, hip_stream >>>(tex_y, tex_uv, p_dst, p_dst_uv,
dst_pitch, dst_width, dst_height, 1.0f * src_width / dst_width, 1.0f * src_height / dst_height);
HIP_API_CALL(hipDestroyTextureObject(tex_y));
HIP_API_CALL(hipDestroyTextureObject(tex_uv));
#else
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);
#endif
}
void ResizeNv12(unsigned char *p_dst_nv12, int dst_pitch, int dst_width, int dst_height, unsigned char *p_src_nv12,
@@ -171,34 +94,6 @@ void ResizeP016(unsigned char *p_dst_p016, int dst_pitch, int dst_width, int dst
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);
}
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
static __global__ void Scale_tex2D(hipTextureObject_t tex_src, 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;
}
*(unsigned char*)(p_dst + (y * pitch) + x) = (unsigned char)(fminf((tex2D<float>(tex_src, x * fx_scale, y * fy_scale)) * 255.0f, 255.0f));
}
static __global__ void Scale_UV_tex2D(hipTextureObject_t tex_src, 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;
}
float2 uv = tex2D<float2>(tex_src, x * fx_scale, y * fy_scale);
uchar2 dst_uv = uchar2{ (unsigned char)(fminf(uv.x * 255.0f, 255.0f)), (unsigned char)(fminf(uv.y * 255.0f, 255.0f)) };
*(uchar2*)(p_dst + (y * pitch) + 2 * x) = dst_uv;
}
#endif
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,
@@ -244,41 +139,6 @@ static __global__ void Scale_UV(uint8_t *p_src, int src_pitch, uint8_t *p_dst, i
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) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
hipResourceDesc res_desc = {};
res_desc.resType = hipResourceTypePitch2D;
res_desc.res.pitch2D.devPtr = dp_src;
res_desc.res.pitch2D.desc = b_resize_uv ? hipCreateChannelDesc<uchar2>() : hipCreateChannelDesc<unsigned char>();
res_desc.res.pitch2D.width = src_width;
res_desc.res.pitch2D.height = src_height;
res_desc.res.pitch2D.pitchInBytes = src_pitch;
hipTextureDesc tex_desc = {};
tex_desc.filterMode = hipFilterModeLinear;
tex_desc.readMode = hipReadModeNormalizedFloat;
tex_desc.addressMode[0] = hipAddressModeClamp;
tex_desc.addressMode[1] = hipAddressModeClamp;
tex_desc.addressMode[2] = hipAddressModeClamp;
hipTextureObject_t tex_src = 0;
HIP_API_CALL(hipCreateTextureObject(&tex_src, &res_desc, &tex_desc, NULL));
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_tex2D <<<gridSize, blockSize, 0, hip_stream >>>(tex_src, dp_dst,
dst_pitch, dst_width, dst_height, 1.0f * src_width / dst_width, 1.0f * src_height / dst_height);
}
else{
Scale_tex2D <<<gridSize, blockSize, 0, hip_stream >>>(tex_src, dp_dst,
dst_pitch, dst_width, dst_height, 1.0f * src_width / dst_width, 1.0f * src_height / dst_height);
}
HIP_API_CALL(hipGetLastError());
HIP_API_CALL(hipDestroyTextureObject(tex_src));
#else
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);
@@ -290,8 +150,6 @@ void ResizeYUVHipLaunchKernel(uint8_t *dp_dst, int dst_pitch, int dst_width, int
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);
}
#endif
}
void ResizeYUV420(uint8_t *p_dst_y,