diff --git a/projects/rocjpeg/CHANGELOG.md b/projects/rocjpeg/CHANGELOG.md index 31febf271879..8365c98aa487 100644 --- a/projects/rocjpeg/CHANGELOG.md +++ b/projects/rocjpeg/CHANGELOG.md @@ -3,12 +3,16 @@ Documentation for rocJPEG is available at [https://rocm.docs.amd.com/projects/rocJPEG/en/latest/](https://rocm.docs.amd.com/projects/rocJPEG/en/latest/) -## (unreleased) rocJPEG 1.10.0 +## (unreleased) rocJPEG 1.11.0 ### Changed * Changed the HIP YUV-to-RGB color conversion kernels from the BT.709 standard to BT.601 for color space conversion. +### Added + +* Added HIP color-conversion kernels in the batched decode path + ## rocJPEG 1.9.0 ### Added diff --git a/projects/rocjpeg/CMakeLists.txt b/projects/rocjpeg/CMakeLists.txt index 8804f57ae827..b84ffd41d62f 100644 --- a/projects/rocjpeg/CMakeLists.txt +++ b/projects/rocjpeg/CMakeLists.txt @@ -56,7 +56,7 @@ endif() # rocjpeg Version # NOTE: package version and rocjpeg_version.h is generated with this version -set(VERSION "1.10.0") +set(VERSION "1.11.0") # Set Project Version and Language project(rocjpeg VERSION ${VERSION} LANGUAGES CXX) diff --git a/projects/rocjpeg/src/rocjpeg_decoder.cpp b/projects/rocjpeg/src/rocjpeg_decoder.cpp index ecf77f357bf9..c9815f80dd98 100644 --- a/projects/rocjpeg/src/rocjpeg_decoder.cpp +++ b/projects/rocjpeg/src/rocjpeg_decoder.cpp @@ -260,7 +260,6 @@ RocJpegStatus RocJpegDecoder::FinalizeDecode(VASurfaceID current_surface_id, con is_roi_valid = false; } - switch (decode_params->output_format) { case ROCJPEG_OUTPUT_NATIVE: // Copy the native decoded output buffers from interop memory directly to the destination buffers @@ -318,16 +317,33 @@ RocJpegStatus RocJpegDecoder::FinalizeDecode(VASurfaceID current_surface_id, con */ RocJpegStatus RocJpegDecoder::FinalizeDecodeBatched(const VASurfaceID *surface_ids, const JpegStreamParameters *jpeg_stream_params, const RocJpegDecodeParams *decode_params, RocJpegImage *destinations, int batch_size) { + // Serialize batched finalization: callers (e.g. DecodeBatchedSync) reach here without + // holding mutex_, and this touches shared state (the per-kernel param buffers and + // hip_stream_). A dedicated lock keeps two concurrent finalizations from interleaving + // without holding mutex_ across the GPU work, so async submits can still proceed. + std::lock_guard finalize_lock(finalize_mutex_); VcnJpegSpec current_vcn_jpeg_spec = jpeg_vaapi_decoder_.GetCurrentVcnJpegSpec(); for (int i = 0; i < batch_size; i += current_vcn_jpeg_spec.num_jpeg_cores) { int sub_batch_end = std::min(i + static_cast(current_vcn_jpeg_spec.num_jpeg_cores), batch_size); int current_sub_batch_size = sub_batch_end - i; + // Pass 1: sync all VA surfaces and collect HIP interop device memory for this group. + // Separating sync from kernel launch allows us to build a stable vector of all + // interop structs before filling the batch params — no struct lifetime ambiguity. + std::vector interop_mems(current_sub_batch_size); + for (int k = 0; k < current_sub_batch_size; k++) { + CHECK_ROCJPEG(jpeg_vaapi_decoder_.SyncSurface(surface_ids[i + k])); + CHECK_ROCJPEG(jpeg_vaapi_decoder_.GetHipInteropMem(surface_ids[i + k], interop_mems[k])); + } + + // Pass 2: accumulate batched kernel params per format; memcpy-only work + // (CopyChannel) is issued inline on the stream. One batched launch per kernel + // type follows the loop. + ResetBatchedParams(); + for (int k = 0; k < current_sub_batch_size; k++) { int idx = i + k; - HipInteropDeviceMem hip_interop_dev_mem = {}; - CHECK_ROCJPEG(jpeg_vaapi_decoder_.SyncSurface(surface_ids[idx])); - CHECK_ROCJPEG(jpeg_vaapi_decoder_.GetHipInteropMem(surface_ids[idx], hip_interop_dev_mem)); + HipInteropDeviceMem &mem = interop_mems[k]; uint16_t chroma_height = 0; uint16_t picture_width = 0; @@ -351,37 +367,37 @@ RocJpegStatus RocJpegDecoder::FinalizeDecodeBatched(const VASurfaceID *surface_i switch (decode_params[idx].output_format) { case ROCJPEG_OUTPUT_NATIVE: - CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, picture_height, chroma_height)); - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_width, picture_height, 0, &destinations[idx], &decode_params[idx], is_roi_valid)); - if (hip_interop_dev_mem.surface_format == VA_FOURCC_NV12) { - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_width, chroma_height, 1, &destinations[idx], &decode_params[idx], is_roi_valid)); - } else if (hip_interop_dev_mem.surface_format == VA_FOURCC_444P || - hip_interop_dev_mem.surface_format == VA_FOURCC_422V) { - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_width, chroma_height, 1, &destinations[idx], &decode_params[idx], is_roi_valid)); - CHECK_ROCJPEG(CopyChannel(hip_interop_dev_mem, picture_width, chroma_height, 2, &destinations[idx], &decode_params[idx], is_roi_valid)); + CHECK_ROCJPEG(GetChromaHeight(mem.surface_format, picture_height, chroma_height)); + CHECK_ROCJPEG(CopyChannel(mem, picture_width, picture_height, 0, &destinations[idx], &decode_params[idx], is_roi_valid)); + if (mem.surface_format == VA_FOURCC_NV12) { + CHECK_ROCJPEG(CopyChannel(mem, picture_width, chroma_height, 1, &destinations[idx], &decode_params[idx], is_roi_valid)); + } else if (mem.surface_format == VA_FOURCC_444P || + mem.surface_format == VA_FOURCC_422V) { + CHECK_ROCJPEG(CopyChannel(mem, picture_width, chroma_height, 1, &destinations[idx], &decode_params[idx], is_roi_valid)); + CHECK_ROCJPEG(CopyChannel(mem, picture_width, chroma_height, 2, &destinations[idx], &decode_params[idx], is_roi_valid)); } break; case ROCJPEG_OUTPUT_YUV_PLANAR: - CHECK_ROCJPEG(GetChromaHeight(hip_interop_dev_mem.surface_format, picture_height, chroma_height)); - CHECK_ROCJPEG(GetPlanarYUVOutputFormat(hip_interop_dev_mem, picture_width, - picture_height, chroma_height, &destinations[idx], &decode_params[idx], is_roi_valid)); + CHECK_ROCJPEG(GetChromaHeight(mem.surface_format, picture_height, chroma_height)); + CHECK_ROCJPEG(AccumulatePlanarYUVOutputFormat(mem, picture_width, picture_height, chroma_height, &destinations[idx], &decode_params[idx], is_roi_valid)); break; case ROCJPEG_OUTPUT_Y: - CHECK_ROCJPEG(GetYOutputFormat(hip_interop_dev_mem, picture_width, - picture_height, &destinations[idx], &decode_params[idx], is_roi_valid)); + CHECK_ROCJPEG(AccumulateYOutputFormat(mem, picture_width, picture_height, &destinations[idx], &decode_params[idx], is_roi_valid)); break; case ROCJPEG_OUTPUT_RGB: - CHECK_ROCJPEG(ColorConvertToRGB(hip_interop_dev_mem, picture_width, - picture_height, &destinations[idx], &decode_params[idx], is_roi_valid)); + CHECK_ROCJPEG(AccumulateColorConvertToRGB(mem, picture_width, picture_height, &destinations[idx], &decode_params[idx], is_roi_valid)); break; case ROCJPEG_OUTPUT_RGB_PLANAR: - CHECK_ROCJPEG(ColorConvertToRGBPlanar(hip_interop_dev_mem, picture_width, - picture_height, &destinations[idx], &decode_params[idx], is_roi_valid)); + CHECK_ROCJPEG(AccumulateColorConvertToRGBPlanar(mem, picture_width, picture_height, &destinations[idx], &decode_params[idx], is_roi_valid)); break; default: break; } } + + // One batched launch per kernel type for everything collected in this group. + CHECK_ROCJPEG(LaunchBatchedParams()); + CHECK_HIP(hipStreamSynchronize(hip_stream_)); for (int k = 0; k < current_sub_batch_size; k++) { jpeg_vaapi_decoder_.SetSurfaceAsIdle(surface_ids[i + k]); @@ -452,6 +468,7 @@ RocJpegStatus RocJpegDecoder::DecodeBatched(RocJpegStreamHandle *jpeg_streams, i FunctionExitLog(g_rocjpeg_logger); return ROCJPEG_STATUS_SUCCESS; } + /** * @brief Submits a batch of JPEG decode operations and stores pending state for all images. */ @@ -794,6 +811,12 @@ RocJpegStatus RocJpegDecoder::GetChromaHeight(uint32_t surface_format, uint16_t * is successful. Returns ROCJPEG_STATUS_JPEG_NOT_SUPPORTED if the surface format is not supported. */ RocJpegStatus RocJpegDecoder::ColorConvertToRGB(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + // roi_offset is the generic ROI byte offset, computed against the luma pitch. It is applied to the + // Y component of every supported YUV surface format, and also to the U/V planes of the formats whose + // chroma is neither vertically nor horizontally subsampled (e.g. YUV444 / VA_FOURCC_444P). + // roi_uv_offset is the chroma-specific ROI byte offset, computed against the chroma pitch with the + // row index halved. It is applied only to the U/V data of the vertically-subsampled formats: + // YUV440 (VA_FOURCC_422V surface) and YUV420 (VA_FOURCC_NV12 surface). uint32_t roi_offset = 0; uint32_t roi_uv_offset = 0; int16_t top = decode_params->crop_rectangle.top; @@ -809,11 +832,15 @@ RocJpegStatus RocJpegDecoder::ColorConvertToRGB(HipInteropDeviceMem& hip_interop switch (hip_interop_dev_mem.surface_format) { case VA_FOURCC_444P: ColorConvertYUV444ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1] + roi_offset, hip_interop_dev_mem.offset[2] + roi_offset); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1] + roi_offset, + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[2] + roi_offset); break; case VA_FOURCC_422V: ColorConvertYUV440ToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1] /*+ roi_uv_offset*/, hip_interop_dev_mem.offset[2] /*+ roi_uv_offset*/); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1] + roi_uv_offset, + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[2] + roi_uv_offset); break; case VA_FOURCC_YUY2: ColorConvertYUYVToRGB(hip_stream_, picture_width, picture_height, destination->channel[0], destination->pitch[0], @@ -855,6 +882,12 @@ RocJpegStatus RocJpegDecoder::ColorConvertToRGB(HipInteropDeviceMem& hip_interop * Returns ROCJPEG_STATUS_JPEG_NOT_SUPPORTED if the surface format is not supported. */ RocJpegStatus RocJpegDecoder::ColorConvertToRGBPlanar(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + // roi_offset is the generic ROI byte offset, computed against the luma pitch. It is applied to the + // Y component of every supported YUV surface format, and also to the U/V planes of the formats whose + // chroma is neither vertically nor horizontally subsampled (e.g. YUV444 / VA_FOURCC_444P). + // roi_uv_offset is the chroma-specific ROI byte offset, computed against the chroma pitch with the + // row index halved. It is applied only to the U/V data of the vertically-subsampled formats: + // YUV440 (VA_FOURCC_422V surface) and YUV420 (VA_FOURCC_NV12 surface). uint32_t roi_offset = 0; uint32_t roi_uv_offset = 0; int16_t top = decode_params->crop_rectangle.top; @@ -870,11 +903,15 @@ RocJpegStatus RocJpegDecoder::ColorConvertToRGBPlanar(HipInteropDeviceMem& hip_i switch (hip_interop_dev_mem.surface_format) { case VA_FOURCC_444P: ColorConvertYUV444ToRGBPlanar(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1] + roi_offset, hip_interop_dev_mem.offset[2] + roi_offset); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1] + roi_offset, + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[2] + roi_offset); break; case VA_FOURCC_422V: ColorConvertYUV440ToRGBPlanar(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], destination->pitch[0], - hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], hip_interop_dev_mem.offset[1] /*+ roi_uv_offset*/, hip_interop_dev_mem.offset[2] /*+ roi_uv_offset*/); + hip_interop_dev_mem.hip_mapped_device_mem + roi_offset, hip_interop_dev_mem.pitch[0], + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[1] + roi_uv_offset, + hip_interop_dev_mem.hip_mapped_device_mem + hip_interop_dev_mem.offset[2] + roi_uv_offset); break; case VA_FOURCC_YUY2: ColorConvertYUYVToRGBPlanar(hip_stream_, picture_width, picture_height, destination->channel[0], destination->channel[1], destination->channel[2], destination->pitch[0], @@ -920,6 +957,10 @@ RocJpegStatus RocJpegDecoder::ColorConvertToRGBPlanar(HipInteropDeviceMem& hip_i * @return The status of the operation. Returns ROCJPEG_STATUS_SUCCESS if successful. */ RocJpegStatus RocJpegDecoder::GetPlanarYUVOutputFormat(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + // Only one ROI offset is needed here, and which plane it addresses depends on the surface format: + // for YUV420 (VA_FOURCC_NV12) it is the chroma offset into the interleaved UV plane (row index halved, + // chroma pitch) since the luma plane is handled by CopyChannel; for YUY2 it is the offset into the + // packed plane (left doubled, because each YUYV pixel occupies 2 bytes). uint32_t roi_offset = 0; if (is_roi_valid) { int16_t top = decode_params->crop_rectangle.top; @@ -964,7 +1005,10 @@ RocJpegStatus RocJpegDecoder::GetPlanarYUVOutputFormat(HipInteropDeviceMem& hip_ * @return The status of the operation. Returns ROCJPEG_STATUS_SUCCESS if successful. */ RocJpegStatus RocJpegDecoder::GetYOutputFormat(HipInteropDeviceMem& hip_interop_dev_mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { - uint32_t roi_offset = 0; + // Only the Y component is produced here, so a single ROI offset into the source luma data suffices. + // It is used for YUY2 only (left doubled, because each YUYV pixel occupies 2 bytes); every other + // format takes the CopyChannel path, which computes its own offset. + uint32_t roi_offset = 0; if (hip_interop_dev_mem.surface_format == VA_FOURCC_YUY2) { // calculate offset and add to hip_mapped_device_mem if (is_roi_valid) { @@ -980,3 +1024,281 @@ RocJpegStatus RocJpegDecoder::GetYOutputFormat(HipInteropDeviceMem& hip_interop_ } return ROCJPEG_STATUS_SUCCESS; } + +/** + * @brief Clears every batched-param scratch buffer at the start of a group. + */ +void RocJpegDecoder::ResetBatchedParams() { + b_rgba_rgb_.Reset(); + b_yuyv_yuvp_.Reset(); + b_nv12_uvp_.Reset(); + b_yuyv_y_.Reset(); + b_yuv_rgb_.Reset(); + b_yuv_rgbp_.Reset(); +} + +/** + * @brief Uploads and launches one batched kernel per non-empty scratch buffer. + */ +RocJpegStatus RocJpegDecoder::LaunchBatchedParams() { + CHECK_ROCJPEG(LaunchBatchedBuffer(b_rgba_rgb_, ColorConvertRGBAToRGBBatched)); + CHECK_ROCJPEG(LaunchBatchedBuffer(b_yuyv_yuvp_, ConvertPackedYUYVToPlanarYUVBatched)); + CHECK_ROCJPEG(LaunchBatchedBuffer(b_nv12_uvp_, ConvertInterleavedUVToPlanarUVBatched)); + CHECK_ROCJPEG(LaunchBatchedBuffer(b_yuyv_y_, ExtractYFromPackedYUYVBatched)); + CHECK_ROCJPEG(LaunchBatchedBuffer(b_yuv_rgb_, ColorConvertYUVToRGBBatched)); + CHECK_ROCJPEG(LaunchBatchedBuffer(b_yuv_rgbp_, ColorConvertYUVToRGBPlanarBatched)); + return ROCJPEG_STATUS_SUCCESS; +} + +/** + * @brief Batched counterpart of ColorConvertToRGB: appends this image's params to + * the appropriate packed-RGB kernel buffer instead of launching per image. + * Mirrors the pointer/offset arithmetic of ColorConvertToRGB exactly. + */ +RocJpegStatus RocJpegDecoder::AccumulateColorConvertToRGB(HipInteropDeviceMem& mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + // roi_offset is the generic ROI byte offset, computed against the luma pitch. It is applied to the + // Y component of every supported YUV surface format, and also to the U/V planes of the formats whose + // chroma is neither vertically nor horizontally subsampled (e.g. YUV444 / VA_FOURCC_444P). + // roi_uv_offset is the chroma-specific ROI byte offset, computed against the chroma pitch with the + // row index halved. It is applied only to the U/V data of the vertically-subsampled formats: + // YUV440 (VA_FOURCC_422V surface) and YUV420 (VA_FOURCC_NV12 surface). + uint32_t roi_offset = 0; + uint32_t roi_uv_offset = 0; + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left; + if (is_roi_valid) { + if (mem.surface_format == VA_FOURCC_422V || mem.surface_format == VA_FOURCC_NV12) { + roi_uv_offset = (top >> 1) * mem.pitch[1] + left; + } else if (mem.surface_format == VA_FOURCC_YUY2) { + left *= 2; + } + roi_offset = top * mem.pitch[0] + left; + } + // Each thread of the unified kernel converts an 8-pixel-wide by 2-row block. + uint32_t dst_width_in_8px_blocks = (picture_width + 7) / 8; + uint32_t dst_height_in_2row_blocks = (picture_height + 1) / 2; + // The five YUV surface formats share one unified batched kernel; only the + // per-image source layout (format tag + which src pointers are set) differs. + if (mem.surface_format == VA_FOURCC_444P || mem.surface_format == VA_FOURCC_422V || + mem.surface_format == VA_FOURCC_YUY2 || mem.surface_format == VA_FOURCC_NV12 || + mem.surface_format == VA_FOURCC_Y800) { + YUVToRGBBatchParams bp{}; + bp.dst_image = destination->channel[0]; + bp.dst_image_stride_in_bytes = destination->pitch[0]; + bp.dst_image_row_pair_stride_in_bytes = destination->pitch[0] * 2; + bp.src_y_image = mem.hip_mapped_device_mem + roi_offset; + bp.src_y_image_stride_in_bytes = mem.pitch[0]; + bp.src_y_image_row_pair_stride_in_bytes = mem.pitch[0] * 2; + bp.dst_width_in_8px_blocks = dst_width_in_8px_blocks; + bp.dst_height_in_2row_blocks = dst_height_in_2row_blocks; + switch (mem.surface_format) { + case VA_FOURCC_444P: + bp.layout = YUV_LAYOUT_YUV444; + bp.src_u_image = mem.hip_mapped_device_mem + mem.offset[1] + roi_offset; + bp.src_v_image = mem.hip_mapped_device_mem + mem.offset[2] + roi_offset; + break; + case VA_FOURCC_422V: + bp.layout = YUV_LAYOUT_YUV440; + bp.src_u_image = mem.hip_mapped_device_mem + mem.offset[1] + roi_uv_offset; + bp.src_v_image = mem.hip_mapped_device_mem + mem.offset[2] + roi_uv_offset; + break; + case VA_FOURCC_YUY2: + bp.layout = YUV_LAYOUT_YUYV; + break; + case VA_FOURCC_NV12: + bp.layout = YUV_LAYOUT_NV12; + bp.src_chroma_image = mem.hip_mapped_device_mem + mem.offset[1] + roi_uv_offset; + bp.src_chroma_image_stride_in_bytes = mem.pitch[1]; + break; + case VA_FOURCC_Y800: + bp.layout = YUV_LAYOUT_YUV400; + break; + } + b_yuv_rgb_.Add(bp, dst_width_in_8px_blocks, dst_height_in_2row_blocks); + return ROCJPEG_STATUS_SUCCESS; + } + switch (mem.surface_format) { + case VA_FOURCC_RGBA: { + RGBAToRGBBatchParams bp; + bp.dst_width = picture_width; + bp.dst_height = picture_height; + bp.dst_image = destination->channel[0]; + bp.dst_image_stride_in_bytes = destination->pitch[0]; + bp.src_image = mem.hip_mapped_device_mem + roi_offset; + bp.src_image_stride_in_bytes = mem.pitch[0]; + b_rgba_rgb_.Add(bp, (picture_width + 7) >> 3, picture_height); + break; + } + default: + ErrorLog(g_rocjpeg_logger, "Surface format is not supported!"); + return ROCJPEG_STATUS_JPEG_NOT_SUPPORTED; + } + return ROCJPEG_STATUS_SUCCESS; +} + +/** + * @brief Batched counterpart of ColorConvertToRGBPlanar. RGBP (already-planar RGB) + * stays a per-image memcpy; all convertible formats accumulate kernel params. + */ +RocJpegStatus RocJpegDecoder::AccumulateColorConvertToRGBPlanar(HipInteropDeviceMem& mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + // roi_offset is the generic ROI byte offset, computed against the luma pitch. It is applied to the + // Y component of every supported YUV surface format, and also to the U/V planes of the formats whose + // chroma is neither vertically nor horizontally subsampled (e.g. YUV444 / VA_FOURCC_444P). + // roi_uv_offset is the chroma-specific ROI byte offset, computed against the chroma pitch with the + // row index halved. It is applied only to the U/V data of the vertically-subsampled formats: + // YUV440 (VA_FOURCC_422V surface) and YUV420 (VA_FOURCC_NV12 surface). + uint32_t roi_offset = 0; + uint32_t roi_uv_offset = 0; + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left; + if (is_roi_valid) { + if (mem.surface_format == VA_FOURCC_422V || mem.surface_format == VA_FOURCC_NV12) { + roi_uv_offset = (top >> 1) * mem.pitch[1] + left; + } else if (mem.surface_format == VA_FOURCC_YUY2) { + left *= 2; + } + roi_offset = top * mem.pitch[0] + left; + } + // Each thread of the unified kernel converts an 8-pixel-wide by 2-row block. + uint32_t dst_width_in_8px_blocks = (picture_width + 7) / 8; + uint32_t dst_height_in_2row_blocks = (picture_height + 1) / 2; + // Same five YUV formats share one unified batched planar kernel. + if (mem.surface_format == VA_FOURCC_444P || mem.surface_format == VA_FOURCC_422V || + mem.surface_format == VA_FOURCC_YUY2 || mem.surface_format == VA_FOURCC_NV12 || + mem.surface_format == VA_FOURCC_Y800) { + YUVToRGBPlanarBatchParams bp{}; + bp.dst_image_r = destination->channel[0]; + bp.dst_image_g = destination->channel[1]; + bp.dst_image_b = destination->channel[2]; + bp.dst_image_stride_in_bytes = destination->pitch[0]; + bp.dst_image_row_pair_stride_in_bytes = destination->pitch[0] * 2; + bp.src_y_image = mem.hip_mapped_device_mem + roi_offset; + bp.src_y_image_stride_in_bytes = mem.pitch[0]; + bp.src_y_image_row_pair_stride_in_bytes = mem.pitch[0] * 2; + bp.dst_width_in_8px_blocks = dst_width_in_8px_blocks; + bp.dst_height_in_2row_blocks = dst_height_in_2row_blocks; + switch (mem.surface_format) { + case VA_FOURCC_444P: + bp.layout = YUV_LAYOUT_YUV444; + bp.src_u_image = mem.hip_mapped_device_mem + mem.offset[1] + roi_offset; + bp.src_v_image = mem.hip_mapped_device_mem + mem.offset[2] + roi_offset; + break; + case VA_FOURCC_422V: + bp.layout = YUV_LAYOUT_YUV440; + bp.src_u_image = mem.hip_mapped_device_mem + mem.offset[1] + roi_uv_offset; + bp.src_v_image = mem.hip_mapped_device_mem + mem.offset[2] + roi_uv_offset; + break; + case VA_FOURCC_YUY2: + bp.layout = YUV_LAYOUT_YUYV; + break; + case VA_FOURCC_NV12: + bp.layout = YUV_LAYOUT_NV12; + bp.src_chroma_image = mem.hip_mapped_device_mem + mem.offset[1] + roi_uv_offset; + bp.src_chroma_image_stride_in_bytes = mem.pitch[1]; + break; + case VA_FOURCC_Y800: + bp.layout = YUV_LAYOUT_YUV400; + break; + } + b_yuv_rgbp_.Add(bp, dst_width_in_8px_blocks, dst_height_in_2row_blocks); + return ROCJPEG_STATUS_SUCCESS; + } + switch (mem.surface_format) { + case VA_FOURCC_RGBP: + // Already-planar RGB: three plane copies, no kernel. + for (uint8_t channel_index = 0; channel_index < 3; channel_index++) { + CHECK_ROCJPEG(CopyChannel(mem, picture_width, picture_height, channel_index, destination, decode_params, is_roi_valid)); + } + break; + default: + ErrorLog(g_rocjpeg_logger, "Surface format is not supported!"); + return ROCJPEG_STATUS_JPEG_NOT_SUPPORTED; + } + return ROCJPEG_STATUS_SUCCESS; +} + +/** + * @brief Batched counterpart of GetPlanarYUVOutputFormat. The luma/chroma plane + * copies stay per-image memcpys; the YUY2→planar and NV12 interleaved-UV + * conversions accumulate kernel params. + */ +RocJpegStatus RocJpegDecoder::AccumulatePlanarYUVOutputFormat(HipInteropDeviceMem& mem, uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + // Only one ROI offset is needed here, and which plane it addresses depends on the surface format: + // for YUV420 (VA_FOURCC_NV12) it is the chroma offset into the interleaved UV plane (row index halved, + // chroma pitch) since the luma plane is handled by CopyChannel; for YUY2 it is the offset into the + // packed plane (left doubled, because each YUYV pixel occupies 2 bytes). + uint32_t roi_offset = 0; + if (is_roi_valid) { + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left; + if (mem.surface_format == VA_FOURCC_NV12) { + roi_offset = (top >> 1) * mem.pitch[1] + left; + } else if (mem.surface_format == VA_FOURCC_YUY2) { + roi_offset = top * mem.pitch[0] + (left * 2); + } + } + if (mem.surface_format == VA_FOURCC_YUY2) { + PackedYUYVToPlanarYUVBatchParams bp; + bp.dst_height = picture_height; + bp.destination_y = destination->channel[0]; + bp.destination_u = destination->channel[1]; + bp.destination_v = destination->channel[2]; + bp.dst_luma_stride_in_bytes = destination->pitch[0]; + bp.dst_chroma_stride_in_bytes = destination->pitch[1]; + bp.src_image = mem.hip_mapped_device_mem + roi_offset; + bp.src_image_stride_in_bytes = mem.pitch[0]; + bp.dst_width_in_8px_blocks = (picture_width + 7) / 8; + b_yuyv_yuvp_.Add(bp, (picture_width + 7) >> 3, picture_height); + } else { + // Copy Luma + CHECK_ROCJPEG(CopyChannel(mem, picture_width, picture_height, 0, destination, decode_params, is_roi_valid)); + if (mem.surface_format == VA_FOURCC_NV12) { + uint32_t uv_width = picture_width >> 1; + uint32_t uv_height = picture_height >> 1; + InterleavedUVToPlanarUVBatchParams bp; + bp.dst_width = uv_width; + bp.dst_height = uv_height; + bp.dst_image1 = destination->channel[1]; + bp.dst_image2 = destination->channel[2]; + bp.dst_image_stride_in_bytes = destination->pitch[1]; + bp.src_image = mem.hip_mapped_device_mem + mem.offset[1] + roi_offset; + bp.src_image_stride_in_bytes = mem.pitch[1]; + b_nv12_uvp_.Add(bp, (uv_width + 7) >> 3, uv_height); + } else if (mem.surface_format == VA_FOURCC_444P || + mem.surface_format == VA_FOURCC_422V) { + CHECK_ROCJPEG(CopyChannel(mem, picture_width, chroma_height, 1, destination, decode_params, is_roi_valid)); + CHECK_ROCJPEG(CopyChannel(mem, picture_width, chroma_height, 2, destination, decode_params, is_roi_valid)); + } + } + return ROCJPEG_STATUS_SUCCESS; +} + +/** + * @brief Batched counterpart of GetYOutputFormat. YUY2 accumulates the Y-extract + * kernel; other formats copy the luma plane per image. + */ +RocJpegStatus RocJpegDecoder::AccumulateYOutputFormat(HipInteropDeviceMem& mem, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid) { + // Only the Y component is produced here, so a single ROI offset into the source luma data suffices. + // It is used for YUY2 only (left doubled, because each YUYV pixel occupies 2 bytes); every other + // format takes the CopyChannel path, which computes its own offset. + uint32_t roi_offset = 0; + if (mem.surface_format == VA_FOURCC_YUY2) { + if (is_roi_valid) { + int16_t top = decode_params->crop_rectangle.top; + int16_t left = decode_params->crop_rectangle.left * 2; + roi_offset = top * mem.pitch[0] + left; + } + YFromPackedYUYVBatchParams bp; + bp.dst_height = picture_height; + bp.destination_y = destination->channel[0]; + bp.dst_luma_stride_in_bytes = destination->pitch[0]; + bp.src_image = mem.hip_mapped_device_mem + roi_offset; + bp.src_image_stride_in_bytes = mem.pitch[0]; + bp.dst_width_in_8px_blocks = (picture_width + 7) / 8; + b_yuyv_y_.Add(bp, (picture_width + 7) >> 3, picture_height); + } else { + // Copy Luma + CHECK_ROCJPEG(CopyChannel(mem, picture_width, picture_height, 0, destination, decode_params, is_roi_valid)); + } + return ROCJPEG_STATUS_SUCCESS; +} diff --git a/projects/rocjpeg/src/rocjpeg_decoder.h b/projects/rocjpeg/src/rocjpeg_decoder.h index dd8ca57c30c2..65ac42afc908 100644 --- a/projects/rocjpeg/src/rocjpeg_decoder.h +++ b/projects/rocjpeg/src/rocjpeg_decoder.h @@ -34,6 +34,61 @@ THE SOFTWARE. #include "rocjpeg_vaapi_decoder.h" #include "rocjpeg_hip_kernels.h" +/** + * @brief Reusable scratch buffer for one batched kernel's per-image parameter array. + * + * Accumulates a host vector of per-image parameter structs for a single group, + * tracks the maximum grid extents across the group, and lazily grows a device + * buffer that the params are uploaded into before the batched launch. + */ +template +class BatchedKernelParams { +public: + BatchedKernelParams() = default; + ~BatchedKernelParams() { if (dev_) { (void)hipFree(dev_); } } + + // Owns a raw device buffer; copying/moving would double-free it. + BatchedKernelParams(const BatchedKernelParams &) = delete; + BatchedKernelParams &operator=(const BatchedKernelParams &) = delete; + BatchedKernelParams(BatchedKernelParams &&) = delete; + BatchedKernelParams &operator=(BatchedKernelParams &&) = delete; + + void Reset() { host_.clear(); max_gx_ = 0; max_gy_ = 0; } + + // Append one image's params; gx/gy are that image's pre-block-division grid extents. + void Add(const T ¶ms, uint32_t gx, uint32_t gy) { + host_.push_back(params); + if (gx > max_gx_) max_gx_ = gx; + if (gy > max_gy_) max_gy_ = gy; + } + + bool Empty() const { return host_.empty(); } + uint32_t Count() const { return static_cast(host_.size()); } + uint32_t MaxGx() const { return max_gx_; } + uint32_t MaxGy() const { return max_gy_; } + + // Grow the device buffer on demand and upload the host params on the given stream. + hipError_t Upload(hipStream_t stream, T **dev_out) { + size_t bytes = host_.size() * sizeof(T); + if (host_.size() > capacity_) { + if (dev_) { (void)hipFree(dev_); dev_ = nullptr; } + capacity_ = 0; + hipError_t status = hipMalloc(reinterpret_cast(&dev_), bytes); + if (status != hipSuccess) return status; + capacity_ = host_.size(); + } + *dev_out = dev_; + return hipMemcpyAsync(dev_, host_.data(), bytes, hipMemcpyHostToDevice, stream); + } + +private: + std::vector host_; + T *dev_ = nullptr; + size_t capacity_ = 0; + uint32_t max_gx_ = 0; + uint32_t max_gy_ = 0; +}; + /** * @class RocJpegDecoder * @brief The RocJpegDecoder class represents a JPEG decoder. @@ -214,14 +269,47 @@ class RocJpegDecoder { */ RocJpegStatus GetYOutputFormat(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); + // Batched-path accumulators: mirror the per-image color-convert/output helpers above, + // but instead of launching a kernel per image they append params to the per-kernel + // scratch buffers below. Memcpy-only work (CopyChannel) is still issued inline. + RocJpegStatus AccumulateColorConvertToRGB(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); + RocJpegStatus AccumulateColorConvertToRGBPlanar(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); + RocJpegStatus AccumulatePlanarYUVOutputFormat(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, uint16_t chroma_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); + RocJpegStatus AccumulateYOutputFormat(HipInteropDeviceMem& hip_interop, uint32_t picture_width, uint32_t picture_height, RocJpegImage *destination, const RocJpegDecodeParams *decode_params, bool is_roi_valid); + void ResetBatchedParams(); + RocJpegStatus LaunchBatchedParams(); + + // Uploads one accumulated buffer (if non-empty) and issues its batched launch. + template + RocJpegStatus LaunchBatchedBuffer(BatchedKernelParams &buf, + void (*launch)(hipStream_t, uint32_t, uint32_t, const T *, uint32_t)) { + if (buf.Empty()) return ROCJPEG_STATUS_SUCCESS; + T *dev = nullptr; + CHECK_HIP(buf.Upload(hip_stream_, &dev)); + launch(hip_stream_, buf.MaxGx(), buf.MaxGy(), dev, buf.Count()); + return ROCJPEG_STATUS_SUCCESS; + } + int num_devices_; // Number of available devices int device_id_; // ID of the device to be used hipDeviceProp_t hip_dev_prop_; // HIP device properties hipStream_t hip_stream_; // HIP stream std::mutex mutex_; // Mutex for thread safety + std::mutex finalize_mutex_; // Serializes FinalizeDecodeBatched (shared param buffers + hip_stream_) across concurrent syncs RocJpegBackend backend_; // RocJpeg backend RocJpegVappiDecoder jpeg_vaapi_decoder_; // RocJpeg VAAPI decoder object std::unordered_map pending_decodes_; // Map of pending asynchronous decodes keyed by destination + + // Per-kernel scratch buffers for the batched output path, one per batched kernel. + // Filled per group in FinalizeDecodeBatched; each device buffer grows on demand. + BatchedKernelParams b_rgba_rgb_; + BatchedKernelParams b_yuyv_yuvp_; + BatchedKernelParams b_nv12_uvp_; + BatchedKernelParams b_yuyv_y_; + // Unified YUV->RGB / YUV->RGB_PLANAR buffers: a single kernel each handles all of + // NV12/YUV444/YUV440/YUYV/YUV400, selected per image by a surface-layout tag. + BatchedKernelParams b_yuv_rgb_; + BatchedKernelParams b_yuv_rgbp_; }; #endif //ROC_JPEG_DECODER_H_ \ No newline at end of file diff --git a/projects/rocjpeg/src/rocjpeg_hip_kernels.cpp b/projects/rocjpeg/src/rocjpeg_hip_kernels.cpp index 27869c6123a0..67afd1f3adfb 100644 --- a/projects/rocjpeg/src/rocjpeg_hip_kernels.cpp +++ b/projects/rocjpeg/src/rocjpeg_hip_kernels.cpp @@ -60,7 +60,6 @@ __global__ void ColorConvertYUV444ToRGBKernel(uint8_t *dst_image, uint32_t dst_i uint32_t src_y0_idx = y * src_yuv_image_stride_in_bytes_comp + (x << 3); uint32_t src_y1_idx = src_y0_idx + src_yuv_image_stride_in_bytes; - uint2 y0 = *((uint2 *)(&src_y_image[src_y0_idx])); uint2 y1 = *((uint2 *)(&src_y_image[src_y1_idx])); @@ -233,11 +232,12 @@ __global__ void ColorConvertYUV444ToRGBKernel(uint8_t *dst_image, uint32_t dst_i * @param dst_image_stride_in_bytes The stride (in bytes) of the destination RGB image buffer. * @param src_yuv_image Pointer to the source YUV444 image buffer. * @param src_yuv_image_stride_in_bytes The stride (in bytes) of the source YUV444 image buffer. - * @param src_u_image_offset The offset (in bytes) to the U component in the source YUV444 image buffer. + * @param src_u_image Pointer to the U plane of the source YUV444 image (already ROI-adjusted). + * @param src_v_image Pointer to the V plane of the source YUV444 image (already ROI-adjusted). */ void ColorConvertYUV444ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset) { + uint32_t src_yuv_image_stride_in_bytes, const uint8_t *src_u_image, const uint8_t *src_v_image) { int32_t local_threads_x = 16; int32_t local_threads_y = 4; @@ -251,8 +251,8 @@ void ColorConvertYUV444ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t ds ColorConvertYUV444ToRGBKernel<<(global_threads_x) / local_threads_x), ceil(static_cast(global_threads_y) / local_threads_y)), dim3(local_threads_x, local_threads_y), 0, stream>>>(dst_image, - dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_yuv_image + src_u_image_offset, - src_yuv_image + src_v_image_offset, src_yuv_image_stride_in_bytes, + dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_u_image, + src_v_image, src_yuv_image_stride_in_bytes, dst_width_comp, dst_height_comp, src_yuv_image_stride_in_bytes_comp); } @@ -267,7 +267,6 @@ __global__ void ColorConvertYUV444ToRGBPlanarKernel(uint8_t *dst_image_r, uint8_ uint32_t src_y0_idx = y * src_yuv_image_stride_in_bytes_comp + (x << 3); uint32_t src_y1_idx = src_y0_idx + src_yuv_image_stride_in_bytes; - uint2 y0 = *((uint2 *)(&src_y_image[src_y0_idx])); uint2 y1 = *((uint2 *)(&src_y_image[src_y1_idx])); @@ -448,7 +447,6 @@ __global__ void ColorConvertYUV444ToRGBPlanarKernel(uint8_t *dst_image_r, uint8_ } } - /** * @brief Converts YUV444 image to RGB planar format. * @@ -463,11 +461,12 @@ __global__ void ColorConvertYUV444ToRGBPlanarKernel(uint8_t *dst_image_r, uint8_ * @param dst_image_stride_in_bytes The stride (in bytes) of the destination image. * @param src_yuv_image Pointer to the source YUV image. * @param src_yuv_image_stride_in_bytes The stride (in bytes) of the source YUV image. - * @param src_u_image_offset The offset (in bytes) to the U channel in the source YUV image. + * @param src_u_image Pointer to the U plane of the source YUV image (already ROI-adjusted). + * @param src_v_image Pointer to the V plane of the source YUV image (already ROI-adjusted). */ void ColorConvertYUV444ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image_r, uint8_t *dst_image_g, uint8_t *dst_image_b, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset) { + uint32_t src_yuv_image_stride_in_bytes, const uint8_t *src_u_image, const uint8_t *src_v_image) { int32_t local_threads_x = 16; int32_t local_threads_y = 4; @@ -481,8 +480,8 @@ void ColorConvertYUV444ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint3 ColorConvertYUV444ToRGBPlanarKernel<<(global_threads_x) / local_threads_x), ceil(static_cast(global_threads_y) / local_threads_y)), dim3(local_threads_x, local_threads_y), 0, stream>>>(dst_image_r, dst_image_g, dst_image_b, - dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_yuv_image + src_u_image_offset, - src_yuv_image + src_v_image_offset, src_yuv_image_stride_in_bytes, + dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_u_image, + src_v_image, src_yuv_image_stride_in_bytes, dst_width_comp, dst_height_comp, src_yuv_image_stride_in_bytes_comp); } @@ -667,11 +666,12 @@ __global__ void ColorConvertYUV440ToRGBKernel(uint8_t *dst_image, uint32_t dst_i * @param dst_image_stride_in_bytes The stride (in bytes) of the destination RGB image buffer. * @param src_yuv_image Pointer to the source YUV440 image buffer. * @param src_yuv_image_stride_in_bytes The stride (in bytes) of the source YUV440 image buffer. - * @param src_u_image_offset The offset (in bytes) to the U component in the source YUV440 image buffer. + * @param src_u_image Pointer to the U plane of the source YUV440 image (already ROI-adjusted). + * @param src_v_image Pointer to the V plane of the source YUV440 image (already ROI-adjusted). */ void ColorConvertYUV440ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset) { + uint32_t src_yuv_image_stride_in_bytes, const uint8_t *src_u_image, const uint8_t *src_v_image) { int32_t local_threads_x = 16; int32_t local_threads_y = 4; @@ -685,8 +685,8 @@ void ColorConvertYUV440ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t ds ColorConvertYUV440ToRGBKernel<<(global_threads_x) / local_threads_x), ceil(static_cast(global_threads_y) / local_threads_y)), dim3(local_threads_x, local_threads_y), 0, stream>>>(dst_image, - dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_yuv_image + src_u_image_offset, - src_yuv_image + src_v_image_offset, src_yuv_image_stride_in_bytes, + dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_u_image, + src_v_image, src_yuv_image_stride_in_bytes, dst_width_comp, dst_height_comp, src_yuv_image_stride_in_bytes_comp); } @@ -879,7 +879,6 @@ __global__ void ColorConvertYUV440ToRGBPlanarKernel(uint8_t *dst_image_r, uint8_ } } - /** * @brief Converts YUV440 image to RGB planar format. * @@ -894,11 +893,12 @@ __global__ void ColorConvertYUV440ToRGBPlanarKernel(uint8_t *dst_image_r, uint8_ * @param dst_image_stride_in_bytes The stride (in bytes) of the destination image. * @param src_yuv_image Pointer to the source YUV image. * @param src_yuv_image_stride_in_bytes The stride (in bytes) of the source YUV image. - * @param src_u_image_offset The offset (in bytes) to the U channel in the source YUV image. + * @param src_u_image Pointer to the U plane of the source YUV image (already ROI-adjusted). + * @param src_v_image Pointer to the V plane of the source YUV image (already ROI-adjusted). */ void ColorConvertYUV440ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image_r, uint8_t *dst_image_g, uint8_t *dst_image_b, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset) { + uint32_t src_yuv_image_stride_in_bytes, const uint8_t *src_u_image, const uint8_t *src_v_image) { int32_t local_threads_x = 16; int32_t local_threads_y = 4; @@ -912,8 +912,8 @@ void ColorConvertYUV440ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint3 ColorConvertYUV440ToRGBPlanarKernel<<(global_threads_x) / local_threads_x), ceil(static_cast(global_threads_y) / local_threads_y)), dim3(local_threads_x, local_threads_y), 0, stream>>>(dst_image_r, dst_image_g, dst_image_b, - dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_yuv_image + src_u_image_offset, - src_yuv_image + src_v_image_offset, src_yuv_image_stride_in_bytes, + dst_image_stride_in_bytes, dst_image_stride_in_bytes_comp, src_yuv_image, src_u_image, + src_v_image, src_yuv_image_stride_in_bytes, dst_width_comp, dst_height_comp, src_yuv_image_stride_in_bytes_comp); } @@ -2243,4 +2243,472 @@ void ConvertPackedYUYVToPlanarYUV(hipStream_t stream, uint32_t dst_width, uint32 ConvertPackedYUYVToPlanarYUVKernel<<(global_threads_x) / local_threads_x), ceil(static_cast(global_threads_y) / local_threads_y)), dim3(local_threads_x, local_threads_y), 0, stream>>>(dst_height, destination_y, destination_u, destination_v, dst_luma_stride_in_bytes, dst_chroma_stride_in_bytes, src_image, src_image_stride_in_bytes, dst_width_comp); +} + +// ============================================================================ +// Batched kernels: one kernel launch converts a group of images. hipBlockIdx_z +// selects the per-image parameter struct. Bodies are copied verbatim from the +// single-image kernels above; per-image args are aliased to struct fields. +// ============================================================================ + +__global__ void ColorConvertRGBAToRGBBatchedKernel(const RGBAToRGBBatchParams *params, uint32_t n_images) { + if (hipBlockIdx_z >= n_images) return; + const RGBAToRGBBatchParams &p = params[hipBlockIdx_z]; + uint32_t dst_width = p.dst_width; + uint32_t dst_height = p.dst_height; + uint8_t *dst_image = p.dst_image; + uint32_t dst_image_stride_in_bytes = p.dst_image_stride_in_bytes; + const uint8_t *src_image = p.src_image; + uint32_t src_image_stride_in_bytes = p.src_image_stride_in_bytes; + + uint32_t x = (hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x) * 8; + uint32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + if (x >= dst_width || y >= dst_height) { + return; + } + + uint32_t src_idx = y * src_image_stride_in_bytes + (x << 2); + uint32_t dst_idx = y * dst_image_stride_in_bytes + (x * 3); + + DUINT8 src = *((DUINT8 *)(&src_image[src_idx])); + DUINT6 dst; + + dst.data[0] = hipPack(make_float4(hipUnpack0(src.data[0]), hipUnpack1(src.data[0]), hipUnpack2(src.data[0]), hipUnpack0(src.data[1]))); + dst.data[1] = hipPack(make_float4(hipUnpack1(src.data[1]), hipUnpack2(src.data[1]), hipUnpack0(src.data[2]), hipUnpack1(src.data[2]))); + dst.data[2] = hipPack(make_float4(hipUnpack2(src.data[2]), hipUnpack0(src.data[3]), hipUnpack1(src.data[3]), hipUnpack2(src.data[3]))); + dst.data[3] = hipPack(make_float4(hipUnpack0(src.data[4]), hipUnpack1(src.data[4]), hipUnpack2(src.data[4]), hipUnpack0(src.data[5]))); + dst.data[4] = hipPack(make_float4(hipUnpack1(src.data[5]), hipUnpack2(src.data[5]), hipUnpack0(src.data[6]), hipUnpack1(src.data[6]))); + dst.data[5] = hipPack(make_float4(hipUnpack2(src.data[6]), hipUnpack0(src.data[7]), hipUnpack1(src.data[7]), hipUnpack2(src.data[7]))); + + *((DUINT6 *)(&dst_image[dst_idx])) = dst; +} + +void ColorConvertRGBAToRGBBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const RGBAToRGBBatchParams *d_params, uint32_t n_images) { + uint32_t grid_x = static_cast(ceil(static_cast(max_gx) / 16)); + uint32_t grid_y = static_cast(ceil(static_cast(max_gy) / 16)); + ColorConvertRGBAToRGBBatchedKernel<<>>(d_params, n_images); +} + +__global__ void ConvertInterleavedUVToPlanarUVBatchedKernel(const InterleavedUVToPlanarUVBatchParams *params, uint32_t n_images) { + if (hipBlockIdx_z >= n_images) return; + const InterleavedUVToPlanarUVBatchParams &p = params[hipBlockIdx_z]; + uint32_t dst_width = p.dst_width; + uint32_t dst_height = p.dst_height; + uint8_t *dst_image1 = p.dst_image1; + uint8_t *dst_image2 = p.dst_image2; + uint32_t dst_image_stride_in_bytes = p.dst_image_stride_in_bytes; + const uint8_t *src_image = p.src_image; + uint32_t src_image_stride_in_bytes = p.src_image_stride_in_bytes; + + uint32_t x = (hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x) * 8; + uint32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + if (x >= dst_width || y >= dst_height) { + return; + } + + uint32_t src_idx = y * src_image_stride_in_bytes + x + x; + uint32_t dst_idx = y * dst_image_stride_in_bytes + x; + + uint4 src = *((uint4 *)(&src_image[src_idx])); + uint2 dst1, dst2; + + dst1.x = hipPack(make_float4(hipUnpack0(src.x), hipUnpack2(src.x), hipUnpack0(src.y), hipUnpack2(src.y))); + dst1.y = hipPack(make_float4(hipUnpack0(src.z), hipUnpack2(src.z), hipUnpack0(src.w), hipUnpack2(src.w))); + dst2.x = hipPack(make_float4(hipUnpack1(src.x), hipUnpack3(src.x), hipUnpack1(src.y), hipUnpack3(src.y))); + dst2.y = hipPack(make_float4(hipUnpack1(src.z), hipUnpack3(src.z), hipUnpack1(src.w), hipUnpack3(src.w))); + + *((uint2 *)(&dst_image1[dst_idx])) = dst1; + *((uint2 *)(&dst_image2[dst_idx])) = dst2; +} + +void ConvertInterleavedUVToPlanarUVBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const InterleavedUVToPlanarUVBatchParams *d_params, uint32_t n_images) { + uint32_t grid_x = static_cast(ceil(static_cast(max_gx) / 16)); + uint32_t grid_y = static_cast(ceil(static_cast(max_gy) / 16)); + ConvertInterleavedUVToPlanarUVBatchedKernel<<>>(d_params, n_images); +} + +__global__ void ExtractYFromPackedYUYVBatchedKernel(const YFromPackedYUYVBatchParams *params, uint32_t n_images) { + if (hipBlockIdx_z >= n_images) return; + const YFromPackedYUYVBatchParams &p = params[hipBlockIdx_z]; + uint32_t dst_height = p.dst_height; + uint8_t *destination_y = p.destination_y; + uint32_t dst_luma_stride_in_bytes = p.dst_luma_stride_in_bytes; + const uint8_t *src_image = p.src_image; + uint32_t src_image_stride_in_bytes = p.src_image_stride_in_bytes; + uint32_t dst_width_in_8px_blocks = p.dst_width_in_8px_blocks; + + uint32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + uint32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + if (x < dst_width_in_8px_blocks && y < dst_height) { + uint32_t src_idx = y * src_image_stride_in_bytes + (x << 4); + uint32_t dst_idx = y * dst_luma_stride_in_bytes + (x << 3); + + uint4 src = *((uint4 *)(&src_image[src_idx])); + uint2 dst_y; + dst_y.x = hipPack(make_float4(hipUnpack0(src.x), hipUnpack2(src.x), hipUnpack0(src.y), hipUnpack2(src.y))); + dst_y.y = hipPack(make_float4(hipUnpack0(src.z), hipUnpack2(src.z), hipUnpack0(src.w), hipUnpack2(src.w))); + + *((uint2 *)(&destination_y[dst_idx])) = dst_y; + } +} + +void ExtractYFromPackedYUYVBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const YFromPackedYUYVBatchParams *d_params, uint32_t n_images) { + uint32_t grid_x = static_cast(ceil(static_cast(max_gx) / 16)); + uint32_t grid_y = static_cast(ceil(static_cast(max_gy) / 4)); + ExtractYFromPackedYUYVBatchedKernel<<>>(d_params, n_images); +} + +__global__ void ConvertPackedYUYVToPlanarYUVBatchedKernel(const PackedYUYVToPlanarYUVBatchParams *params, uint32_t n_images) { + if (hipBlockIdx_z >= n_images) return; + const PackedYUYVToPlanarYUVBatchParams &p = params[hipBlockIdx_z]; + uint32_t dst_height = p.dst_height; + uint8_t *destination_y = p.destination_y; + uint8_t *destination_u = p.destination_u; + uint8_t *destination_v = p.destination_v; + uint32_t dst_luma_stride_in_bytes = p.dst_luma_stride_in_bytes; + uint32_t dst_chroma_stride_in_bytes = p.dst_chroma_stride_in_bytes; + const uint8_t *src_image = p.src_image; + uint32_t src_image_stride_in_bytes = p.src_image_stride_in_bytes; + uint32_t dst_width_in_8px_blocks = p.dst_width_in_8px_blocks; + + uint32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + uint32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + if ((x < dst_width_in_8px_blocks && y < dst_height)) { + uint32_t src_idx = y * src_image_stride_in_bytes + (x << 4); + uint32_t dst_y_idx = y * dst_luma_stride_in_bytes + (x << 3); + uint32_t dst_uv_idx = y * dst_chroma_stride_in_bytes + (x << 2); + + uint4 src = *((uint4 *)(&src_image[src_idx])); + uint2 dst_y; + uint32_t dst_u, dst_v; + + dst_y.x = hipPack(make_float4(hipUnpack0(src.x), hipUnpack2(src.x), hipUnpack0(src.y), hipUnpack2(src.y))); + dst_y.y = hipPack(make_float4(hipUnpack0(src.z), hipUnpack2(src.z), hipUnpack0(src.w), hipUnpack2(src.w))); + dst_u = hipPack(make_float4(hipUnpack1(src.x), hipUnpack1(src.y), hipUnpack1(src.z), hipUnpack1(src.w))); + dst_v = hipPack(make_float4(hipUnpack3(src.x), hipUnpack3(src.y), hipUnpack3(src.z), hipUnpack3(src.w))); + + *((uint2 *)(&destination_y[dst_y_idx])) = dst_y; + *((uint32_t *)(&destination_u[dst_uv_idx])) = dst_u; + *((uint32_t *)(&destination_v[dst_uv_idx])) = dst_v; + } +} + +void ConvertPackedYUYVToPlanarYUVBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const PackedYUYVToPlanarYUVBatchParams *d_params, uint32_t n_images) { + uint32_t grid_x = static_cast(ceil(static_cast(max_gx) / 16)); + uint32_t grid_y = static_cast(ceil(static_cast(max_gy) / 4)); + ConvertPackedYUYVToPlanarYUVBatchedKernel<<>>(d_params, n_images); +} + +// ============================================================================ +// Unified YUV->RGB batched kernels. A per-image `layout` tag (block-uniform, +// since hipBlockIdx_z fixes the image for a whole block) selects the source-load +// path; the color-conversion and the write are shared. The conversion body and the +// planar-deinterleave write are copied verbatim from ColorConvertYUV444ToRGBKernel +// and ColorConvertYUV444ToRGBPlanarKernel respectively. +// ============================================================================ + +// Fills the six per-2x8 sample vectors (y0,y1,u0,u1,v0,v1) from whichever source +// layout the `layout` tag denotes. All five layouts share the same 8-luma-per-thread, +// 2-rows-per-thread geometry, so only the load differs. +__device__ __forceinline__ void LoadYUVSamples(uint32_t layout, + const uint8_t *src_y_image, uint32_t src_y_stride, uint32_t src_y_row_pair_stride, + const uint8_t *src_u_image, const uint8_t *src_v_image, + const uint8_t *src_chroma_image, uint32_t src_chroma_stride, + uint32_t x, uint32_t y, + uint2 &y0, uint2 &y1, uint2 &u0, uint2 &u1, uint2 &v0, uint2 &v1) { + + if (layout == YUV_LAYOUT_YUYV) { + uint32_t l0_idx = y * src_y_row_pair_stride + (x << 4); + uint32_t l1_idx = l0_idx + src_y_stride; + uint4 l0 = *((uint4 *)(&src_y_image[l0_idx])); + uint4 l1 = *((uint4 *)(&src_y_image[l1_idx])); + y0.x = hipPack(make_float4(hipUnpack0(l0.x), hipUnpack2(l0.x), hipUnpack0(l0.y), hipUnpack2(l0.y))); + y0.y = hipPack(make_float4(hipUnpack0(l0.z), hipUnpack2(l0.z), hipUnpack0(l0.w), hipUnpack2(l0.w))); + y1.x = hipPack(make_float4(hipUnpack0(l1.x), hipUnpack2(l1.x), hipUnpack0(l1.y), hipUnpack2(l1.y))); + y1.y = hipPack(make_float4(hipUnpack0(l1.z), hipUnpack2(l1.z), hipUnpack0(l1.w), hipUnpack2(l1.w))); + u0.x = hipPack(make_float4(hipUnpack1(l0.x), hipUnpack1(l0.x), hipUnpack1(l0.y), hipUnpack1(l0.y))); + u0.y = hipPack(make_float4(hipUnpack1(l0.z), hipUnpack1(l0.z), hipUnpack1(l0.w), hipUnpack1(l0.w))); + u1.x = hipPack(make_float4(hipUnpack1(l1.x), hipUnpack1(l1.x), hipUnpack1(l1.y), hipUnpack1(l1.y))); + u1.y = hipPack(make_float4(hipUnpack1(l1.z), hipUnpack1(l1.z), hipUnpack1(l1.w), hipUnpack1(l1.w))); + v0.x = hipPack(make_float4(hipUnpack3(l0.x), hipUnpack3(l0.x), hipUnpack3(l0.y), hipUnpack3(l0.y))); + v0.y = hipPack(make_float4(hipUnpack3(l0.z), hipUnpack3(l0.z), hipUnpack3(l0.w), hipUnpack3(l0.w))); + v1.x = hipPack(make_float4(hipUnpack3(l1.x), hipUnpack3(l1.x), hipUnpack3(l1.y), hipUnpack3(l1.y))); + v1.y = hipPack(make_float4(hipUnpack3(l1.z), hipUnpack3(l1.z), hipUnpack3(l1.w), hipUnpack3(l1.w))); + return; + } + + // Y / luma plane load, common to NV12 / YUV444 / YUV440 / YUV400. + uint32_t src_y0_idx = y * src_y_row_pair_stride + (x << 3); + uint32_t src_y1_idx = src_y0_idx + src_y_stride; + y0 = *((uint2 *)(&src_y_image[src_y0_idx])); + y1 = *((uint2 *)(&src_y_image[src_y1_idx])); + + if (layout == YUV_LAYOUT_YUV444) { + u0 = *((uint2 *)(&src_u_image[src_y0_idx])); + u1 = *((uint2 *)(&src_u_image[src_y1_idx])); + v0 = *((uint2 *)(&src_v_image[src_y0_idx])); + v1 = *((uint2 *)(&src_v_image[src_y1_idx])); + } else if (layout == YUV_LAYOUT_YUV440) { + uint32_t src_chroma_idx = y * src_y_stride + (x << 3); + u0 = *((uint2 *)(&src_u_image[src_chroma_idx])); + v0 = *((uint2 *)(&src_v_image[src_chroma_idx])); + u1 = u0; + v1 = v0; + } else if (layout == YUV_LAYOUT_NV12) { + uint32_t src_uv_idx = y * src_chroma_stride + (x << 3); + uint2 uv = *((uint2 *)(&src_chroma_image[src_uv_idx])); + float4 f; + f.x = hipUnpack0(uv.x); f.y = f.x; f.z = hipUnpack2(uv.x); f.w = f.z; u0.x = hipPack(f); + f.x = hipUnpack0(uv.y); f.y = f.x; f.z = hipUnpack2(uv.y); f.w = f.z; u0.y = hipPack(f); + u1.x = u0.x; u1.y = u0.y; + f.x = hipUnpack1(uv.x); f.y = f.x; f.z = hipUnpack3(uv.x); f.w = f.z; v0.x = hipPack(f); + f.x = hipUnpack1(uv.y); f.y = f.x; f.z = hipUnpack3(uv.y); f.w = f.z; v0.y = hipPack(f); + v1.x = v0.x; v1.y = v0.y; + } else { // YUV_LAYOUT_YUV400: no chroma -> neutral 128 gives R=G=B=Y. + uint2 neutral; + neutral.x = hipPack(make_float4(128.0f, 128.0f, 128.0f, 128.0f)); + neutral.y = neutral.x; + u0 = neutral; u1 = neutral; v0 = neutral; v1 = neutral; + } +} + +// Shared YUV->RGB conversion (verbatim from ColorConvertYUV444ToRGBKernel). +__device__ __forceinline__ void ConvertYUVToRGB( + uint2 y0, uint2 y1, uint2 u0, uint2 u1, uint2 v0, uint2 v1, + DUINT6 &rgb0, DUINT6 &rgb1) { + float2 cr = make_float2(CC_CR0, CC_CR1); + float2 cg = make_float2(CC_CG0, CC_CG1); + float2 cb = make_float2(CC_CB0, CC_CB1); + float3 yuv; + float4 f; + + yuv = make_float3(hipUnpack0(y0.x), hipUnpack0(u0.x), hipUnpack0(v0.x)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.x = fmaf(cr.y, yuv.z, yuv.x); + f.y = fmaf(cg.x, yuv.y, yuv.x); + f.y = fmaf(cg.y, yuv.z, f.y); + f.z = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack1(y0.x), hipUnpack1(u0.x), hipUnpack1(v0.x)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.w = fmaf(cr.y, yuv.z, yuv.x); + rgb0.data[0] = hipPack(f); + + f.x = fmaf(cg.x, yuv.y, yuv.x); + f.x = fmaf(cg.y, yuv.z, f.x); + f.y = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack2(y0.x), hipUnpack2(u0.x), hipUnpack2(v0.x)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.z = fmaf(cr.y, yuv.z, yuv.x); + f.w = fmaf(cg.x, yuv.y, yuv.x); + f.w = fmaf(cg.y, yuv.z, f.w); + rgb0.data[1] = hipPack(f); + + f.x = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack3(y0.x), hipUnpack3(u0.x), hipUnpack3(v0.x)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.y = fmaf(cr.y, yuv.z, yuv.x); + f.z = fmaf(cg.x, yuv.y, yuv.x); + f.z = fmaf(cg.y, yuv.z, f.z); + f.w = fmaf(cb.x, yuv.y, yuv.x); + rgb0.data[2] = hipPack(f); + + yuv = make_float3(hipUnpack0(y0.y), hipUnpack0(u0.y), hipUnpack0(v0.y)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.x = fmaf(cr.y, yuv.z, yuv.x); + f.y = fmaf(cg.x, yuv.y, yuv.x); + f.y = fmaf(cg.y, yuv.z, f.y); + f.z = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack1(y0.y), hipUnpack1(u0.y), hipUnpack1(v0.y)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.w = fmaf(cr.y, yuv.z, yuv.x); + rgb0.data[3] = hipPack(f); + + f.x = fmaf(cg.x, yuv.y, yuv.x); + f.x = fmaf(cg.y, yuv.z, f.x); + f.y = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack2(y0.y), hipUnpack2(u0.y), hipUnpack2(v0.y)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.z = fmaf(cr.y, yuv.z, yuv.x); + f.w = fmaf(cg.x, yuv.y, yuv.x); + f.w = fmaf(cg.y, yuv.z, f.w); + rgb0.data[4] = hipPack(f); + + f.x = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack3(y0.y), hipUnpack3(u0.y), hipUnpack3(v0.y)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.y = fmaf(cr.y, yuv.z, yuv.x); + f.z = fmaf(cg.x, yuv.y, yuv.x); + f.z = fmaf(cg.y, yuv.z, f.z); + f.w = fmaf(cb.x, yuv.y, yuv.x); + rgb0.data[5] = hipPack(f); + + yuv = make_float3(hipUnpack0(y1.x), hipUnpack0(u1.x), hipUnpack0(v1.x)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.x = fmaf(cr.y, yuv.z, yuv.x); + f.y = fmaf(cg.x, yuv.y, yuv.x); + f.y = fmaf(cg.y, yuv.z, f.y); + f.z = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack1(y1.x), hipUnpack1(u1.x), hipUnpack1(v1.x)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.w = fmaf(cr.y, yuv.z, yuv.x); + rgb1.data[0] = hipPack(f); + + f.x = fmaf(cg.x, yuv.y, yuv.x); + f.x = fmaf(cg.y, yuv.z, f.x); + f.y = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack2(y1.x), hipUnpack2(u1.x), hipUnpack2(v1.x)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.z = fmaf(cr.y, yuv.z, yuv.x); + f.w = fmaf(cg.x, yuv.y, yuv.x); + f.w = fmaf(cg.y, yuv.z, f.w); + rgb1.data[1] = hipPack(f); + + f.x = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack3(y1.x), hipUnpack3(u1.x), hipUnpack3(v1.x)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.y = fmaf(cr.y, yuv.z, yuv.x); + f.z = fmaf(cg.x, yuv.y, yuv.x); + f.z = fmaf(cg.y, yuv.z, f.z); + f.w = fmaf(cb.x, yuv.y, yuv.x); + rgb1.data[2] = hipPack(f); + + yuv = make_float3(hipUnpack0(y1.y), hipUnpack0(u1.y), hipUnpack0(v1.y)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.x = fmaf(cr.y, yuv.z, yuv.x); + f.y = fmaf(cg.x, yuv.y, yuv.x); + f.y = fmaf(cg.y, yuv.z, f.y); + f.z = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack1(y1.y), hipUnpack1(u1.y), hipUnpack1(v1.y)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.w = fmaf(cr.y, yuv.z, yuv.x); + rgb1.data[3] = hipPack(f); + + f.x = fmaf(cg.x, yuv.y, yuv.x); + f.x = fmaf(cg.y, yuv.z, f.x); + f.y = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack2(y1.y), hipUnpack2(u1.y), hipUnpack2(v1.y)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.z = fmaf(cr.y, yuv.z, yuv.x); + f.w = fmaf(cg.x, yuv.y, yuv.x); + f.w = fmaf(cg.y, yuv.z, f.w); + rgb1.data[4] = hipPack(f); + + f.x = fmaf(cb.x, yuv.y, yuv.x); + yuv = make_float3(hipUnpack3(y1.y), hipUnpack3(u1.y), hipUnpack3(v1.y)); + yuv.y -= 128.0f; + yuv.z -= 128.0f; + f.y = fmaf(cr.y, yuv.z, yuv.x); + f.z = fmaf(cg.x, yuv.y, yuv.x); + f.z = fmaf(cg.y, yuv.z, f.z); + f.w = fmaf(cb.x, yuv.y, yuv.x); + rgb1.data[5] = hipPack(f); +} + +__global__ void ColorConvertYUVToRGBBatchedKernel(const YUVToRGBBatchParams *params, uint32_t n_images) { + if (hipBlockIdx_z >= n_images) return; + const YUVToRGBBatchParams &p = params[hipBlockIdx_z]; + + uint32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + uint32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + if ((x < p.dst_width_in_8px_blocks) && (y < p.dst_height_in_2row_blocks)) { + uint2 y0, y1, u0, u1, v0, v1; + LoadYUVSamples(p.layout, p.src_y_image, p.src_y_image_stride_in_bytes, p.src_y_image_row_pair_stride_in_bytes, + p.src_u_image, p.src_v_image, p.src_chroma_image, p.src_chroma_image_stride_in_bytes, + x, y, y0, y1, u0, u1, v0, v1); + DUINT6 rgb0, rgb1; + ConvertYUVToRGB(y0, y1, u0, u1, v0, v1, rgb0, rgb1); + uint32_t rgb0_idx = y * p.dst_image_row_pair_stride_in_bytes + (x * 24); + uint32_t rgb1_idx = rgb0_idx + p.dst_image_stride_in_bytes; + *((DUINT6 *)(&p.dst_image[rgb0_idx])) = rgb0; + *((DUINT6 *)(&p.dst_image[rgb1_idx])) = rgb1; + } +} + +void ColorConvertYUVToRGBBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const YUVToRGBBatchParams *d_params, uint32_t n_images) { + uint32_t grid_x = static_cast(ceil(static_cast(max_gx) / 16)); + uint32_t grid_y = static_cast(ceil(static_cast(max_gy) / 4)); + ColorConvertYUVToRGBBatchedKernel<<>>(d_params, n_images); +} + +__global__ void ColorConvertYUVToRGBPlanarBatchedKernel(const YUVToRGBPlanarBatchParams *params, uint32_t n_images) { + if (hipBlockIdx_z >= n_images) return; + const YUVToRGBPlanarBatchParams &p = params[hipBlockIdx_z]; + + uint32_t x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + uint32_t y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + if ((x < p.dst_width_in_8px_blocks) && (y < p.dst_height_in_2row_blocks)) { + uint2 y0, y1, u0, u1, v0, v1; + LoadYUVSamples(p.layout, p.src_y_image, p.src_y_image_stride_in_bytes, p.src_y_image_row_pair_stride_in_bytes, + p.src_u_image, p.src_v_image, p.src_chroma_image, p.src_chroma_image_stride_in_bytes, + x, y, y0, y1, u0, u1, v0, v1); + DUINT6 rgb0, rgb1; + ConvertYUVToRGB(y0, y1, u0, u1, v0, v1, rgb0, rgb1); + + uint8_t *dst_image_r = p.dst_image_r; + uint8_t *dst_image_g = p.dst_image_g; + uint8_t *dst_image_b = p.dst_image_b; + uint32_t rgb0_idx = y * p.dst_image_row_pair_stride_in_bytes + (x * 8); + uint32_t rgb1_idx = rgb0_idx + p.dst_image_stride_in_bytes; + uint2 red0, red1, green0, green1, blue0, blue1; + red0.x = hipPack(make_float4(hipUnpack0(rgb0.data[0]), hipUnpack3(rgb0.data[0]), hipUnpack2(rgb0.data[1]), hipUnpack1(rgb0.data[2]))); + red0.y = hipPack(make_float4(hipUnpack0(rgb0.data[3]), hipUnpack3(rgb0.data[3]), hipUnpack2(rgb0.data[4]), hipUnpack1(rgb0.data[5]))); + red1.x = hipPack(make_float4(hipUnpack0(rgb1.data[0]), hipUnpack3(rgb1.data[0]), hipUnpack2(rgb1.data[1]), hipUnpack1(rgb1.data[2]))); + red1.y = hipPack(make_float4(hipUnpack0(rgb1.data[3]), hipUnpack3(rgb1.data[3]), hipUnpack2(rgb1.data[4]), hipUnpack1(rgb1.data[5]))); + + green0.x = hipPack(make_float4(hipUnpack1(rgb0.data[0]), hipUnpack0(rgb0.data[1]), hipUnpack3(rgb0.data[1]), hipUnpack2(rgb0.data[2]))); + green0.y = hipPack(make_float4(hipUnpack1(rgb0.data[3]), hipUnpack0(rgb0.data[4]), hipUnpack3(rgb0.data[4]), hipUnpack2(rgb0.data[5]))); + green1.x = hipPack(make_float4(hipUnpack1(rgb1.data[0]), hipUnpack0(rgb1.data[1]), hipUnpack3(rgb1.data[1]), hipUnpack2(rgb1.data[2]))); + green1.y = hipPack(make_float4(hipUnpack1(rgb1.data[3]), hipUnpack0(rgb1.data[4]), hipUnpack3(rgb1.data[4]), hipUnpack2(rgb1.data[5]))); + + blue0.x = hipPack(make_float4(hipUnpack2(rgb0.data[0]), hipUnpack1(rgb0.data[1]), hipUnpack0(rgb0.data[2]), hipUnpack3(rgb0.data[2]))); + blue0.y = hipPack(make_float4(hipUnpack2(rgb0.data[3]), hipUnpack1(rgb0.data[4]), hipUnpack0(rgb0.data[5]), hipUnpack3(rgb0.data[5]))); + blue1.x = hipPack(make_float4(hipUnpack2(rgb1.data[0]), hipUnpack1(rgb1.data[1]), hipUnpack0(rgb1.data[2]), hipUnpack3(rgb1.data[2]))); + blue1.y = hipPack(make_float4(hipUnpack2(rgb1.data[3]), hipUnpack1(rgb1.data[4]), hipUnpack0(rgb1.data[5]), hipUnpack3(rgb1.data[5]))); + + *((uint2 *)(&dst_image_r[rgb0_idx])) = red0; + *((uint2 *)(&dst_image_r[rgb1_idx])) = red1; + *((uint2 *)(&dst_image_g[rgb0_idx])) = green0; + *((uint2 *)(&dst_image_g[rgb1_idx])) = green1; + *((uint2 *)(&dst_image_b[rgb0_idx])) = blue0; + *((uint2 *)(&dst_image_b[rgb1_idx])) = blue1; + } +} + +void ColorConvertYUVToRGBPlanarBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const YUVToRGBPlanarBatchParams *d_params, uint32_t n_images) { + uint32_t grid_x = static_cast(ceil(static_cast(max_gx) / 16)); + uint32_t grid_y = static_cast(ceil(static_cast(max_gy) / 4)); + ColorConvertYUVToRGBPlanarBatchedKernel<<>>(d_params, n_images); } \ No newline at end of file diff --git a/projects/rocjpeg/src/rocjpeg_hip_kernels.h b/projects/rocjpeg/src/rocjpeg_hip_kernels.h index e43a456fc889..855d3666e8f8 100644 --- a/projects/rocjpeg/src/rocjpeg_hip_kernels.h +++ b/projects/rocjpeg/src/rocjpeg_hip_kernels.h @@ -75,11 +75,12 @@ THE SOFTWARE. * @param dst_image_stride_in_bytes The stride (in bytes) of the destination RGB image buffer. * @param src_yuv_image Pointer to the source YUV444 image buffer. * @param src_yuv_image_stride_in_bytes The stride (in bytes) of the source YUV444 image buffer. - * @param src_u_image_offset The offset (in bytes) of the U component in the source YUV444 image buffer. + * @param src_u_image Pointer to the U plane of the source YUV444 image (already ROI-adjusted). + * @param src_v_image Pointer to the V plane of the source YUV444 image (already ROI-adjusted). */ void ColorConvertYUV444ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset); + uint32_t src_yuv_image_stride_in_bytes, const uint8_t *src_u_image, const uint8_t *src_v_image); /** * @brief Converts YUV440 image to RGB image. @@ -93,11 +94,12 @@ void ColorConvertYUV444ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t ds * @param dst_image_stride_in_bytes The stride (in bytes) of the destination RGB image buffer. * @param src_yuv_image Pointer to the source YUV440 image buffer. * @param src_yuv_image_stride_in_bytes The stride (in bytes) of the source YUV440 image buffer. - * @param src_u_image_offset The offset (in bytes) of the U component in the source YUV440 image buffer. + * @param src_u_image Pointer to the U plane of the source YUV440 image (already ROI-adjusted). + * @param src_v_image Pointer to the V plane of the source YUV440 image (already ROI-adjusted). */ void ColorConvertYUV440ToRGB(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset); + uint32_t src_yuv_image_stride_in_bytes, const uint8_t *src_u_image, const uint8_t *src_v_image); /** * @brief Converts an image in YUYV format to RGB format. @@ -188,11 +190,12 @@ void ColorConvertRGBAToRGB(hipStream_t stream, uint32_t dst_width, uint32_t dst_ * @param dst_image_stride_in_bytes The stride (in bytes) of the destination RGB image. * @param src_yuv_image Pointer to the source YUV444 image. * @param src_yuv_image_stride_in_bytes The stride (in bytes) of the source YUV444 image. - * @param src_u_image_offset The offset (in bytes) of the U plane in the source YUV444 image. + * @param src_u_image Pointer to the U plane of the source YUV444 image (already ROI-adjusted). + * @param src_v_image Pointer to the V plane of the source YUV444 image (already ROI-adjusted). */ void ColorConvertYUV444ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image_r, uint8_t *dst_image_g, uint8_t *dst_image_b, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset); + uint32_t src_yuv_image_stride_in_bytes, const uint8_t *src_u_image, const uint8_t *src_v_image); /** * @brief Converts YUV440 image to RGB planar format. @@ -209,11 +212,12 @@ void ColorConvertYUV444ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint3 * @param dst_image_stride_in_bytes The stride (in bytes) of the destination RGB image. * @param src_yuv_image Pointer to the source YUV440 image. * @param src_yuv_image_stride_in_bytes The stride (in bytes) of the source YUV440 image. - * @param src_u_image_offset The offset (in bytes) of the U plane in the source YUV440 image. + * @param src_u_image Pointer to the U plane of the source YUV440 image (already ROI-adjusted). + * @param src_v_image Pointer to the V plane of the source YUV440 image (already ROI-adjusted). */ void ColorConvertYUV440ToRGBPlanar(hipStream_t stream, uint32_t dst_width, uint32_t dst_height, uint8_t *dst_image_r, uint8_t *dst_image_g, uint8_t *dst_image_b, uint32_t dst_image_stride_in_bytes, const uint8_t *src_yuv_image, - uint32_t src_yuv_image_stride_in_bytes, uint32_t src_u_image_offset, uint32_t src_v_image_offset); + uint32_t src_yuv_image_stride_in_bytes, const uint8_t *src_u_image, const uint8_t *src_v_image); /** * Converts a YUYV image to RGB planar format. @@ -337,6 +341,110 @@ void ConvertPackedYUYVToPlanarYUV(hipStream_t stream, uint32_t dst_width, uint32 uint8_t *destination_y, uint8_t *destination_u, uint8_t *destination_v, uint32_t dst_luma_stride_in_bytes, uint32_t dst_chroma_stride_in_bytes, const uint8_t *src_image, uint32_t src_image_stride_in_bytes); +// ---- batched param structs + prototypes ---- + +struct RGBAToRGBBatchParams { + uint32_t dst_width; + uint32_t dst_height; + uint8_t *dst_image; + uint32_t dst_image_stride_in_bytes; + const uint8_t *src_image; + uint32_t src_image_stride_in_bytes; +}; +void ColorConvertRGBAToRGBBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const RGBAToRGBBatchParams *d_params, uint32_t n_images); + +struct InterleavedUVToPlanarUVBatchParams { + uint32_t dst_width; + uint32_t dst_height; + uint8_t *dst_image1; + uint8_t *dst_image2; + uint32_t dst_image_stride_in_bytes; + const uint8_t *src_image; + uint32_t src_image_stride_in_bytes; +}; +void ConvertInterleavedUVToPlanarUVBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const InterleavedUVToPlanarUVBatchParams *d_params, uint32_t n_images); + +struct YFromPackedYUYVBatchParams { + uint32_t dst_height; + uint8_t *destination_y; + uint32_t dst_luma_stride_in_bytes; + const uint8_t *src_image; + uint32_t src_image_stride_in_bytes; + uint32_t dst_width_in_8px_blocks; // Destination width in 8-pixel-wide blocks; each thread writes 8 luma samples of one row. +}; +void ExtractYFromPackedYUYVBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const YFromPackedYUYVBatchParams *d_params, uint32_t n_images); + +struct PackedYUYVToPlanarYUVBatchParams { + uint32_t dst_height; + uint8_t *destination_y; + uint8_t *destination_u; + uint8_t *destination_v; + uint32_t dst_luma_stride_in_bytes; + uint32_t dst_chroma_stride_in_bytes; + const uint8_t *src_image; + uint32_t src_image_stride_in_bytes; + uint32_t dst_width_in_8px_blocks; // Destination width in 8-pixel-wide blocks; each thread writes 8 luma samples of one row. +}; +void ConvertPackedYUYVToPlanarYUVBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const PackedYUYVToPlanarYUVBatchParams *d_params, uint32_t n_images); + +// ---- Unified YUV->RGB batched kernels (cover NV12/YUV444/YUV440/YUYV/YUV400) ---- + +// Per-image source layout selector for the unified kernels. +enum YUVSurfaceLayout : uint32_t { + YUV_LAYOUT_NV12 = 0, // 2-plane: luma + interleaved UV (4:2:0) + YUV_LAYOUT_YUV444, // 3-plane 4:4:4 + YUV_LAYOUT_YUV440, // 3-plane 4:4:0 (422V), chroma half-height + YUV_LAYOUT_YUYV, // 1-plane packed YUYV (4:2:2) + YUV_LAYOUT_YUV400, // 1-plane luma only (grayscale) +}; + +// Superset of per-image source params; which fields are read depends on `layout`. +// Each thread of the unified kernels converts an 8-pixel-wide by 2-row block, so the +// "*_in_8px_blocks" / "*_in_2row_blocks" extents are the per-image grid bounds and the +// "*_row_pair_stride_in_bytes" values are the byte distance between consecutive row pairs. +struct YUVToRGBBatchParams { + uint32_t layout; + uint8_t *dst_image; + uint32_t dst_image_stride_in_bytes; + uint32_t dst_image_row_pair_stride_in_bytes; // dst_image_stride_in_bytes * 2 (one thread writes 2 rows) + const uint8_t *src_y_image; + uint32_t src_y_image_stride_in_bytes; + uint32_t src_y_image_row_pair_stride_in_bytes; // src_y_image_stride_in_bytes * 2 (one thread reads 2 rows) + const uint8_t *src_u_image; // YUV444 / YUV440 + const uint8_t *src_v_image; // YUV444 / YUV440 + const uint8_t *src_chroma_image; // NV12 interleaved UV + uint32_t src_chroma_image_stride_in_bytes; // NV12 + uint32_t dst_width_in_8px_blocks; // ceil(dst_width / 8) + uint32_t dst_height_in_2row_blocks; // ceil(dst_height / 2) +}; + +struct YUVToRGBPlanarBatchParams { + uint32_t layout; + uint8_t *dst_image_r; + uint8_t *dst_image_g; + uint8_t *dst_image_b; + uint32_t dst_image_stride_in_bytes; + uint32_t dst_image_row_pair_stride_in_bytes; // dst_image_stride_in_bytes * 2 (one thread writes 2 rows) + const uint8_t *src_y_image; + uint32_t src_y_image_stride_in_bytes; + uint32_t src_y_image_row_pair_stride_in_bytes; // src_y_image_stride_in_bytes * 2 (one thread reads 2 rows) + const uint8_t *src_u_image; + const uint8_t *src_v_image; + const uint8_t *src_chroma_image; + uint32_t src_chroma_image_stride_in_bytes; + uint32_t dst_width_in_8px_blocks; // ceil(dst_width / 8) + uint32_t dst_height_in_2row_blocks; // ceil(dst_height / 2) +}; + +void ColorConvertYUVToRGBBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const YUVToRGBBatchParams *d_params, uint32_t n_images); +void ColorConvertYUVToRGBPlanarBatched(hipStream_t stream, uint32_t max_gx, uint32_t max_gy, + const YUVToRGBPlanarBatchParams *d_params, uint32_t n_images); + /** * @brief Structure representing an array of 6 unsigned integers. * diff --git a/projects/rocjpeg/src/rocjpeg_vaapi_decoder.cpp b/projects/rocjpeg/src/rocjpeg_vaapi_decoder.cpp index da7286f807a4..55de13cdf390 100644 --- a/projects/rocjpeg/src/rocjpeg_vaapi_decoder.cpp +++ b/projects/rocjpeg/src/rocjpeg_vaapi_decoder.cpp @@ -193,9 +193,13 @@ bool RocJpegVaapiMemoryPool::DeleteIdleEntry() { * @brief Adds a pool entry to the memory pool for a specific surface format. * * This function adds a pool entry to the memory pool for a specific surface format. - * If the memory pool for the given surface format is not full, the new entry is added to the pool. - * If the memory pool is full, the oldest entry is removed from the pool and replaced with the new entry. - * If the removed entry has associated resources (VA context, VA surface, HIP memory), they are destroyed and freed. + * + * max_pool_size_ is a reuse/eviction target, not a hard cap on live allocations. When the pool has + * reached the target, an idle entry is evicted first so memory is recycled rather than grown. If no + * idle entry exists, every entry is still in flight and the pool is allowed to grow past the target: + * on a GPU with a single JPEG core the batched decoder submits one image per call, so each image in + * the batch holds its own busy entry and a batch larger than the target legitimately needs more live + * entries than the target allows. Failing here would reject valid work. * * @param surface_format The surface format for which the pool entry is being added. * @param pool_entry The pool entry to be added. @@ -205,16 +209,11 @@ RocJpegStatus RocJpegVaapiMemoryPool::AddPoolEntry(uint32_t surface_format, cons std::lock_guard lock(pool_mutex_); size_t total_mem_pool_size = GetTotalMemPoolSize(); auto& entries = mem_pool_[surface_format]; - if (total_mem_pool_size < max_pool_size_) { - entries.push_back(pool_entry); - } else { - if (DeleteIdleEntry()) { - entries.push_back(pool_entry); - } else { - ErrorLog(g_rocjpeg_logger, "Cannot find an idle entry in the the memory pool!"); - return ROCJPEG_STATUS_INVALID_PARAMETER; - } + if (total_mem_pool_size >= max_pool_size_) { + // Best-effort eviction to stay near the target; growing past it is valid when all entries are busy. + DeleteIdleEntry(); } + entries.push_back(pool_entry); return ROCJPEG_STATUS_SUCCESS; }