Repository navigation
fix(hipblaslt): wait for LDS operands in handwritten gfx950 final loops - #12676
Merged
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Collaborator
|
can we replace those custom kernels with the CMS kernels? Do we still need those? |
nakajee
reviewed
Sep 29, 2026
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (38.92%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #12676 +/- ##
========================================
Coverage 71.93% 71.93%
========================================
Files 2863 2863
Lines 472701 472701
Branches 69827 69827
========================================
Hits 340000 340000
Misses 109118 109118
Partials 23583 23583
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
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.
nakajee
approved these changes
Sep 29, 2026
nakajee
left a comment
nakajee
left a comment
Contributor
There was a problem hiding this comment.
Thanks for fixing the issue.
The fix looks OK.
Optimizing waitcnt for Opt NLL only will not make much difference and no need to add it for now.
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
force-pushed
the
users/newling/fix-gfx950-custom-final-loop-wait
branch
from
September 29, 2026 18:10
dd28da3 to
7786214
Compare
newling
enabled auto-merge (squash)
September 29, 2026 18:10
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
force-pushed
the
users/newling/fix-gfx950-custom-final-loop-wait
branch
from
September 29, 2026 22:36
7786214 to
e5db896
Compare
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
force-pushed
the
users/newling/fix-gfx950-custom-final-loop-wait
branch
from
October 5, 2026 15:33
e5db896 to
3dad4dd
Compare
…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
force-pushed
the
users/newling/fix-gfx950-custom-final-loop-wait
branch
from
October 7, 2026 16:00
3dad4dd to
8c78825
Compare
Contributor
TheRock Submodule Bump ActivityNewest first
|
newling
added a commit
that referenced
this pull request
Oct 9, 2026
## Description Fix parsing bug in rocisa, and is so doing eliminate a race condition (see [here](https://gist.github.com/0d7982619e7a0b927bd5e780e98f5e94)) On gfx950 (with hipBLASLt 1.4.1 and on develop) this buggy kernel was selected by ``` hipblaslt-bench --precision bf16_r --compute_type f32_r --transA N --transB T -m 3072 -n 1152 -k 8192 --alpha 1 --beta 0 ``` The equivalent pytorch operation is ``` torch.randn(8192, 1152, device="cuda", dtype=torch.bfloat16).T @ torch.randn(8192, 3072, device="cuda", dtype=torch.bfloat16) ``` While it is clear from the trace from the race detection tool that a real race **was** present, the number of instructions between the load from LDS and the read from the register is very large, so I wouldn't call this a scary/close race... ## Details (from AI) hipBLASLt generates assembly for GPU matrix-multiplication kernels. It builds parts of that assembly from reusable instruction templates, called macros. Arguments tell each macro which loop stage it represents, and conditions inside the macro choose which instructions to emit. The rocisa library evaluates these conditions during kernel generation, before the assembly is compiled. In the affected kernels, computation overlaps memory transfers: while the GPU multiplies one set of inputs, it prepares another. Loads from shared memory, called LDS, place inputs in registers for the matrix instructions to consume. Those loads are asynchronous, so the generated code must wait for them to finish before using their results. The required waits depend on whether the loop is still prefetching inputs or is finishing the remaining work. The scheduler expresses that distinction with conditions such as: ```text \useGR == 0 && \usePLR == 1 ``` Here, `useGR` enables global-memory reads and `usePLR` enables local-memory read prefetching. Both arguments are zero in the final loop, so this condition must be false. `evalMacroCondition()` evaluates conditions in two stages. First, it appends each comparison's result, `0` or `1`, to a vector named `results`, inserting `2` as a marker for each `&&`. For the example above, this should produce `[1, 2, 0]`, representing `true && false`. A second loop starts with the first result and combines a later result only when it encounters an `&&` marker. The bug is in recording those markers. The parser's `tokenIdx` cycles through positions 0–3: left operand, comparison operator, right operand, and `&&`. However, the conjunction handler was in unreachable `case 4`. Both comparisons were evaluated, but the vector contained `[1, 0]` instead of `[1, 2, 0]`. With no marker, the second loop never combined the false result with the first result, so the condition incorrectly returned true. Changing the handler to `case 3` records the missing markers and restores the final-loop branch already specified by the scheduler. The incorrect branch selected weaker waits, allowing matrix instructions to consume registers before their LDS loads had completed. The observed effect is in gfx950 BF16 and FP16 kernels with A untransposed, B transposed, and a 192×320×64 tile. Regenerating these kernels changes five waits per kernel. For example: ```diff -s_waitcnt lgkmcnt(4) +s_waitcnt lgkmcnt(0) ``` The old instruction allows four operations counted by `lgkmcnt` to remain outstanding; the new one drains that counter, ensuring the LDS loads have completed. The [assembly excerpt and full race reports](https://gist.github.com/0d7982619e7a0b927bd5e780e98f5e94) show the load, waits, and consuming matrix instruction together. The new macro regressions fail before the fix and pass afterward. Regenerated BF16/FP16 kernels assemble and pass numerical validation under rocJITsu functional GPU emulation. At M=3072, N=1152, K=8192, the BF16 kernel's two reported hazards disappear. Physical-GPU correctness and performance validation remain pending. Related investigation: #12676 addresses final-loop waits in handwritten kernels. JIRA ID : N/A — no separate ticket has been assigned for this finding.
1 task
tony-davis
added a commit
that referenced
this pull request
Oct 30, 2026
#12676 added the missing LDS operand waits to the handwritten gfx950 MT256x256x64 final loops, which fixes the wrong bf16 TN values at K = 64. Davood's retest on 59fb065 found no issue, so matmul_fast_check_short_k runs again on gfx950 as a regression gate. Co-authored-by: Cursor <cursoragent@cursor.com>
Merged
1 task done
tony-davis
added a commit
that referenced
this pull request
Nov 5, 2026
… by #14147 and #12676 (#14390) JIRA ID : ROCM-32343, ROCM-32344 ## Motivation The gfx950 custom-schedule `MT256x160x64` kernels used to write wrong values in the TN every-solution tests (ROCM-32343 for bf16, ROCM-32153 for f16). We quarantined the failing cases in `known_bugs.yaml` so the nightlies could stay green. #14147 fixed the cause: the schedule's prologue issued the A loads before the B loads, and the main loop's partial wait counts assumed B first. That fix is on develop, so the quarantines have done their job. This PR removes the four entries they left behind, so the cases run again as real regression gates. The ROCM-32344 entry can go too. It quarantined `matmul_fast_check_short_k` (bf16 TN) because the custom `MT256x256x64` kernels gave wrong values at K = 64. #12676 fixed those kernels on Oct 7, and the quarantine landed a day later in #13022. Davood's retest on 59fb065 found no issue. ## Technical Details - `known_bugs.yaml` loses the two bf16 entries (`matmul_fast_check_every_solution_repeat` and `matmul_fast_check_streamk_every_solution_repeat`, `a_type: bf16_r`) and the two f16 entries (the same two names, `a_type: f16_r`, `M: 4099`, TN) on `gfx950`, along with their comments. - `known_bugs.yaml` also loses the ROCM-32344 entry (`matmul_fast_check_short_k`, `a_type: bf16_r`, `transA: T`) on `gfx950`, along with its comment. - Test data only. No library, kernel, or client code changes. ## Test Plan - Compare the test records `hipblaslt_gentest.py` generates on develop and on this branch. Only the previously quarantined cases should change. - Build develop plus this change for gfx950 and run the formerly quarantined cases on MI350X and MI355X. ## Test Result - `hipblaslt_gentest.py` generates the same 143,631 records on both sides. 24 records change, and all of them just lose `gfx950` from `known_bug_platforms`: the two every-solution names at 9 records each (8 bf16 plus the f16 TN 4,099 case) and `matmul_fast_check_short_k` at 6 (bf16 TN, K = 1, 15, 16, 33, 64 and 129). Nothing else changes. - MI350X (Device 75a0): `*fast_check_every_solution_repeat*` and `*fast_check_streamk_every_solution_repeat*` (without the large-K and xf32 variants) ran 48 tests, 48 passed. - MI355X (Device 75a3): the same filter ran 48 tests, 48 passed. - For comparison, before the fix my A/B on #13276 showed the bf16 TN defaults wrong in 33 of 41 runs on MI350X on a develop that predated #14147, and 0 of 70 with the fix. - I did not run these on MI350P (Device 75a8). The cluster has none. The 75a8 tables carry copies of the same schedule (2 f16 and 19 bf16 solutions), which #14147 should fix the same way. The first check of those is the MI350P nightly after TheRock's rocm-libraries pin moves past #14147 (it is at 565c4c7 today). - Required checks passed before this update. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test data only, and these cases are nightly-category, so PR checks are unaffected. If the MI350P nightly still shows failures on the 75a8 copies, we put the quarantine back for that device's cases and look at those kernels again. MI350X and MI355X already pass. ## Risk Acceptance / Waivers None needed. Related: ROCM-32153, #13269, #13793, #13940 (closed), #14147. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
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.
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 useslgkmcnt(7)before matrix fused multiply-add (MFMA) index 0 andlgkmcnt(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. 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.
Publication status: GitHub currently contains the earlier full-drain version (
lgkmcnt(0)at final-loop entry). The validated partial-wait revision above is local at2f2c2c80dc3and awaits push; its exact patch is in the linked gist.JIRA ID : AIHPBLAS-4866