Implement asynchronous LDS loads for MI350 #5348
Open
avbokovoy wants to merge 5 commits into
Open
Conversation
Contributor
spcyppt
reviewed
Feb 13, 2026
|
|
||
| asm volatile("cp.async.wait_group %0;\n" ::"n"(N)); | ||
| #elif defined(USE_ROCM) && \ | ||
| (ROCM_VERSION_MAJOR <= 7 && ROCM_VERSION_MINOR < 2) && defined(__gfx950__) |
Contributor
There was a problem hiding this comment.
Is this supposed to be supported for rocm version < 7.2?
If so, it should be
(ROCM_VERSION_MAJOR < 7 || (ROCM_VERSION_MAJOR == 7 && ROCM_VERSION_MINOR < 2))?
Contributor
Author
There was a problem hiding this comment.
Indeed. Your comment + some further tweaks with adjustment to future lds intrinsic API are addressed in 2c739ab
|
|
||
| asm volatile("cp.async.wait_all;\n" ::); | ||
| #elif defined(USE_ROCM) && \ | ||
| (ROCM_VERSION_MAJOR <= 7 && ROCM_VERSION_MINOR < 2) && defined(__gfx950__) |
spcyppt
reviewed
Feb 24, 2026
| const uint4* row_v[kRowUnroll]; | ||
| int32_t idx_v[kRowUnroll]; | ||
| int32_t cache_idx_v[kRowUnroll]; | ||
| bool row_valid_v[kRowUnroll]; |
Contributor
There was a problem hiding this comment.
could you ensure all the changes only affect rocm?
Contributor
Author
There was a problem hiding this comment.
The whole block (lines 162-228) is under if is_rocm jinja guard
| } | ||
| {% if weighted %} | ||
| #pragma unroll | ||
| for (uint32_t inner_i = 0; inner_i < kRowUnroll; inner_i++) { |
Contributor
There was a problem hiding this comment.
Please guard the changes to be ROCM only? We see small regression in NVIDIA.
Contributor
Author
There was a problem hiding this comment.
The whole block (lines 162-228) is under if is_rocm jinja guard
cp_async_zfill_cg is async on Ampere+ and gfx950 but synchronous elsewhere. Inlining the sync fallback into the per-iteration row-load loop kills load pipelining (load->store dependency forces N waitcnts instead of one) and adds wave divergence on mixed-validity warps. Measured up to -19% BW on MI300 (gfx942) for weighted L=20/L=50. Wrap the row-store section in a #if matching the helper's dispatch: gfx950/Ampere keep the fused cp_async_zfill_cg loop; everything else gets the original two-loop pattern (load all -> masked store). Helper and gfx950 paths untouched. Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
…ents about pipelining of memory ops
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
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.
This PR implements direct HBM->LDS stores in tbe inference kernel. There are 2 major changes:
Due to pre-7.2 ROCm features, we are forced to used assembly inline to get 16B loads to work, so manual synchronization was added. In case of ROCm >= 7.2, we use proper intrinsics to handle memory synchronization.
This change brings ~10% performance boost on average for weighted and unweighted cases. We may try to push it further by doing async loads for indices weights.