Repository navigation
fix(tensilelite): generate aux output E in persistent kernels - #12869
Merged
jaopaulolc merged 9 commits intoOct 7, 2026
Merged
Conversation
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❌ Patch coverage is ❌ Your project check has failed because the head coverage (49.31%) 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 #12869 +/- ##
===========================================
- Coverage 71.93% 70.95% -0.97%
===========================================
Files 2863 2946 +83
Lines 472701 485433 +12732
Branches 69827 70078 +251
===========================================
+ Hits 340000 344425 +4425
- Misses 109118 117059 +7941
- Partials 23583 23949 +366
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
A persistent TileProcessingStrategy (StreamK or DataParallel) sets GlobalSplitU to 0. The epilogue already emitted E stores for persistent kernels, but the E address registers, row offset and row advance only checked GlobalSplitU == 1, so generating any persistent kernel with E failed with a TypeError. Set up the E address and row offset for the persistent final store with isPersistent, the same way bias and ScaleAlphaVec already do. In the StreamK tree-reduction fixup, check out the 3 SGPRs the flag wait uses instead of 2; with E, the third one was the low word of the E buffer address and the kernel faulted. Reject E with StoreRemapVectorWidth for persistent kernels, as for GlobalSplitU == 1. Add gfx950 common tests that validate D and E for StreamK with each WorkAssignment, with and without tree reduction, and for DataParallel, each with a gradient case that reads E.
jaopaulolc
force-pushed
the
users/jolabega/fix-tensilelite-streamk-aux-codegen
branch
from
October 2, 2026 17:14
35fa576 to
74e2139
Compare
jaopaulolc
changed the title
jaopaulolc
marked this pull request as ready for review
October 2, 2026 18:19
jaopaulolc
added
the
rocm:gemm algos
AlexBrownAMD
reviewed
Oct 5, 2026
AlexBrownAMD
left a comment
AlexBrownAMD
left a comment
Contributor
There was a problem hiding this comment.
Does this also work with Stream-K + ParallelReduction (post kerenl code path)?
The sk_hgemm_use_e sizes all select tree reduction. Add 512x512x8192: with 16 tiles and 128 iterations per tile, origami selects parallel reduction for StaticGrid and Hybrid, so the GEMM kernel writes split partials to the workspace and the post kernel applies bias and the activation and writes E.
Contributor
Author
|
jaopaulolc
added a commit
that referenced
this pull request
Oct 5, 2026
TensileLite now generates Stream-K kernels that write the auxiliary output E (#12869's fix), so the Origami predictor no longer drops the catalog's Hybrid Stream-K candidates when the request has one. The knowledge test expects an auxiliary output to keep every candidate, and the docs drop the limitation.
AlexBrownAMD
approved these changes
Oct 5, 2026
jaopaulolc
enabled auto-merge (squash)
October 5, 2026 21:28
jaopaulolc
deleted the
users/jolabega/fix-tensilelite-streamk-aux-codegen
branch
October 7, 2026 21:30
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 : AIHPBLAS-4984
Motivation
TensileLite generates the GEMM kernels that hipBLASLt runs. Some also write an auxiliary output, E, holding the values before the activation. Persistent kernels (
TileProcessingStrategyStreamKorDataParallel) loop over output tiles. With E enabled, generating any persistent kernel failed with aTypeError.Technical Details
GlobalSplitUto 0. The epilogue already stored E for them, but the E address and row offset setup inAsmStoreState.py,AsmAddressCalculation.py,ComputeStoreVgprs.pyandLSU.pyrequiredGlobalSplitU1. It now also acceptsisPersistent, as bias does.StreamK.py: the tree-reduction fixup reserved 2 SGPRs but used 3. With E, the third held part of the E buffer address, and the kernel faulted. It now reserves 3.Solution.pyrejects E withStoreRemapVectorWidthfor persistent kernels, as forGlobalSplitU1.Tensile/Tests/common/streamk/gfx950/, FP16 with bias and E plus a gradient case:sk_hgemm_use_e.yamlcovers eachWorkAssignment, with and without tree reduction, anddata_parallel_static_hgemm_use_e.yamlcoversDataParallel.Test Plan
From
projects/hipblaslt/tensilelite, runtox -e py3 -- -m common -k hgemm_use_e.Test Result
On an MI355X (gfx950), both tests fail on
developwith theTypeErrorand pass with this change (176 of 176 cases, D and E validated). The 50 existing persistent, aux output and gradient common tests pass. Unit tests show no new failures. Automated GitHub checks are pending.Submission Checklist
Risk level
Low (2/5): only persistent kernels with E change, and those failed to generate before. Tree-reduction kernels without E generate identical code.