Repository navigation
fix(tensilelite): Only defer FUSI next-loop local reads on sub-tile overlap - #12886
Merged
mihnea-chirila merged 7 commits intoOct 8, 2026
Merged
Conversation
…verlap PR #6944 deferred all ForceUnrollSubIter next-loop local reads past the MFMAs for every non-F32X kernel with a single VGPR buffer. The deferral is only needed when one local read spans more tile elements than a sub-tile (lrvwTile > MIWaveTile/numSubTiles): that read then overwrites sub-tile 1 registers the last MFMAs still read. FP8/BF16/F16 FUSI kernels use lrvwTile == 1, so they paid the serialization cost without a hazard (~3-6% slowdown, e.g. FP8 TN 2112x16384x7168). Gate the deferral on the actual overlap, which restores the pre-#6944 schedule for those kernels and keeps the deferral for the SGEMM/CGEMM kernels in bufferLoadStore0.yaml and cgemm.yaml. Also drop the redundant sync1LdsMfmaIndex override added to SIA3. Fixes ROCM-31274
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #12886 +/- ##
===========================================
- Coverage 71.93% 70.92% -1.01%
===========================================
Files 2863 2946 +83
Lines 472710 485747 +13037
Branches 69829 70150 +321
===========================================
+ Hits 340010 344489 +4479
- Misses 109117 117302 +8185
- Partials 23583 23956 +373
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Add a CPU-only codegen test that emits SGEMM ForceUnrollSubIter kernels
(MIWaveTile 4x4, SIA1/SIA3, VectorWidthA/B in {2, 4}) and checks the last
sub-iteration of the main loop:
- no next-loop local read overwrites a register a later MFMA still reads;
- the reads are deferred past the last MFMA exactly when one read spans
more than a sub-tile (VectorWidth 4), and interleaved otherwise.
Drop the `iteration > isBarrier` suppression in the SIA3 path: with
ForceUnrollSubIter numItersPLR is 1, so isBarrier is the last iteration
unless doFullPackCodePrefetch moves it, which requires F32X emulation and
is excluded by the gate. The branch could never be taken.
mihnea-chirila
enabled auto-merge (squash)
October 7, 2026 20:10
Contributor
TheRock Submodule Bump ActivityNewest first
|
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 : ROCM-31274
Motivation
PR #6944 (4068136, "feat(tensilelite): Added Global Read/Write Support") regressed large FP8 TN GEMMs on gfx950 by ~5.8%. The regression also affects BF16 and F16 ForceUnrollSubIter (FUSI) kernels.
The cause is in
_makeSubIterSchedule. #6944 defers all next-loop local reads until after the MFMAs for every non-F32X FUSI kernel with a single VGPR buffer (numVgprBuffer == 1). The deferral is only needed when the reads would overwrite registers that are still in use. For FP8/BF16/F16 kernels they don't, so the scheduler gives up local-read/MFMA overlap for nothing.Technical Details
With FUSI, each operand's X0 buffer holds two sub-tiles. During the last quadrant, the next-loop local reads refill sub-tile 0 while the MFMAs still read sub-tile 1.
bufferLoadStore0.yaml/cgemm.yaml). WithTransposeLDS=0andVectorWidth=4, oneds_read_b128writesValuA_X0_I0+0:3, but a sub-tile is only 2 VGPRs wide. The read clobbers sub-tile 1 before the last MFMAs consume it (WAR). Without the deferral these kernels fail validation (NaNs).lrvwTile == 1(unroll-major LDS orVectorWidth == 1), so each read stays inside one sub-tile. In the FP8 TN kernel, for example, the early reads writeA+0..15/B+0..71while the last MFMAs readA+16..31/B+72..143.This PR gates the deferral on that overlap:
It also removes the
sync1LdsMfmaIndexoverride that #6944 added to the SIA3 non-scheduled-local-write branch. It produces identical code for the #6944 test and can't apply to the regressed kernels (they all useScheduleLocalWrite=1with PGR >= 1).It also drops the
iteration > isBarriersuppression that #6944 added to the SIA3 path. With ForceUnrollSubIter,numItersPLRis 1, soisBarrieris the last iteration unlessdoFullPackCodePrefetchmoves it. That needs F32X emulation, which the gate excludes, so the branch could never be taken.Scope: of the 1,488 FUSI solutions in
asm_fulllogic (all gfx950), about 1,220 are FP8, 54 BF16/F16 and 215 F32X. None of the non-F32X ones meet the overlap condition, so all of them get the pre-#6944 schedule back.Test Plan
All runs on one gfx950 GPU (
HIP_VISIBLE_DEVICES=2).test_config.pywithbufferLoadStore0.yamlandcgemm.yaml(the feat(tensilelite): Added Global Read/Write Support #6944 tests).Tensile/Tests/unit, including the new CPU-only codegen testtest_fusi_next_loop_read_war.py. It emits SGEMM FUSI kernels (MIWaveTile 4x4, SIA1/SIA3, VectorWidthA/B in {2, 4}) and checks the last sub-iteration of the main loop: no next-loop local read overwrites a register that a later MFMA still reads, and the reads are deferred past the last MFMA exactly when one read spans more than a sub-tile. The test fails with feat(tensilelite): Added Global Read/Write Support #6944's original condition (2 cases) and with a data-type-only gate (12 cases).ed5614bf1c6), and the kernel is timed.Test Result
bufferLoadStore0.yamlandcgemm.yamlpass.Performance as relative change in kernel time (negative is faster). Each kernel has a main loop byte-identical to the parent's:
F16 NT sol 197 in detail:
develop: this PR only reorders 18 instructions in the main loop and 18 in the last-iteration (NLL) loop, with no other change. Kernel time drops by 6.4%.Separate, unrelated regression found during validation (not fixed here)
On current
develop, F16 NT sol 197 is still ~3.4% slower than the pre-#6944 parent. Agit bisect(9 steps, the #6944 fix applied where needed) points to b20a30d (#11948, "feat(origami): Update Origami Model to Select Across More Kernel Parameters", AIHPBLAS-1926). The commit is 3.6% slower than its parent524afdea981, with byte-identical kernel code. The change is in the host-side StreamK grid selection:correct_sk_grid_for_partial_tiles(shared/origami/src/origami/streamk.cpp) now replaces the SK grid of 238 with a DP grid of 832.develop+ this PR, F16 NT 8192^3. Kernel time relative to the old SK grid (238) in the same run:TENSILE_STREAMK_FIXED_GRIDThis belongs to the Origami heuristic owners and should be tracked separately. This PR doesn't change it.
Submission Checklist