diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 79d61a27..e1566158 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -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 +) diff --git a/utils/resize_kernels.cpp b/utils/resize_kernels.cpp index 13b96cb9..c93114dc 100644 --- a/utils/resize_kernels.cpp +++ b/utils/resize_kernels.cpp @@ -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 -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(tex_y, x * fx_scale, y * fy_scale) * max_yuv_value), - (YuvUnit)(tex2D(tex_y, (x + 1) * fx_scale, y * fy_scale) * max_yuv_value) - }; - y++; - *(YuvUnitx2 *)(p_dst + y * pitch + x * sizeof(YuvUnit)) = YuvUnitx2 { - (YuvUnit)(tex2D(tex_y, x * fx_scale, y * fy_scale) * max_yuv_value), - (YuvUnit)(tex2D(tex_y, (x + 1) * fx_scale, y * fy_scale) * max_yuv_value) - }; - float2 uv = tex2D(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 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(); - 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(); - 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 <<>>(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 <<>>(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(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(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(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() : hipCreateChannelDesc(); - 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 <<>>(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 <<>>(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 <<>>(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,