Repository navigation
fix(stinkytofu): keep one split-barrier signal in flight per barrier id in the CDNA5 DAG - #14636
Merged
Merged
Conversation
…id in the CDNA5 DAG gfx1250 split barriers (s_barrier_signal / s_barrier_wait) hold one signal per barrier id in flight: after a signal, the next barrier on that id must be a wait. Barrier pairs only get DAG edges through their own LDS tokens, so two pairs on different tokens had no edge between them and the scheduler could nest them (sig A, sig B, wait B, wait A on id -1). The bf16 TN StreamK solutions removed in #14540 (MT256x256x64 MIWT16_4 and MT320x256x64 MIWT20_4, WG16_8_1) have three such sites each and race or hang. The scheduler now tracks the open barrier ids. A signal opens its id; a wait on it closes it. While an id is open, isBarrierEligibleNow() rejects every barrier on that id that is not a wait; the wait that is ready is that signal's own, since it shares the signal's tokens. Other ids are free, so -1 pairs inside an open -3 cluster signal stay legal. The barrier phase's drain path now uses isBarrierEligibleNow() like the promoted and cluster paths, so no path bypasses the rule. If only blocked barriers remain, pickOne() stops with an error instead of emitting the wrong order. raw_asm_token_annotation.s feeds nested pairs; its checks now require them to alternate (it fails on the old scheduler). The two #14540 kernels drop from 3 nested sites to 0, with the same 52 signals and 52 waits. stinkytofu ctest 1810/1810.
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
shared: stinkytofu
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. |
hcman2
approved these changes
Oct 29, 2026
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #14636 +/- ##
========================================
Coverage 77.27% 77.27%
========================================
Files 2892 2892
Lines 485461 485463 +2
Branches 73342 73342
========================================
+ Hits 375101 375104 +3
+ Misses 85603 85602 -1
Partials 24757 24757
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
darrenhsieh-amd
added a commit
that referenced
this pull request
Oct 29, 2026
… next barrier is due - develop's per-id pairing rule (#14636) holds a due barrier while a signal on its id is open; the barrier then waited for the drain at region end, pushing its tensor_load to the loop tail - dueBarrier() returns that signal's wait instead when it is eligible: the wait closes the id, so the due barrier group goes out next at its threshold - promoteDueBarrier() and extractForcedBarrier() both take dueBarrier(); a due barrier held for WMMA cover now also holds the barriers behind it instead of letting a later one go first - test: DueBarrierHeldByAnOpenSignalPromotesThatSignalsWait
1 task
Contributor
|
JIRA ID : AIHPBLAS-5229 |
geotseng-amd
added a commit
that referenced
this pull request
Oct 29, 2026
…tions" (#14650) ## Motivation #14540 removed two gfx1250 bf16 TN StreamK solutions because they computed wrong results and intermittently hung the GPU: - `MT256x256x64_..._MIWT16_4_..._WG16_8_1` - `MT320x256x64_..._MIWT20_4_..._WG16_8_1` The root cause was not in the logic files. It was in stinkytofu's CDNA5 DAG scheduler, which could nest gfx1250 split-barrier pairs on the same barrier id (`s_barrier_signal -1`, `s_barrier_signal -1`, `s_barrier_wait -1`, `s_barrier_wait -1`). Each kernel had three such sites. #14636 (AIHPBLAS-5375) fixed the scheduler: it allows only one signal in flight per barrier id, and it fails instead of emitting the wrong order. With that fix, these two kernels regenerate with 0 nested sites, down from 3 (same 52 signals / 52 waits). Since the kernels are now generated correctly, this PR restores the two solutions so the heuristic can select them again. ## Technical Details - Pure `git revert` of #14540 (5dccfe1). No other changes. - Restores the two solutions in both copies of `gfx1250_Cijk_Alik_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs.yaml`: - `asm_full/gfx1250/gfx1250/Origami/` - `asm_full/gfx1250/gfx1250_96cu/Origami/` Each copy goes from 698 back to 700 solutions, with the original `SolutionIndex` numbering. - Both files are byte-identical to their state before #14540. No other commit has touched them since, so the revert applies cleanly. - Depends on #14636, which is already in `develop` and in this branch's base. ## Test Plan Clean build of this branch (485eece, includes #14636) in a fresh worktree, on gfx1250 (ASIC revision 1) with ROCm 10.3.0a20261029. No `/opt/rocm` in the environment. - stinkytofu ctest. - hipBLASLt build (3 build configurations) and install layout check. - Smoke test with `HSA_DISABLE_GFX12_STRICT=0` and `=1`. - The three bf16 TN tests that failed or hung before #14540, run twice. - `hipblaslt-bench` over all 662 applicable solutions on two problem sizes, checking `norm_error`. - Full `hipblaslt-test`. ## Test Result | Step | Result | |---|---| | stinkytofu ctest | 1828/1828 passed | | hipBLASLt build | rc=0, 3 builds in 1266 s, layout OK | | Smoke (`=0` and `=1`) | 1152 passed in each run, loading `library/gfx1250` | | bf16 TN tests (3 tests × 2 runs) | 6/6 passed, no alarm timeout | | `hipblaslt-bench`, all 662 solutions × 2 problem sizes | 0 wrong results; both restored solutions (2033, 2118) have `norm_error` 0 | | Full `hipblaslt-test` | **69704 passed, 0 failed, 185 skipped** in 18 min, no alarm timeout | The full `hipblaslt-test` run includes the two restored solutions, so restoring them does not bring back the nightly failures on current `develop`. Reverts #14540. Fixed by #14636.
1 task
Contributor
TheRock Submodule Bump ActivityNewest first
|
darrenhsieh-amd
added a commit
that referenced
this pull request
Oct 30, 2026
JIRA ID : AIPYTORCH-1308 ## Motivation support MXBlockA/B = 128 and 2D MX scale tiles on gfx1250. Each thread's scale reaches the WMMA through a splat v_perm, and LDS returns data in order, so the scheduler needs to issue ds_loads in the order the WMMAs need them. - Supported MX block 128 configurations Scale block | MacroTile (free dim) | Requirement -- | -- | -- 1D 128 | any | none 2D [128, 128] | multiple of 128 (e.g. 128, 256) | 128 % (MI x VW) == 0 (always true for mxf8: MI 16, VW ≤ 8) 2D [128, 128] | divisor of 128 (e.g. 32, 64) | none 2D [128, 128] | other (e.g. 192, 224) | gcd(MT, 128) % (MI x VW) == 0 (mxf8: MT 192 → VW ≤ 4, MT 224 → VW ≤ 2); no UseSubtileImpl, PrefetchAcrossPersistent or DirectToVgprMXS; at most 2 LDS buffers; scale LDS read offsets below 64K 2D [free, K], free or K != 128 | - | rejected ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> - TensileLite host side: MXBlock can now be written as a list, plus the matching predicate, reference and data-generation changes. - TensileLite codegen: just-in-time splat v_perm, 2D TDM/GL2 address fixes, and reads for splat-fed tensors issuing first. - StinkyTofu scheduler: ds_loads ordered by when their WMMA needs them, a gate so later WMMAs' loads wait for earlier ones, and the new Step 3.5 in WmmaHideBudgetAnalysis that removes the burst of 52 back-to-back ds_loads. - (dropped, It's been addressed in #14636 )Bug: The CDNA5 scheduler could interleave split barriers (sig1 sig0 wait0 … wait1), letting tensor_load overwrite LDS still being read. - Fix: Block new barriers while an all-wave signal is open, pulling its wait forward so signal/wait strictly alternate. <!-- Explain any relevant testing done to verify this PR. --> ## Test Result tox passed (excluding existing failed sk cases like CI shows) hipblaslt-test passed ctest passed performance regression test done - all maf, mab, medium codegen are the same - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. [AIPYTORCH-1308]: https://amd-hub.atlassian.net/browse/AIPYTORCH-1308?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ --------- Co-authored-by: mengzcai <mengzcai@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Anne Yang <Anne.Yang@amd.com> Co-authored-by: hcman2 <menghung@amd.com> Co-authored-by: darrenhsieh-amd <darren.hsieh2@amd.com> Co-authored-by: dahsieh <dahsieh@amd.com>
ScottTodd
pushed a commit
that referenced
this pull request
Nov 2, 2026
…id in the CDNA5 DAG [bkc 20261027] (#14863) Cherry-picks: - #14636 ## Motivation Cherry-pick request for the BKC 26.14 release branch (`release/bkc/therock-10.3-20261027`). gfx1250 split barriers (`s_barrier_signal` / `s_barrier_wait`) hold one signal per barrier id in flight: after a signal, the next barrier on that id must be a wait. In the CDNA5 DAG, barrier pairs only get edges through their own LDS tokens, so the scheduler can nest two pairs on different tokens (signal A, signal B, wait B, wait A on the same id), which can race or hang. The fix makes the scheduler track open barrier ids and reject any non-wait barrier on an id that is still open. ## Technical Details - Cherry-pick (`git cherry-pick -x`) of the merged develop commit `bcf5892eec` (#14636), two files: `shared/stinkytofu/src/transforms/asm/dag/CDNA5.hpp` and `shared/stinkytofu/tests/filecheck/raw_asm_token_annotation.s`. - One trivial conflict in the `#include` block of `CDNA5.hpp`: this branch lacks the unrelated develop change that added `<span>`, so only `<set>` is included. The functional diff is identical to #14636. - Landed on develop after the 26.14 branch was cut (from the 2026-10-27 nightly), so it is not in the 26.14 snap build. ## Issue Tracking JIRA ID: AIHPBLAS-5375 ## Test Plan / Result Validated on develop in the original PR (#14636): the nested-pair filecheck test fails on the old scheduler and passes with the fix, the affected kernels go from 3 nested sites to 0 with unchanged signal/wait counts, and stinkytofu ctest passed 1810/1810. No additional testing was done on the release branch. ## Review @KKyang is the owner of the original change; please review this cherry-pick. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: KKyang <YANGWEN.HUANG@AMD.COM>
1 task done
ScottTodd
added a commit
to ROCm/TheRock
that referenced
this pull request
Nov 2, 2026
## Motivation * This includes ROCm/rocm-libraries#14863 which cherry-picks ROCm/rocm-libraries#14636 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/CONTRIBUTING.md.
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-5375
Problem
gfx1250 split barriers (
s_barrier_signal/s_barrier_wait) hold one signal per barrier id in flight: after a signal, the next barrier on that id must be a wait. In the CDNA5 DAG, barrier pairs only get edges through their own LDS tokens, so two pairs on different tokens have no edge between them and the scheduler can nest them:The two gfx1250 bf16 TN StreamK solutions removed in #14540 (
MT256x256x64 ... MIWT16_4 ... WG16_8_1,MT320x256x64 ... MIWT20_4 ... WG16_8_1) have three such sites each.Change
isBarrierEligibleNow()rejects every barrier on that id that is not a wait. The ready wait is that signal's own, since it shares the signal's tokens. Other ids are free, so-1pairs inside an open-3cluster signal stay legal.isBarrierEligibleNow()like the promoted and cluster paths, so no path bypasses the rule.pickOne()stops with an error instead of emitting the wrong order.Tests
raw_asm_token_annotation.sfeeds nested pairs; its checks now require them to alternate. It fails on the old scheduler.