Repository navigation
fix(stinkytofu): keep one split-barrier signal in flight per barrier id in the CDNA5 DAG [bkc 20261027] - #14863
Merged
ScottTodd merged 1 commit intoNov 2, 2026
Conversation
…id in the CDNA5 DAG (ROCm#14636) JIRA ID : AIHPBLAS-5375 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: ``` s_barrier_signal -1 (pair A) ... s_barrier_signal -1 (pair B) <- second signal while A is open s_barrier_wait -1 (pair B) ... s_barrier_wait -1 (pair A) ``` The two gfx1250 bf16 TN StreamK solutions removed in ROCm#14540 (`MT256x256x64 ... MIWT16_4 ... WG16_8_1`, `MT320x256x64 ... MIWT20_4 ... WG16_8_1`) have three such sites each. - The scheduler 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 ready wait 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 ROCm#14540 kernels, regenerated from the pre-ROCm#14540 logic file: nested sites 3 -> 0 in each, with the same 52 signals and 52 waits. - stinkytofu ctest 1810/1810. (cherry picked from commit bcf5892)
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
shared: stinkytofu
labels
Author
assistant-librarian
Bot
added
the
external contribution
Author
|
@KKyang @hcman2, since you are the author and the approving reviewer of the original change (#14636), could one of you please approve this cherry-pick on GitHub? It targets the 26.14 branch ( |
ThanHenderson
approved these changes
Nov 2, 2026
babakpst
approved these changes
Nov 2, 2026
ScottTodd
approved these changes
Nov 2, 2026
ScottTodd
merged commit Nov 2, 2026
9ca7988
into
ROCm:release/bkc/therock-10.3-20261027
32 of 38 checks passed
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.
Cherry-picks:
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
git cherry-pick -x) of the merged develop commitbcf5892eec(fix(stinkytofu): keep one split-barrier signal in flight per barrier id in the CDNA5 DAG #14636), two files:shared/stinkytofu/src/transforms/asm/dag/CDNA5.hppandshared/stinkytofu/tests/filecheck/raw_asm_token_annotation.s.#includeblock ofCDNA5.hpp: this branch lacks the unrelated develop change that added<span>, so only<set>is included. The functional diff is identical to fix(stinkytofu): keep one split-barrier signal in flight per barrier id in the CDNA5 DAG #14636.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