Repository navigation
fix(hipblaslt): complete the K=64 prefetch in handwritten gfx950 kernels - #12679
Merged
newling merged 1 commit intoSep 29, 2026
Conversation
nakajee
approved these changes
Sep 29, 2026
nakajee
left a comment
nakajee
left a comment
Contributor
There was a problem hiding this comment.
The code change is reasonable.
OK to merge this change.
(as long as we keep these custom kernels)
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## users/newling/fix-gfx950-custom-final-loop-wait #12679 +/- ##
===================================================================================
- Coverage 70.57% 70.43% -0.14%
===================================================================================
Files 2932 2932
Lines 482397 482397
Branches 69542 69542
===================================================================================
- Hits 340429 339749 -680
- Misses 118174 118787 +613
- Partials 23794 23861 +67
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
newling
merged commit Sep 29, 2026
dd28da3
into
users/newling/fix-gfx950-custom-final-loop-wait
213 of 227 checks passed
newling
added a commit
that referenced
this pull request
Sep 29, 2026
…els (#12679) JIRA ID : AIHPBLAS-4866 ## Motivation Three handwritten gfx950 BF16/FP16 kernels can read local data share (LDS) before its global-memory loads complete at `K=64`. Skipping the second prefetch leaves only 16 loads, but the existing `vmcnt(24)` and `vmcnt(16)` waits assume 32 and do not guarantee their completion. ## Technical Details Add `s_waitcnt vmcnt(0)` on the path that skips the second prefetch, before the existing barriers and first LDS reads. The path with two prefetches bypasses the new wait. This follow-up is based on [#12676](#12676), which fixes a separate LDS-to-register dependency in the final compute loop. ## Test Plan Compare baseline, fixed and same-size NOP controls in gfx950 functional emulation at `M=N=256, K=64, batch=1`, with `alpha=1`, `beta=0` and `GSU=1`. Run the prefetch regressions and K=0/K=128/beta=1 controls. ## Test Result All three kernels go from **2 reports to 0**; replacing the new wait with `s_nop 0` restores both reports. All 19 emulator runs pass full-output numerical validation, and all nine focused regressions pass. Baseline numerical validation also passes; incorrect output was not reproduced. [Full detector trace, detailed results and reproduction commands](https://gist.github.com/newling/e688488a90362116cddc26ebb0a61927). ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium: changes kernel synchronization. Physical gfx950 validation and performance remain unmeasured; waitcheck cannot complete this DirectToLDS control-flow analysis.
newling
added a commit
that referenced
this pull request
Sep 29, 2026
…els (#12679) JIRA ID : AIHPBLAS-4866 ## Motivation Three handwritten gfx950 BF16/FP16 kernels can read local data share (LDS) before its global-memory loads complete at `K=64`. Skipping the second prefetch leaves only 16 loads, but the existing `vmcnt(24)` and `vmcnt(16)` waits assume 32 and do not guarantee their completion. ## Technical Details Add `s_waitcnt vmcnt(0)` on the path that skips the second prefetch, before the existing barriers and first LDS reads. The path with two prefetches bypasses the new wait. This follow-up is based on [#12676](#12676), which fixes a separate LDS-to-register dependency in the final compute loop. ## Test Plan Compare baseline, fixed and same-size NOP controls in gfx950 functional emulation at `M=N=256, K=64, batch=1`, with `alpha=1`, `beta=0` and `GSU=1`. Run the prefetch regressions and K=0/K=128/beta=1 controls. ## Test Result All three kernels go from **2 reports to 0**; replacing the new wait with `s_nop 0` restores both reports. All 19 emulator runs pass full-output numerical validation, and all nine focused regressions pass. Baseline numerical validation also passes; incorrect output was not reproduced. [Full detector trace, detailed results and reproduction commands](https://gist.github.com/newling/e688488a90362116cddc26ebb0a61927). ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium: changes kernel synchronization. Physical gfx950 validation and performance remain unmeasured; waitcheck cannot complete this DirectToLDS control-flow analysis.
newling
added a commit
that referenced
this pull request
Oct 5, 2026
…els (#12679) JIRA ID : AIHPBLAS-4866 ## Motivation Three handwritten gfx950 BF16/FP16 kernels can read local data share (LDS) before its global-memory loads complete at `K=64`. Skipping the second prefetch leaves only 16 loads, but the existing `vmcnt(24)` and `vmcnt(16)` waits assume 32 and do not guarantee their completion. ## Technical Details Add `s_waitcnt vmcnt(0)` on the path that skips the second prefetch, before the existing barriers and first LDS reads. The path with two prefetches bypasses the new wait. This follow-up is based on [#12676](#12676), which fixes a separate LDS-to-register dependency in the final compute loop. ## Test Plan Compare baseline, fixed and same-size NOP controls in gfx950 functional emulation at `M=N=256, K=64, batch=1`, with `alpha=1`, `beta=0` and `GSU=1`. Run the prefetch regressions and K=0/K=128/beta=1 controls. ## Test Result All three kernels go from **2 reports to 0**; replacing the new wait with `s_nop 0` restores both reports. All 19 emulator runs pass full-output numerical validation, and all nine focused regressions pass. Baseline numerical validation also passes; incorrect output was not reproduced. [Full detector trace, detailed results and reproduction commands](https://gist.github.com/newling/e688488a90362116cddc26ebb0a61927). ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium: changes kernel synchronization. Physical gfx950 validation and performance remain unmeasured; waitcheck cannot complete this DirectToLDS control-flow analysis.
newling
added a commit
that referenced
this pull request
Oct 7, 2026
…els (#12679) JIRA ID : AIHPBLAS-4866 ## Motivation Three handwritten gfx950 BF16/FP16 kernels can read local data share (LDS) before its global-memory loads complete at `K=64`. Skipping the second prefetch leaves only 16 loads, but the existing `vmcnt(24)` and `vmcnt(16)` waits assume 32 and do not guarantee their completion. ## Technical Details Add `s_waitcnt vmcnt(0)` on the path that skips the second prefetch, before the existing barriers and first LDS reads. The path with two prefetches bypasses the new wait. This follow-up is based on [#12676](#12676), which fixes a separate LDS-to-register dependency in the final compute loop. ## Test Plan Compare baseline, fixed and same-size NOP controls in gfx950 functional emulation at `M=N=256, K=64, batch=1`, with `alpha=1`, `beta=0` and `GSU=1`. Run the prefetch regressions and K=0/K=128/beta=1 controls. ## Test Result All three kernels go from **2 reports to 0**; replacing the new wait with `s_nop 0` restores both reports. All 19 emulator runs pass full-output numerical validation, and all nine focused regressions pass. Baseline numerical validation also passes; incorrect output was not reproduced. [Full detector trace, detailed results and reproduction commands](https://gist.github.com/newling/e688488a90362116cddc26ebb0a61927). ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium: changes kernel synchronization. Physical gfx950 validation and performance remain unmeasured; waitcheck cannot complete this DirectToLDS control-flow analysis.
newling
added a commit
that referenced
this pull request
Oct 7, 2026
…ps (#12676) Three handwritten gfx950 kernels in hipBLASLt's production solution tables can consume matrix operands before their asynchronous loads are guaranteed to have completed. The affected kernels are the two BF16 TN variants and the FP16 TN variant with a 256×256×64 tile. The optimized final loop runs for `alpha=1`, `beta=0`, an unsplit reduction (`GSU=1`), and complete output tiles. The proposed update uses `lgkmcnt(7)` before matrix fused multiply-add (MFMA) index 0 and `lgkmcnt(8)` before index 8, matching the ordinary no-load loop. These waits complete each group of operands before use while allowing later reads to remain pending. This PR also includes the K=64 fix from [#12679](#12679). When the second global-memory prefetch is skipped, `vmcnt(0)` completes the sole prefetch before the barrier and subsequent local data share (LDS) reads. RocJITsu reports ten final-loop hazards per kernel on develop and zero with the two partial waits. Replacing both waits with NOPs restores ten reports; removing either wait alone also restores reports. The K=64 controls give two reports before the fix and zero afterward. Tests cover all three kernels, reduction-size and output-edge cases, padded strides, and batch=1024; six fresh detector runs after the rebase also report zero. All runs pass numerical validation, and the custom-kernel regression tests pass. These are functional-emulation results; physical gfx950 correctness and performance have not been measured. [Complete detector trace, results, reproduction recipe and proposed patch](https://gist.github.com/newling/d8ed49e5f33de16938482e5b2b61c13f). **Publication status:** GitHub currently contains the earlier full-drain version (`lgkmcnt(0)` at final-loop entry). The validated partial-wait revision above is local at `2f2c2c80dc3` and awaits push; its exact patch is in the linked gist. JIRA ID : AIHPBLAS-4866
1 task
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
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.
JIRA ID : AIHPBLAS-4866
Motivation
Three handwritten gfx950 BF16/FP16 kernels can read local data share (LDS) before its global-memory loads complete at
K=64. Skipping the second prefetch leaves only 16 loads, but the existingvmcnt(24)andvmcnt(16)waits assume 32 and do not guarantee their completion.Technical Details
Add
s_waitcnt vmcnt(0)on the path that skips the second prefetch, before the existing barriers and first LDS reads. The path with two prefetches bypasses the new wait.This follow-up is based on #12676, which fixes a separate LDS-to-register dependency in the final compute loop.
Test Plan
Compare baseline, fixed and same-size NOP controls in gfx950 functional emulation at
M=N=256, K=64, batch=1, withalpha=1,beta=0andGSU=1. Run the prefetch regressions and K=0/K=128/beta=1 controls.Test Result
All three kernels go from 2 reports to 0; replacing the new wait with
s_nop 0restores both reports. All 19 emulator runs pass full-output numerical validation, and all nine focused regressions pass. Baseline numerical validation also passes; incorrect output was not reproduced.Full detector trace, detailed results and reproduction commands.
Submission Checklist
Risk level
Medium: changes kernel synchronization. Physical gfx950 validation and performance remain unmeasured; waitcheck cannot complete this DirectToLDS control-flow analysis.