added testcase for videodecode with resizing and colorconversion to rgb (#608)
* 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>
Этот коммит содержится в:
Rajy Rawther
коммит произвёл
GitHub
GitHub
родитель
a2e1bf1920
Коммит
26f9ec23a6
@@ -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
|
||||
)
|
||||
|
||||
@@ -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,
|
||||
|
||||
Ссылка в новой задаче
Block a user