You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
feat(rocJPEG): batch HIP color-conversion kernels in the batched decode path - #11876
The batched decode path (FinalizeDecodeBatched) previously launched one
color-conversion kernel per image per sub-batch — for a batch of N images,
N launches per output format, so launch overhead dominated for small/medium
images.
This PR restructures the path to accumulate per-image parameters and issue a
single kernel launch per format per sub-batch; hipBlockIdx_z selects the
image inside the kernel. For RGB / RGB_PLANAR, all supported YUV surface layouts
are collapsed into one unified kernel each, so a mixed-format batch is one
launch instead of up to five.
No public API change — behavior is identical, only the number of kernel launches
drops.
Technical Details
Kernels (rocjpeg_hip_kernels.{h,cpp})
ColorConvertYUVToRGBBatched / ColorConvertYUVToRGBPlanarBatched — unified
kernels covering NV12, YUV444, YUV440 (422V), YUYV (422), and YUV400, selected
per image by a YUVSurfaceLayout tag.
Per-format batched kernels for paths that don't fit the unified YUV model: ColorConvertRGBAToRGBBatched, ConvertPackedYUYVToPlanarYUVBatched, ConvertInterleavedUVToPlanarUVBatched, ExtractYFromPackedYUYVBatched.
All single-image kernels are unchanged.
Decoder (rocjpeg_decoder.{h,cpp})
BatchedKernelParams<T> — reusable scratch buffer (host param vector +
lazily-grown device buffer + max grid extents), uploaded once per group.
Copy/move are deleted since it owns a raw device allocation.
FinalizeDecodeBatched is now two-pass per sub-batch: (1) sync all VA surfaces
and collect HIP interop memory; (2) accumulate per-image params, then one launch
per kernel type, then a single hipStreamSynchronize.
Accumulate* helpers mirror the existing per-image conversion helpers'
pointer/offset/ROI math exactly, but push params instead of launching.
Design notes
No warp divergence from the layout branch:hipBlockIdx_z fixes the image
(hence layout) for a whole block, so the per-layout switch in the unified
load is block-uniform.
Bit-identical conversion: the unified kernel's color-conversion body and
planar-deinterleave write are copied verbatim from the YUV444 kernels; only the
source load differs per layout. YUV400 folds in by feeding U=V=128, which the
conversion collapses to R=G=B=Y exactly.
Grid sizing: one grid per launch, sized to the max width/height across the
group — very uneven image sizes waste some threads on the padded axis.
Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*). Current: code file(s) changed: projects/rocjpeg/src/rocjpeg_decoder.cpp, projects/rocjpeg/src/rocjpeg_decoder.h, projects/rocjpeg/src/rocjpeg_hip_kernels.cpp, projects/rocjpeg/src/rocjpeg_hip_kernels.h, projects/rocjpeg/src/rocjpeg_vaapi_decoder.cpp; no test file found
🚫 Draft PR
🔜 To Be Enabled
—
🚩 Feature Flag
🔜 To Be Enabled
—
📊 Code Coverage
🔜 To Be Enabled
—
🎉 All policy checks passed!
📖 Need help? See the Policy FAQ for details on every check and how to fix failures.
The reason will be displayed to describe this comment to others. Learn more.
Warning
Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.
Pull request overview
Copilot reviewed 6 out of 6 changed files in this pull request and generated 2 comments.
Release sub-batch surfaces when parameter setup fails
projects/rocjpeg/src/rocjpeg_decoder.cpp:394
If any newly added parameter upload/allocation fails here, CHECK_ROCJPEG returns before the cleanup loop below, so none of the sub-batch surfaces is passed to SetSurfaceAsIdle. The failed call can therefore leave pool entries permanently busy and make later decodes fail; error exits from this new launch path need to synchronize/ release the surfaces before returning.
The reason will be displayed to describe this comment to others. Learn more.
suggested by claude review
Bit-identity not machine-verified for 444P / 422V / RGBA. These formats aren't in the test image set, and the ctest doesn't MD5-compare (see §5). Given the whole PR rests on "verbatim/bit-identical," consider adding a 444/440 image + an output-hash check, or at least manually diffing single vs batched output for those formats before merge.
LaunchBatchedBuffer ignores the kernel-launch result (rocjpeg_decoder.h:496) — launch(...) returns void and no hipGetLastError() is checked; a bad-config launch only surfaces at the later hipStreamSynchronize. Acceptable, but a post-launch error check would localize failures.
finalize_mutex_ vs hip_stream_: the batched finalize is serialized, but the non-batched FinalizeDecode also uses hip_stream_ without taking finalize_mutex_. If both can run concurrently on one decoder, that's a shared-stream hazard. Worth confirming the intended threading model (async tests pass, so not observed).
Header comment // ---- Auto-generated batched param structs ---- (rocjpeg_hip_kernels.h:1042) — if these aren't actually generated by a tool, drop "Auto-generated" to avoid implying a generator that reviewers should look for.
Minor whitespace inconsistency mem.offset[2]+ roi_uv_offset (rocjpeg_decoder.cpp:300); pre-commit clang-format will likely fix.
Bit-identity validation / MD5 checks: Adding MD5-based verification is beyond the scope of this PR. However, I have manually verified that the batched output matches the single-image path for all currently supported chroma subsampling formats, including the formats covered by the new batched implementation. I agree that adding output hash comparisons would strengthen the test coverage, and we can address this as a testing enhancement in a follow-up PR by adding additional test images (e.g., 444P/422V) along with automated MD5-based validation.
Kernel launch error checking: Agreed. The current implementation relies on the subsequent hipStreamSynchronize() to surface launch failures. While this is consistent with the existing error handling flow, adding a post-launch hipGetLastError() check would improve error localization and can be considered as a follow-up improvement.
finalize_mutex_ vs. hip_stream_: The intended threading model does not allow concurrent execution of batched and non-batched finalize operations on the same decoder instance. This is also consistent with the current usage patterns and test coverage, which is why no issue has been observed. That said, the concern is valid, and we can add a clarifying comment or revisit synchronization if the threading model evolves.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
The batched decode path (
FinalizeDecodeBatched) previously launched onecolor-conversion kernel per image per sub-batch — for a batch of N images,
N launches per output format, so launch overhead dominated for small/medium
images.
This PR restructures the path to accumulate per-image parameters and issue a
single kernel launch per format per sub-batch;
hipBlockIdx_zselects theimage inside the kernel. For RGB / RGB_PLANAR, all supported YUV surface layouts
are collapsed into one unified kernel each, so a mixed-format batch is one
launch instead of up to five.
No public API change — behavior is identical, only the number of kernel launches
drops.
Technical Details
Kernels (
rocjpeg_hip_kernels.{h,cpp})ColorConvertYUVToRGBBatched/ColorConvertYUVToRGBPlanarBatched— unifiedkernels covering NV12, YUV444, YUV440 (422V), YUYV (422), and YUV400, selected
per image by a
YUVSurfaceLayouttag.ColorConvertRGBAToRGBBatched,ConvertPackedYUYVToPlanarYUVBatched,ConvertInterleavedUVToPlanarUVBatched,ExtractYFromPackedYUYVBatched.Decoder (
rocjpeg_decoder.{h,cpp})BatchedKernelParams<T>— reusable scratch buffer (host param vector +lazily-grown device buffer + max grid extents), uploaded once per group.
Copy/move are deleted since it owns a raw device allocation.
FinalizeDecodeBatchedis now two-pass per sub-batch: (1) sync all VA surfacesand collect HIP interop memory; (2) accumulate per-image params, then one launch
per kernel type, then a single
hipStreamSynchronize.Accumulate*helpers mirror the existing per-image conversion helpers'pointer/offset/ROI math exactly, but push params instead of launching.
Design notes
hipBlockIdx_zfixes the image(hence layout) for a whole block, so the per-layout
switchin the unifiedload is block-uniform.
planar-deinterleave write are copied verbatim from the YUV444 kernels; only the
source load differs per layout. YUV400 folds in by feeding U=V=128, which the
conversion collapses to R=G=B=Y exactly.
group — very uneven image sizes waste some threads on the padded axis.
Issue Tracking
JIRA ID: AICV-302
Test Plan
rocjpeg ctest.
Test Result
rocjpeg ctest should pass.
Submission Checklist