feat(stinkytofu): per-window WMMA hide budget analysis - #11301
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
shared: stinkytofu
labels
4 tasks
KKyang
changed the title
KKyang
marked this pull request as ready for review
August 26, 2026 13:20
KKyang
force-pushed
the
users/kkyang/wmma-hide-budget-prescan
branch
from
August 26, 2026 13:20
e71a6d3 to
2154251
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (35.22%) 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 #11301 +/- ##
===========================================
- Coverage 69.32% 69.32% -0.00%
===========================================
Files 2810 2810
Lines 465520 465534 +14
Branches 68685 68687 +2
===========================================
Hits 322706 322706
- Misses 119250 119264 +14
Partials 23564 23564
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
KKyang
force-pushed
the
users/kkyang/wmma-hide-budget-prescan
branch
from
August 27, 2026 00:04
2154251 to
96d541e
Compare
KKyang
changed the title
hcman2
approved these changes
Aug 31, 2026
KKyang
force-pushed
the
users/kkyang/wmma-hide-budget-prescan
branch
from
August 31, 2026 07:53
96d541e to
6b545d2
Compare
github-actions
Bot
added
github actions
project: miopen-provider
project: hip-kernel-provider
project: hipdnn-integration-tests
rocke-lib: attention
✅ 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. |
Groundwork for teaching the CDNA5 scheduler which WMMA windows can afford to give up an issue slot and which cannot. This lands the measurement; nothing gates on it yet, and it does not run unless asked for. No behaviour change. The pre-scan is off by default, and with it on the scheduling output is byte-identical -- it measures and reports, it does not decide. ## 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. ## The knob dagFeatures.enableWmmaHideBudgetPrescan, default false, reachable only through stinkytofu-opt --enable-wmma-hide-budget-prescan. There is deliberately no ModuleOptions mirror, so TensileLite cannot turn it on -- same treatment as mergeBarrierThreshold. The budget is a pure measurement so far, so a production compile would pay for a verdict nobody reads. ## 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 the gated call: the budget is a local in onInitRegion, reported and dropped, because nothing in the pick paths reads it yet. The follow-up that calls extraIssueFor() from pickOne() is what promotes it to per-region state -- until then a member would be state nobody reads, holding StinkyInstruction* past the region boundary. Its isBlockedCycle() now delegates to isBlockedWindowCycle() in the analysis header, so the end-anchored-mask convention has one definition instead of two; that helper is inline because advanceTime(), computeValuAdvanceCycles() and freeCoIssueSpace() call it once per window cycle. The queue needs the region DAG to answer "which loads does this WMMA wait on". The pass already builds that graph 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 from def-use chains, which carry RAW edges only and are function- rather than region-scoped. It is passed before the region's policy edges are merged in, so a pre-scan sees real data dependencies rather than heuristic orderings. Rather than make that a fifth positional parameter, onInitRegion now takes a RegionDependencies: the dependency edges the region already has (the DAG, in) and the ones the queue would like added (requestedConstraints, out, previously its own parameter). Both are edges of the same graph, which is what makes them one argument rather than a bag. It also retires a hazard the old signature carried -- three of the five parameters were IRList::iterator, so transposing two of them compiled silently. One override, one call site; RegionDAG stays 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 - New dag_wmma_hide_budget_prescan.stir, four 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 LD_SCALE 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. - v_wmma_f32_16x16x4_f32, which carries coIssueWindow 0x0000: capacityValu is 0 there, but nothing was blocked, so the LD_SCALE remark must stay silent. - ctest 1130/1130, no failures (112 FileCheck tests included, -Werror build). - Verified off by default: without the flag the region emits no remarks, and --print-output is byte-identical with the flag on and off. - No existing test changed: the analysis gates nothing. ## Notes for reviewers - The "lost every VALU co-issue slot" remark counts only windows that HAD co-issue slots and lost them to blockedScaleMask. capacityValu == 0 alone does not mean LD_SCALE took anything: v_wmma_f32_16x16x4_f32 carries coIssueWindow 0x0000, so it never had a VALU slot and its window still offers 15 issue cycles to SALU and memory. Pinned by the last function in the FileCheck test. - 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. A follow-up wires the budget into the pick paths and decides per window.
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
2 participants
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.
Groundwork for teaching the CDNA5 scheduler which WMMA windows can afford to give
up an issue slot and which cannot. This lands the measurement; nothing gates on it
yet, and it does not run unless asked for.
No behaviour change. The pre-scan is off by default, and with it on the scheduling
output is byte-identical -- it measures and reports, it does not decide.
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:
capacityCycles is every window cycle after its own issue slot that
blockedScaleMask does not reserve, and capacityValu the subset coIssueWindow
also names.
it -- its deadline. RegionDAG ids are program indices and its edges run
strictly forward, so one reverse sweep settles every node, O(V+E).
on are floating: they still compete for window space at pick time, but they
can always be deferred, so they never force anything.
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.
The knob
dagFeatures.enableWmmaHideBudgetPrescan, default false, reachable only through
stinkytofu-opt --enable-wmma-hide-budget-prescan. There is deliberately no
ModuleOptions mirror, so TensileLite cannot turn it on -- same treatment as
mergeBarrierThreshold. The budget is a pure measurement so far, so a production
compile would pay for a verdict nobody reads.
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 the gated call: the budget is a local in
onInitRegion, reported and dropped, because nothing in the pick paths reads it yet.
The follow-up that calls extraIssueFor() from pickOne() is what promotes it to
per-region state -- until then a member would be state nobody reads, holding
StinkyInstruction* past the region boundary. Its isBlockedCycle() now delegates to
isBlockedWindowCycle() in the analysis header, so the end-anchored-mask convention
has one definition instead of two; that helper is inline because advanceTime(),
computeValuAdvanceCycles() and freeCoIssueSpace() call it once per window cycle.
The queue needs the region DAG to answer "which loads does this WMMA wait on".
The pass already builds that graph 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 from def-use chains, which carry RAW
edges only and are function- rather than region-scoped. It is passed before the
region's policy edges are merged in, so a pre-scan sees real data dependencies
rather than heuristic orderings.
Rather than make that a fifth positional parameter, onInitRegion now takes a
RegionDependencies: the dependency edges the region already has (the DAG, in) and
the ones the queue would like added (requestedConstraints, out, previously its own
parameter). Both are edges of the same graph, which is what makes them one
argument rather than a bag. It also retires a hazard the old signature carried --
three of the five parameters were IRList::iterator, so transposing two of them
compiled silently. One override, one call site; RegionDAG stays forward-declared
in ReadyQueue.hpp to keep the include one-way.
Observability
Test plan
shadow, they fit, region stays quiet.
fit, and window [hipDNN] Enablement #2 is told to overrun by 2. Same IR, only the opcode differs
-- that isolates what the LD_SCALE blocked-cycle model costs when the work has
a deadline.
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.
there, but nothing was blocked, so the LD_SCALE remark must stay silent.
--print-output is byte-identical with the flag on and off.
Notes for reviewers
The "lost every VALU co-issue slot" remark counts only windows that HAD co-issue
slots and lost them to blockedScaleMask. capacityValu == 0 alone does not mean
LD_SCALE took anything: v_wmma_f32_16x16x4_f32 carries coIssueWindow 0x0000, so it
never had a VALU slot and its window still offers 15 issue cycles to SALU and
memory. Pinned by the last function in the FileCheck test.
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.
A follow-up wires the budget into the pick paths and decides per window.
🤖 Generated with Claude Code