Repository navigation
fix(tensilelite): keep the packed StaggerU across persistent tiles - #14719
Merged
Merged
Conversation
declareStaggerParms decoded the stride shift and mapping from the packed StaggerU sgpr and then overwrote it with the masked amount (and, for PGR>=3, zeroed it on short loops). StaggerU is only unpacked from the kernel args at entry, but persistent kernels rerun declareStaggerParms for every tile, so a workgroup's second and later tiles staggered with mapping 0 and stride shift 0. Under uniform summation order this gave non-identical D rows for identical A rows on Hybrid StreamK kernels in the dynamic work-queue mode (more tiles than workgroups). Decode the amount into the existing temp instead, and have the PGR=2 closeLoop counter decrement test the masked amount. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…ation order Add a row-uniformity case where an SM count target makes the heuristic pick the Hybrid dynamic sub-mode, so a workgroup runs several whole tiles. Every default pick must give identical D rows for identical A rows, or be refused. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #14719 +/- ##
===========================================
- Coverage 77.34% 77.34% -0.00%
===========================================
Files 2893 2893
Lines 486262 486261 -1
Branches 73502 73502
===========================================
- Hits 376095 376094 -1
Misses 85342 85342
Partials 24825 24825
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
AlexBrownAMD
approved these changes
Oct 30, 2026
ThanHenderson
approved these changes
Oct 30, 2026
iq136boy
approved these changes
Oct 30, 2026
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.
Persistent kernels lose their stagger settings after a workgroup's first tile. With uniform summation order on (the default), this gives different results for identical rows of A.
Cause
The host packs the stagger amount, stride and mapping into one register, which the kernel reads once at start.
declareStaggerParmsdecodes it and then overwrites it with just the amount. Persistent kernels rundeclareStaggerParmsagain for each tile, so from the second tile on, the mapping and stride read as 0.Mapping 0 rotates the start of the K loop by the tile's M index, so rows in different M tiles are summed in a different order. Under uniform summation order this breaks row uniformity. It shows up on Hybrid StreamK kernels in the dynamic sub-mode, which give each workgroup many tiles. Example: f32 1536x1536x2048 with
HIPBLASLT_MATMUL_DESC_SM_COUNT_TARGET=160gives 719,328 mismatched elements.Fix
declareStaggerParmsdecodes the amount into a temporary register and leaves the packed register unchanged.Kernels that process one tile per workgroup behave the same as before.
Testing
declareStaggerParmsnever writes the packed register (PGR 1, 2, 3).HybridDynamicStagger_1536x1536x2048_SmCount160, which fails ondevelopand passes with this fix.Jira ID: AIHPBLAS-5430