fix(stinkytofu): put the LD_SCALE blocked-cycle model behind a flag, default off - #11300
Closed
KKyang wants to merge 1 commit into
Closed
fix(stinkytofu): put the LD_SCALE blocked-cycle model behind a flag, default off#11300KKyang wants to merge 1 commit into
KKyang wants to merge 1 commit into
Conversation
…default off `v_wmma_scale_*` / `v_wmma_scale16_*` encode a back-to-back VOP3P pair whose LD_SCALE runs in the last cycle of the latency window, and #11156 taught the CDNA5 DAG scheduler to leave that cycle empty. The model is right about the hardware, but it costs an issue slot in *every* scale-WMMA window, and the scheduler has no way to tell which windows can afford that: it fills one window at a time and only learns the shadow ran out after the fact, by which point the overflow is already sitting in the next window. Put the model behind `dagFeatures.enableWmmaBlockedScaleCycles`, default off, so the shipping schedule is the pre-#11156 one until the scheduler can make that call per window. A follow-up turns the default back on once it can. ## The gate `activeWmmaBlockedScale_` is the single value every blocked-cycle path reads, so gating its one assignment in `pickOneFromWMMA` reverts all of them at once: - `advanceTime`'s roll past a blocked cycle - `computeValuAdvanceCycles`' co-issue slot count - `freeCoIssueSpace`'s hidden-stall shadow All-or-nothing on purpose. `advanceTime`'s roll is what guarantees the timeline never comes to rest on a blocked cycle, and `isValuPickable()` leans on that guarantee instead of re-checking the mask -- so enabling the other two without the roll would drop a VALU straight back into the LD_SCALE cycle, which is the bug #11156 fixed. `HwInstDesc::blockedScaleMask` and the tablegen declarations are untouched: the hardware fact stays in the tables, only the scheduler's use of it is gated. ## Reaching the flag - `ModuleOptions.EnableWmmaBlockedScaleCycles` -> `dagFeatures` (Gfx1250Backend) - `stinkytofu-opt --enable-wmma-blocked-scale` ## Test plan - [x] `ctest` -- 1102/1103. The one failure, `FileCheck.asm_move_propagation`, is pre-existing on develop and unrelated (it runs only `--AsmMovePropagationPass`; `CDNA5.hpp` is included solely by `StinkyDAGSchedulerPass.cpp`). - [x] New `dag_wmma_blocked_scale_default_off.stir` pins the default: the scale pair fits 7 SALU in its latency-8 window, same as the plain WMMA, and the LD_SCALE cycle is filled like any other. - [x] `dag_wmma_scale_blocked_tail.stir` keeps the model covered by passing `--enable-wmma-blocked-scale`; with it on, the scale pair still fits 6 against the plain WMMA's 7. - [x] No other test changed behaviour when the default flipped, which matches #11156's own note that the two scale-WMMA DAG unit tests pass either way. ## Note for reviewers While the default is off, issue #11171 is un-fixed in shipped kernels: the LD_SCALE cycle is fillable again. That is the deliberate trade -- a wasted issue slot in every scale window was the worse of the two -- and it is temporary. The comments on the knob say so, so the default is not read as "the model was wrong".
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. |
KKyang
added a commit
that referenced
this pull request
Aug 26, 2026
Groundwork for turning the LD_SCALE blocked-cycle model back on. #11300 put that model behind a flag and defaulted it off because it costs an issue slot in every scale-WMMA window and the scheduler has no way to tell which windows can afford that. This lands the measurement that will let it tell. No behaviour change: the pre-scan measures and reports, nothing gates on it, and the flag stays off by default. Scheduling output is byte-identical to #11300. ## Why a per-window budget "Can this region hide its work" is the wrong question, because the work is not interchangeable. A ds_load feeding WMMA 6 has to be issued before WMMA 6 whether or not a shadow has room for it; an independent SALU can wait forever. So the question each window answers is "may I issue more than my slot?". `scanRegionHideBudget()` runs from `onInitRegion` and reads the region DAG: 1. Number the matrix ops in program order -- those are the windows. For each, `capacityCycles` is every window cycle after its own issue slot that `blockedScaleMask` does not reserve, and `capacityValu` the subset `coIssueWindow` also names. 2. Give every node the index of the earliest WMMA that transitively depends on it -- its deadline. RegionDAG ids are program indices and its edges run strictly forward, so one reverse sweep settles every node, O(V+E). 3. Charge each filler's issue cycles to its deadline. Nodes no WMMA depends on are floating: they still compete for window space at pick time, but they can always be deferred, so they never force anything. 4. Walk the deadlines in order. Where the shadow before WMMA i falls short, the shortfall is granted to window i-1, the latest window that can still meet the deadline, as `extraIssue`. Granting late is the tightest answer; a consumer may spread the same total earlier, which issues the loads sooner, but not later. Query: `wmmaExtraIssueAllowance(inst)`. Demand is a deliberate lower bound: VALU is counted at its `issueCycles` even though a VALU inside a window can cost more (`computeValuAdvanceCycles` walks to the next co-issue bit), so an exception is raised only where one is owed beyond doubt. Barriers are excluded -- they are not window fillers, and their placement is the barrier-threshold machinery's job. ## Interface `ReadyQueue::onInitRegion` gained a `const RegionDAG&`. The pass already builds that graph three lines earlier and was dropping it on the floor; handing it over means region pre-scans reason about the same graph the scheduler drains instead of rebuilding a weaker view of it. One override, one call site, forward-declared in ReadyQueue.hpp to keep the include one-way. ## Observability - `--remarks` reports only windows that owe an overrun, naming them. - `--debug-pass StinkyDAGSchedulerPass` dumps capacity/extraIssue per window. ## Test plan - [x] `ctest` -- 1103/1104. The one failure, `FileCheck.asm_move_propagation`, is pre-existing on develop and unrelated (it runs only `--AsmMovePropagationPass`). - [x] New `dag_wmma_hide_budget_prescan.stir`. Three regions: - 20 ds_loads all feeding WMMA #3, plain FP8: 3 windows x 7 = 21 cycles of shadow, they fit, region stays quiet. - Identical dependency structure on the scale pair: 3 x 6 = 18, they no longer fit, and window #2 is told to overrun by 2. Same IR, only the opcode differs -- that isolates what the blocked-cycle model costs when the work has a deadline. - tensor_load -> ds_load -> WMMA: the deadline reaches through LDS, not just registers (StinkyBuildImplicitDependencyPass materialises memory tokens as RegType::LDS pseudo-registers, so the DAG carries the edge), so the tensor_loads count as deadlined rather than floating. It runs with `--enable-wmma-blocked-scale`, since the plain/scale contrast is exactly what the model costs and there is nothing to contrast without it. - [x] No existing test changed: the pre-scan gates nothing. ## Notes for reviewers - The pre-scan honours `enableWmmaBlockedScaleCycles` too, so its budget describes the scheduler that will actually run rather than the instruction tables. - Unrelated but adjacent, and not addressed here: `InsertCoexecHazardPass` reads the *unresolved* `getHwInstDesc()->coIssueWindow` while the DAG uses the matrix-format-resolved `inst->coIssueWindow`. On gfx1250 FP4/FP4 those are 0x00C8 and 0x0008, so it inserts two surplus `v_nop` per dependent WMMA pair. Pre-existing, worth its own fix.
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #11300 +/- ##
========================================
Coverage 69.68% 69.68%
========================================
Files 2803 2803
Lines 461507 461507
Branches 68015 68015
========================================
+ Hits 321570 321571 +1
Misses 116671 116671
+ Partials 23266 23265 -1
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
KKyang
added a commit
that referenced
this pull request
Aug 27, 2026
Groundwork for turning the LD_SCALE blocked-cycle model back on. #11300 put that model behind a flag and defaulted it off because it costs an issue slot in every scale-WMMA window and the scheduler has no way to tell which windows can afford that. This lands the measurement that will let it tell. No behaviour change: the analysis measures and reports, nothing gates on it, and the flag stays off by default. Scheduling output is byte-identical to #11300. ## Why a per-window budget "Can this region hide its work" is the wrong question, because the work is not interchangeable. A ds_load feeding WMMA 6 has to be issued before WMMA 6 whether or not a shadow has room for it; an independent SALU can wait forever. So the question each window answers is "may I issue more than my slot?". `analyzeWmmaHideBudget()` reads the region DAG: 1. Number the matrix ops in program order -- those are the windows. For each, `capacityCycles` is every window cycle after its own issue slot that `blockedScaleMask` does not reserve, and `capacityValu` the subset `coIssueWindow` also names. 2. Give every node the index of the earliest WMMA that transitively depends on it -- its deadline. RegionDAG ids are program indices and its edges run strictly forward, so one reverse sweep settles every node, O(V+E). 3. Charge the issue cycles of each filler to its deadline. Nodes no WMMA depends on are floating: they still compete for window space at pick time, but they can always be deferred, so they never force anything. 4. Walk the deadlines in order. Where the shadow before WMMA i falls short, the shortfall is granted to window i-1, the latest window that can still meet the deadline, as `extraIssue`. Granting late is the tightest answer; a consumer may spread the same total earlier, which issues the loads sooner, but not later. Demand is a deliberate lower bound: VALU is counted at its `issueCycles` even though a VALU inside a window can cost more (`computeValuAdvanceCycles` walks to the next co-issue bit), so an overrun is only ever demanded where one is owed beyond doubt. Barriers are excluded -- they are not window fillers, and placing them is the barrier-threshold work, not this. ## Layout The analysis is its own unit under `analysis/asm/`, next to the other asm analyses, rather than more free functions in `CDNA5.hpp` -- which is already 2000 lines of scheduler. The ready queue keeps only what it needs to be a consumer: the result, `wmmaExtraIssueAllowance()`, and the call. Its `isBlockedCycle()` now delegates to `isBlockedWindowCycle()` in the analysis, so the end-anchored-mask convention has one definition instead of two. `ReadyQueue::onInitRegion` gained a `const RegionDAG&`. The pass already builds that graph three lines earlier and was dropping it on the floor; handing it over means region pre-scans reason about the same graph the scheduler drains instead of rebuilding a weaker view of it. One override, one call site, forward-declared in ReadyQueue.hpp to keep the include one-way. ## Observability - `--remarks` reports only windows that owe an overrun, naming them. - `--debug-pass StinkyDAGSchedulerPass` dumps capacity/extraIssue per window. ## Test plan - [x] `ctest` -- 1103/1104. The one failure, `FileCheck.asm_move_propagation`, is pre-existing on develop and unrelated (it runs only `--AsmMovePropagationPass`). - [x] New `dag_wmma_hide_budget_prescan.stir`. Three regions: - 20 ds_loads all feeding WMMA #3, plain FP8: 3 windows x 7 = 21 cycles of shadow, they fit, region stays quiet. - Identical dependency structure on the scale pair: 3 x 6 = 18, they no longer fit, and window #2 is told to overrun by 2. Same IR, only the opcode differs -- that isolates what the blocked-cycle model costs when the work has a deadline. - tensor_load -> ds_load -> WMMA: the deadline reaches through LDS, not just registers (StinkyBuildImplicitDependencyPass materialises memory tokens as RegType::LDS pseudo-registers, so the DAG carries the edge), so the tensor_loads count as deadlined rather than floating. It runs with `--enable-wmma-blocked-scale`, since the plain/scale contrast is exactly what the model costs and there is nothing to contrast without it. - [x] No existing test changed: the analysis gates nothing. ## Notes for reviewers - The analysis honours `enableWmmaBlockedScaleCycles` too, so its budget describes the scheduler that will actually run rather than the instruction tables. - `RegionDAG.hpp` is private to the scheduler (it lives under `src/`, not the installed include tree), so the analysis reaches it the same way `StinkyDAGSchedulerPass.cpp` does: relatively. Only the .cpp needs it; the header forward-declares. - Unrelated but adjacent, and not addressed here: `InsertCoexecHazardPass` reads the *unresolved* `getHwInstDesc()->coIssueWindow` while the DAG uses the matrix-format-resolved `inst->coIssueWindow`. On gfx1250 FP4/FP4 those are 0x00C8 and 0x0008, so it inserts two surplus `v_nop` per dependent WMMA pair. Pre-existing, worth its own fix.
Contributor
Author
|
Closing: not needed. The LD_SCALE blocked-cycle model from #11156 stays on unconditionally — no flag, no default flip, so issue #11171 is never un-fixed on develop. #11301 no longer stacks on this branch. It is rebased straight onto develop and carries only the per-window hide budget measurement, behind its own knob ( The per-window decision this PR was groundwork for still lands the same way: #11301 measures, a follow-up wires the budget into the pick paths. |
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
1 participant
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.
Warning
This is not the finished fix. It is a stopgap so the lost issue slot can be
switched off while the real fix is built. The model being off is a temporary state:
issue #11171 is un-fixed in shipped kernels for as long as this sits on develop
alone, because the LD_SCALE cycle becomes fillable again.
The finished fix decides per window instead of globally: a scale WMMA window gives up
its LD_SCALE cycle unless the work a later WMMA depends on has nowhere else to go.
#11301 lands the measurement that makes that decision possible; a follow-up wires it
into the scheduler and turns the default back on.
v_wmma_scale_*/v_wmma_scale16_*encode a back-to-back VOP3P pair whoseLD_SCALE runs in the last cycle of the latency window, and #11156 taught the
CDNA5 DAG scheduler to leave that cycle empty. The model is right about the
hardware, but it costs an issue slot in every scale-WMMA window, and the
scheduler has no way to tell which windows can afford that: it fills one window
at a time and only learns the shadow ran out after the fact, by which point the
overflow is already sitting in the next window.
Put the model behind
dagFeatures.enableWmmaBlockedScaleCycles, default off, sothe shipping schedule is the pre-#11156 one until the scheduler can make that
call per window. A follow-up turns the default back on once it can.
The gate
activeWmmaBlockedScale_is the single value every blocked-cycle path reads, sogating its one assignment in
pickOneFromWMMAreverts all of them at once:advanceTime's roll past a blocked cyclecomputeValuAdvanceCycles' co-issue slot countfreeCoIssueSpace's hidden-stall shadowAll-or-nothing on purpose.
advanceTime's roll is what guarantees the timelinenever comes to rest on a blocked cycle, and
isValuPickable()leans on thatguarantee instead of re-checking the mask -- so enabling the other two without
the roll would drop a VALU straight back into the LD_SCALE cycle, which is the
bug #11156 fixed.
HwInstDesc::blockedScaleMaskand the tablegen declarations are untouched: thehardware fact stays in the tables, only the scheduler's use of it is gated.
Reaching the flag
ModuleOptions.EnableWmmaBlockedScaleCycles->dagFeatures(Gfx1250Backend)stinkytofu-opt --enable-wmma-blocked-scaleTest plan
ctest-- 1102/1103. The one failure,FileCheck.asm_move_propagation,is pre-existing on develop and unrelated (it runs only
--AsmMovePropagationPass;CDNA5.hppis included solely byStinkyDAGSchedulerPass.cpp).dag_wmma_blocked_scale_default_off.stirpins the default: the scalepair fits 7 SALU in its latency-8 window, same as the plain WMMA, and the
LD_SCALE cycle is filled like any other.
dag_wmma_scale_blocked_tail.stirkeeps the model covered by passing--enable-wmma-blocked-scale; with it on, the scale pair still fits 6against the plain WMMA's 7.
fix(stinkytofu): block the LD_SCALE cycle of scale WMMA co-issue windows #11156's own note that the two scale-WMMA DAG unit tests pass either way.
Note for reviewers
While the default is off, issue #11171 is un-fixed in shipped kernels: the
LD_SCALE cycle is fillable again. That is the deliberate trade -- a wasted issue
slot in every scale window was the worse of the two -- and it is temporary. The
comments on the knob say so, so the default is not read as "the model was wrong".
First of a stack. #11301 adds the per-window hide budget on top of this branch and
changes no behaviour; the default returns only once that budget actually drives the
decision per window.
🤖 Generated with Claude Code