Skip to content

feat(tensilelite): reorder wmma - #12767

Open
chengchingwen wants to merge 1 commit into
ROCm:developfrom
chengchingwen:user/chengchingwen/loopmodel
Open

chengchingwen wants to merge 1 commit into
ROCm:developfrom
chengchingwen:user/chengchingwen/loopmodel

Conversation

@chengchingwen

Copy link
Copy Markdown
Contributor

Motivation

JIRA ID: AIHPBLAS-4229

WMMA GEMM inner loop is scheduled by hand-written logic, which fixes one traversal of the K/M/N axes and one placement for every local read, LDS copy and wait. This adds an opt-in path UseLoopModel that allows to change the order of loop axes.

UseLoopModel=False remains the default and the normal path.

Technical Details

New opt-in parameter. UseLoopModel (default False) routes the inner-loop body through
the new path.

New kernel parameters, all defaulting to the existing behaviour:

parameter meaning
WmmaInnerOrder 1..6 the order the tile axes nest in (KMN, KNM, MNK, MKN, NMK, NKM)
WmmaOuterOrder -1..3 which split axis is hoisted outermost; -1 derives it from the live splits
TDMSplitA / TDMSplitB 0..2 per-operand TDM descriptor split: 0 none, 1 the tile axis, 2 DepthU

Two new folder Tensile/LoopModel holds the Solution->LoopIR plus its checks.
Tensile/Lowering lowers LoopIR to GIR and renders it.

StinkyTofu barrier and wait placement. — GirFrameContractImportPass, GirFencePlacementPass, GirWaitCntInsertionPass and the GirFrameCounterFlow analysis — which read a frame contract emitted beside the GIR and decide fences and s_wait_* counts from the program itself.

Test Plan

Tensile/Tests/common/gemm/gfx12/ulm_gfx1250.yaml

Test Result

Submission Checklist

@therock-pr-bot

therock-pr-bot Bot commented Sep 30, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

  • python -m pip install pre-commit
  • pre-commit install
  • pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

@chengchingwen
chengchingwen force-pushed the user/chengchingwen/loopmodel branch from 84eb697 to 728668b Compare September 30, 2026 04:10
@therock-pr-bot

therock-pr-bot Bot commented Sep 30, 2026 •

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@chengchingwen
chengchingwen force-pushed the user/chengchingwen/loopmodel branch 2 times, most recently from c255630 to 32d0f98 Compare September 30, 2026 15:25
@KKyang
KKyang requested a review from nakajee October 1, 2026 02:21

@KKyang KKyang left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for this. It's a big change, so here is what I'd want before it goes further. I read the changes to existing StinkyTofu and TensileLite code closely and only skimmed the new LoopModel/, Lowering/ and GIR pass bodies.

1. Unit tests

  • The only tests added are GirFrameAnalysisTest.cpp, one filecheck (gir_frame_waitcnt.stir) and the sweep yaml ulm_gfx1250.yaml. The roughly 15k lines of new Python (Tensile/LoopModel, Tensile/Lowering, Components/LoopModel) have no unit tests, and neither do GirFencePlacementPass or GirFrameCounterFlow.
  • The SCC rule rewrite has no test either (see the inline comment).
  • Please also show that existing kernels are unchanged with UseLoopModel=False. The wave-separated TDM descriptor code was replaced for all kernels (initTDMDescriptorWaveSeparated and friends became the ...Group helpers driven by the new tdmFusedGroups and tdmGroupWaveRanges). By my reading the predicate and the default grouping match the old behaviour, but a before/after assembly diff on the existing gfx1250 TDM yamls (tdm_gfx1250, decouple_pgr_tdm_fuse_gfx1250, tdmpluslds_tdm_gfx1250, the streamk TDM ones) would settle it.

2. Description
The "Test Result" section is empty, and the Technical Details are incomplete. Please add:

  • The new StinkyTofu pieces: OrderTokenData and NoWaitCntData, the module options EnableGirFramePipeline, EnableGirFrameWaitCntInsertion, DisableWaitCntRemoval and DisableTensorcntInsertion, and the dag features useMemoryTokenOrdering and preserveTensorLoadOrder.
  • The change to StinkyBuildImplicitDependencyPass and RegionDAG (order-token walls), the WaitDataflow gating, and PB.applyExtensionPoint(EndOfPipeline, ...) in the backend.
  • The SCC rule rewrite and the TDM group refactor of existing code.
  • The rocisa instruction.hpp/.cpp changes.
  • What was run and the results. Also, the StinkyTofu paragraph starts with a stray "—" and "Two new folder" needs fixing.

3. Please split this up if you can
At 137 files and +25k lines in one commit it's hard to review, and a regression on the default path wouldn't be attributable. The StinkyTofu behaviour changes in particular should land separately and first:

  • SCC rule rewrite (StinkyDAGSchedulerPass.cpp).
  • Order-token ordering (OrderTokenData, RegionDAG, StinkyBuildImplicitDependencyPass) and preserveTensorLoadOrder.
  • The WaitDataflow / WaitAwareScheduleRepairPass / backend wiring.

After that, the TDM grouping refactor in TensileLite, and then the loop model itself (LoopModel/, Lowering/, GIR passes, docs).

4. Aborts
See the inline comments. report_fatal_error (and STINKY_UNREACHABLE) end in std::abort(), so input-dependent failures should return errors that Tensile can turn into a solution rejection.

What I checked and found fine: the new StinkyTofu code is inert for existing kernels. OrderTokenData is new and only the adaptor sets it, GirFrameAnalysis returns early when no contract is loaded, and the legacy TDMSplit emitter that was removed was already rejected on the parent (Solution.py). I'd just like that to be demonstrated by tests and a before/after diff rather than my reading.


Generated by Claude Code

// because its reader is waiting on that barrier. So pin to a fixpoint, then lock.
std::vector<std::vector<char>> reach;
std::vector<std::vector<char>> pinned(live.size(), std::vector<char>(barriers.size(), 0));
for (bool changed = true; changed;) {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This rewrites the existing cluster-barrier SCC rule (applyClusterBarrierSccRule is split into pinChainAfterBarriers plus a pin-to-fixpoint loop, and lock decisions now use the post-pin reach). It isn't related to the loop model, and the PR description doesn't mention it.

I traced the single-chain case and it reduces to the old logic. What changes is the multi-chain interplay, which is exactly the situation the new comment describes (a lock and a barrier blocking each other). It changes scheduling for existing kernels whenever ClusterBarrier is on, not only under UseLoopModel.

Two asks:

  • Please move this to its own PR so a scheduling change on the default path is attributable to one commit.
  • Please add a unit test. The PR changes no scheduler test, so nothing exercises the fixpoint or the lock/barrier deadlock it fixes. A DAGSchedulerPassTest case with two chains where pinning one makes the other's reader depend on a barrier would do, and the existing ClusterBarrierSccRule_* tests should be run against it.

Generated by Claude Code

@chengchingwen

Copy link
Copy Markdown
Contributor Author

OrderTokenData and NoWaitCntData related will be removed.

@KKyang

KKyang commented Oct 1, 2026 •

Copy link
Copy Markdown
Contributor

Summarizing the concerns raised across the inline comments, since they add up to one overall point about this PR's design:

  1. A third IR layer (GIR), alongside the existing LoopIR and rocisa Logical IR. This part I don't object to on its own -- whether StinkyTofu should take on a GIR layer at all is a bigger architectural call than one PR review, and should be discussed with @nakajee before this goes further, not decided here.
  2. GirFrameContractImportPass/GirFrameAnalysis -- a second hazard-discovery system, bypassing the existing MemTokenData/StinkyBuildImplicitDependencyPass mechanism instead of extending it.
  3. GirFencePlacementPass -- duplicates InsertClusterBarrierPass (same kWorkgroupBarrierId = -1 barrier, already split/placed by the existing pass in this same file).
  4. GirWaitCntInsertionPass -- duplicates StinkyWaitCntInsertionPass (see the inline comment on this pass's own header: it already strips and re-derives every wait from a clean slate).
  5. Three carve-outs added to existing, general passes (disableWaitCntRemoval, useMemoryTokenOrdering = !EnableGirFramePipeline, RemoveDscntPass skipped) only to keep 3/4 from being clobbered or altered by the originals they duplicate.
  6. The loop-axis-reordering scheme itself (WmmaInnerOrder/WmmaOuterOrder) duplicates what LogicalScheduler/SchedulerConfig already does for the Subtile kernel path (Tensile/Components/Subtile) -- a second answer to the same scheduling-freedom problem, for a different kernel architecture in the same codebase.

Net: 2-6 are the same pattern repeated -- new machinery standing in for something that already exists in this codebase, rather than extending it. I don't think this should merge as one 25k-line commit while that's unresolved; see the earlier review for the split-PR ask. Whether a GIR layer is the right direction at all for StinkyTofu (1) is worth resolving with @nakajee first, since that answer changes what 2-5 should even look like.


Generated by Claude Code

@chengchingwen

chengchingwen commented Oct 1, 2026 •

Copy link
Copy Markdown
Contributor Author
  1. GirFencePlacementPass -- duplicates InsertClusterBarrierPass

The FencePlacementPass add workgroup barrier, not cluster barrier. The pass inserts barrier after instruction scheduling.

  1. GirWaitCntInsertionPass -- duplicates StinkyWaitCntInsertionPass

GirWaitCntInsertionPass use exist StinkyWaitCntInsertionPass for dscnt/vgpr use-def, only the tensorcnt and tdm WAR dscnt is inserted by this gir pass based frame graph/hazard

  1. The loop-axis-reordering scheme itself (WmmaInnerOrder/WmmaOuterOrder) duplicates what LogicalScheduler/SchedulerConfig already does for the Subtile kernel path

The loop axis reordering would cooperate with the PLR/CLR/TDMSplitAB/TDMFuse behavior that is not expressible by what Subtile path currently has, and enable the full ST scheduling.

@KKyang

KKyang commented Oct 1, 2026 •

Copy link
Copy Markdown
Contributor
  1. GirFencePlacementPass -- duplicates InsertClusterBarrierPass

The FencePlacementPass add workgroup barrier, not cluster barrier. The pass inserts barrier after instruction scheduling.

  1. GirWaitCntInsertionPass -- duplicates StinkyWaitCntInsertionPass

GirWaitCntInsertionPass use exist StinkyWaitCntInsertionPass for dscnt/vgpr use-def, only the tensorcnt and tdm WAR dscnt is inserted by this gir pass based frame graph/hazard

  1. The loop-axis-reordering scheme itself (WmmaInnerOrder/WmmaOuterOrder) duplicates what LogicalScheduler/SchedulerConfig already does for the Subtile kernel path

The loop axis reordering would cooperate with the PLR/CLR/TDMSplitAB/TDMFuse behavior that is not expressible by what Subtile path currently has, and enable the full ST scheduling.

  1. It should be renamed to something else instead of GIR.
  2. Don't create a special pass even if you reuse.
  3. They added in the internal repo, although the concept is not present, but we should extend instead of a brand-new path unless there are other issues.

@chengchingwen

Copy link
Copy Markdown
Contributor Author
  1. GIR is short for GemmIR, which is where the frame contract entered that allowing the barrier/waitcnt insertion based on the GEMM semanatics.
  2. The pass is different with existing pipeline. The existing tensorcnt insertion is based on the barrier placement. The GIR path remove barrier and insert barrier after instruction scheduling. barrier placement and waitcnt placement would be co-decided.
  3. It is a separate layer. They can still convert to GIR and reuse the gir passes.

@chengchingwen
chengchingwen force-pushed the user/chengchingwen/loopmodel branch 3 times, most recently from 07303b0 to aeeedfb Compare October 3, 2026 16:48
@codecov

codecov Bot commented Oct 3, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 83.48474% with 1418 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...slt/tensilelite/tensilelite/LoopModel/traversal.py 85.71% 83 Missing and 39 partials ⚠️
...ilelite/tensilelite/LoopModel/loopir/functional.py 71.66% 83 Missing and 38 partials ⚠️
...laslt/tensilelite/tensilelite/LoopModel/decoder.py 90.26% 62 Missing and 34 partials ⚠️
...t/tensilelite/tensilelite/LoopModel/loopir/expr.py 81.21% 66 Missing and 18 partials ⚠️
.../hipblaslt/tensilelite/tensilelite/KernelWriter.py 63.00% 55 Missing and 19 partials ⚠️
...silelite/tensilelite/LoopModel/loopir/semantics.py 81.43% 38 Missing and 27 partials ⚠️
...lite/tensilelite/Components/LoopModel/TdmRegion.py 63.03% 61 Missing ⚠️
...aslt/tensilelite/tensilelite/LoopModel/schedule.py 76.99% 39 Missing and 16 partials ⚠️
...lelite/tensilelite/Components/LoopModel/Program.py 23.94% 52 Missing and 2 partials ⚠️
...lite/tensilelite/Components/LoopModel/Registers.py 75.00% 32 Missing and 14 partials ⚠️
... and 41 more
Additional details and impacted files
@@             Coverage Diff             @@
##           develop   #12767      +/-   ##
===========================================
- Coverage    77.77%   77.72%   -0.05%     
===========================================
  Files         2905     2974      +69     
  Lines       491723   503457   +11734     
  Branches     74578    76933    +2355     
===========================================
+ Hits        382425   391292    +8867     
- Misses       84295    86525    +2230     
- Partials     25003    25640     +637     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 54.02% <100.00%> (+0.11%) ⬆️
TensileLite-Unit 83.82% <83.37%> (-1.26%) ⬇️
hipBLAS 90.62% <ø> (ø) Carriedforward from 41d662e
hipBLASLt 70.07% <ø> (+<0.01%) ⬆️
hipCUB 82.68% <ø> (ø) Carriedforward from 41d662e
hipDNN 87.25% <ø> (ø) Carriedforward from 41d662e
hipFFT 42.65% <ø> (ø) Carriedforward from 41d662e
hipRAND 76.12% <ø> (ø) Carriedforward from 41d662e
hipSOLVER 69.67% <ø> (ø) Carriedforward from 41d662e
hipSPARSE 87.03% <ø> (ø) Carriedforward from 41d662e
rocBLAS 69.20% <ø> (ø) Carriedforward from 41d662e
rocFFT 46.49% <ø> (ø) Carriedforward from 41d662e
rocRAND 57.32% <ø> (ø) Carriedforward from 41d662e
rocSOLVER 76.65% <ø> (ø) Carriedforward from 41d662e
rocSPARSE 83.60% <ø> (ø) Carriedforward from 41d662e
rocThrust 93.15% <ø> (ø) Carriedforward from 41d662e

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines Coverage Δ
...lite/rocisa/rocisa/src/instruction/instruction.cpp 86.93% <100.00%> (+7.76%) ⬆️
...tensilelite/tensilelite/Common/GlobalParameters.py 96.18% <100.00%> (+0.20%) ⬆️
...nsilelite/tensilelite/Common/RequiredParameters.py 100.00% <ø> (ø)
.../tensilelite/tensilelite/Common/ValidParameters.py 100.00% <ø> (ø)
.../tensilelite/tensilelite/Components/ClusterLoad.py 99.16% <100.00%> (ø)
...ipblaslt/tensilelite/tensilelite/Components/GSU.py 74.53% <ø> (ø)
...lelite/tensilelite/Components/LraTileAssignment.py 92.35% <100.00%> (+0.12%) ⬆️
...silelite/tensilelite/Components/TensorDataMover.py 91.45% <100.00%> (-0.32%) ⬇️
...lt/tensilelite/tensilelite/KernelWriterAssembly.py 79.34% <ø> (-1.42%) ⬇️
...nsilelite/tensilelite/LoopModel/loopir/__init__.py 100.00% <100.00%> (ø)
... and 81 more

... and 1 file with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

@chengchingwen
chengchingwen force-pushed the user/chengchingwen/loopmodel branch 2 times, most recently from 2feee35 to 30fa5a0 Compare October 12, 2026 07:17
@chengchingwen
chengchingwen force-pushed the user/chengchingwen/loopmodel branch 6 times, most recently from 3aab103 to 726cfd4 Compare October 20, 2026 09:26
@chengchingwen
chengchingwen force-pushed the user/chengchingwen/loopmodel branch 10 times, most recently from 777670a to b5a4383 Compare October 27, 2026 18:09
///
/// Both halves are pinned. `owner.anchor` names the fence by its signal, and the token edges all
/// land on the wait, so pinning the anchor alone gives the halves different windows to drift in.
static void pinGirFenceWindows(RegionDAG& dag, const GirFrameHazardAnalysis::Result& hazards,

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@KKyang I don't think this part is necessary after #14636

passFeatureConfig.dagFeatures.girFramePipeline =
moduleOptions.EnableGirFramePipeline;
passFeatureConfig.dagFeatures.clusterBarrierTensorLoadDelayCycles =
moduleOptions.EnableGirFramePipeline && moduleOptions.ClusterBarrier ? 200 : 0;

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why use this magic number to seperate the tensor_load and barrier? Could you find a rule when the tensor load can be issued?

return true;
};

// Merge a nearby later fence into the first fence when at most two configured WMMA queues

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We may discuss how to merge the token for all kinds of kernels. This looks only for the large MT type. cc @aazz44ss

@chengchingwen
chengchingwen force-pushed the user/chengchingwen/loopmodel branch from 0d27126 to fa1fc22 Compare November 6, 2026 08:09

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants