diff --git a/projects/rocdecode/samples/videoDecodeRGB/videodecrgb.cpp b/projects/rocdecode/samples/videoDecodeRGB/videodecrgb.cpp index 11655c8760..c36b7b6d22 100644 --- a/projects/rocdecode/samples/videoDecodeRGB/videodecrgb.cpp +++ b/projects/rocdecode/samples/videoDecodeRGB/videodecrgb.cpp @@ -77,75 +77,75 @@ void DumpRGBImage(std::string outputfileName, void* pdevMem, OutputSurfaceInfo * } } -void ColorConvertYUV2RGB(uint8_t *p_src, OutputSurfaceInfo *surf_info, uint8_t *rgb_dev_mem_ptr, OutputFormatEnum e_output_format) { +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(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == bgra) YUV444ToColor32(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 4 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgb) YUV444ToColor24(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgba) YUV444ToColor32(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 4 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + 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(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == bgra) Nv12ToColor32(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 4 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgb) Nv12ToColor24(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgba) Nv12ToColor32(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 4 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + 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(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgb) YUV444P16ToColor24(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == bgr48) YUV444P16ToColor48(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 6 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgb48) YUV444P16ToColor48(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 6 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == bgra64) YUV444P16ToColor64(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 8 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgba64) YUV444P16ToColor64(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 8 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + 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(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgb) P016ToColor24(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 3 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == bgr48) P016ToColor48(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 6 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgb48) P016ToColor48(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 6 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == bgra64) P016ToColor64(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 8 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); else if (e_output_format == rgba64) P016ToColor64(p_src, surf_info->output_pitch, static_cast(rgb_dev_mem_ptr), 8 * rgb_width, surf_info->output_width, - surf_info->output_height, surf_info->output_vstride, 0); + surf_info->output_height, surf_info->output_vstride, 0, hip_stream); } } @@ -281,7 +281,7 @@ int main(int argc, char **argv) { return -1; } } - ColorConvertYUV2RGB(p_frame, surf_info, p_rgb_dev_mem, e_output_format); + ColorConvertYUV2RGB(p_frame, surf_info, p_rgb_dev_mem, e_output_format, viddec.GetStream()); } if (dump_output_frames) { if (convert_to_rgb) diff --git a/projects/rocdecode/utils/colorspace_kernels.cpp b/projects/rocdecode/utils/colorspace_kernels.cpp index 49e35740fa..fbd1ad63bf 100644 --- a/projects/rocdecode/utils/colorspace_kernels.cpp +++ b/projects/rocdecode/utils/colorspace_kernels.cpp @@ -283,209 +283,209 @@ __global__ static void Yuv444ToRgbPlanarKernel(uint8_t *dp_yuv, int yuv_pitch, u } template -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) { +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 - <<>> + <<>> (dp_nv12, nv12_pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_nv12, nv12_pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_p016, p016_pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_p016, p016_pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_nv12, nv12_pitch, dp_bgrp, nBgrpPitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_p016, p016_pitch, dp_bgrp, nBgrpPitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgrp, nBgrpPitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgrp, nBgrpPitch, width, height, v_pitch); } // Explicit Instantiation: for RGB32/BGR32 and RGB64/BGR64 formats -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); -template 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); +template 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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); +template 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); +template 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); +template 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); +template 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); +template 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); +template 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); +template 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); template -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) { +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 - <<>> + <<>> (dp_nv12, nv12_pitch, dp_bgr, bgr_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_nv12, nv12_pitch, dp_bgr, bgr_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgr, bgr_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgr, bgr_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_p016, p016_pitch, dp_bgr, bgr_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_p016, p016_pitch, dp_bgr, bgr_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgra, bgra_pitch, width, height, v_pitch); } template -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) { +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 - <<>> + <<>> (dp_yuv_444, pitch, dp_bgr, bgr_pitch, width, height, v_pitch); } // Explicit Instantiation: for RGB24/BGR24 and RGB48/BGR48 formats -template void Nv12ToColor24(uint8_t *dp_nv12, int nv12_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void Nv12ToColor24(uint8_t *dp_nv12, int nv12_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void Nv12ToColor48(uint8_t *dp_nv12, int nv12_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void Nv12ToColor48(uint8_t *dp_nv12, int nv12_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void YUV444ToColor24(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void YUV444ToColor24(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void YUV444ToColor48(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void YUV444ToColor48(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void P016ToColor24(uint8_t *dp_p016, int p016_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void P016ToColor24(uint8_t *dp_p016, int p016_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void P016ToColor48(uint8_t *dp_p016, int p016_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void P016ToColor48(uint8_t *dp_p016, int p016_pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void YUV444P16ToColor24(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void YUV444P16ToColor24(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void YUV444P16ToColor48(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); -template void YUV444P16ToColor48(uint8_t *dp_yuv_444, int pitch, uint8_t *p_bgr, int p_bgrPitch, int width, int height, int v_pitch, int col_standard); +template void Nv12ToColor24(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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 @@ -539,10 +539,10 @@ __global__ static void RgbaToYuvKernel(uint8_t *dp_rgb, int rgba_pitch, uint8_t }; } -void Bgra64ToP016(uint8_t *dp_bgra, int bgra_pitch, uint8_t *dp_p016, int p016_pitch, int width, int height, int col_standard) { +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 - <<>> + <<>> (dp_bgra, bgra_pitch, dp_p016, p016_pitch, width, height); } @@ -585,9 +585,9 @@ __global__ static void RgbToYuvKernel(uint8_t *dp_rgb, int rgb_pitch, uint8_t *d }; } -void Bgr48ToP016(uint8_t *p_bgr, int bgr_pitch, uint8_t *dp_p016, int p016_pitch, int width, int height, int col_standard) { +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 - <<>> + <<>> (p_bgr, bgr_pitch, dp_p016, p016_pitch, width, height); } diff --git a/projects/rocdecode/utils/colorspace_kernels.h b/projects/rocdecode/utils/colorspace_kernels.h index c600de8582..49ad9cdfb5 100644 --- a/projects/rocdecode/utils/colorspace_kernels.h +++ b/projects/rocdecode/utils/colorspace_kernels.h @@ -108,36 +108,36 @@ union RGBA64 { // color-convert hip kernel function definitions template -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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 -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); +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);