Repository navigation
fix(hipblaslt): correct the Stream-K split predicate for uniform summation order - #10896
Merged
jaopaulolc merged 4 commits intoAug 20, 2026
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (35.19%) 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 @@
## users/jolabega/uniform-summation-order-custom-kernels #10896 +/- ##
======================================================================================
Coverage 69.82% 69.82%
======================================================================================
Files 2783 2783
Lines 458625 458625
Branches 67643 67643
======================================================================================
Hits 320221 320221
Misses 115160 115160
Partials 23244 23244
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
This was referenced Aug 18, 2026
Merged
jaopaulolc
added
the
rocm:gemm algos
jaopaulolc
marked this pull request as ready for review
August 18, 2026 19:46
jaopaulolc
force-pushed
the
users/jolabega/streamk-uniform-summation-predicate
branch
from
August 18, 2026 19:51
7f39d17 to
81dcb23
Compare
|
🎉 All checks passed! This PR is ready for review. |
jaopaulolc
force-pushed
the
users/jolabega/streamk-uniform-summation-predicate
branch
from
August 19, 2026 16:48
81dcb23 to
85452f0
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-uniform-summation-predicate
branch
from
August 19, 2026 20:37
85452f0 to
5d73a5f
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-uniform-summation-predicate
branch
from
August 19, 2026 20:47
5d73a5f to
d1b83c1
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-uniform-summation-predicate
branch
from
August 20, 2026 15:15
d1b83c1 to
456fa1a
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-uniform-summation-predicate
branch
from
August 20, 2026 16:06
456fa1a to
86ac929
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-uniform-summation-predicate
branch
from
August 20, 2026 17:37
86ac929 to
4805ed6
Compare
The StreamK=3 and StreamK=5-static kernel-arg packers each wrote out the same skTiles/SKItersPerWG derivation, and CuCount_test.cpp carried a third transcription of it. checkUniformSummationOrder() needs the same numbers to reason about the split the kernel actually performs, which would have made a fourth copy. Introduce streamKStaticSplit() as the single source of truth, plus streamKStaticSplitRowUniform() for the row-uniformity condition over its result, and call the former from all three existing sites. No behaviour change: the helper reproduces the existing arithmetic including the skFullTiles term and the clamp to the tile count.
…he tiles
The uniform-summation-order gate refused any Stream-K launch unless the
resolved grid divided the tile count. That is sufficient but not
necessary. Row uniformity requires only that every output tile be folded
from the same ordered list of chunk lengths, which also holds when the
grid is a multiple of the tile count and every tile is cut into equal
chunks: there is no data-parallel region, so no tile is left whole, and
the cut offsets are identical for all of them.
Test the property directly, over the split streamKStaticSplit() packs:
skTiles == 0 force-DP-only
extraIters == 0 && skItersPerWG != 0
&& (skItersPerWG == itersPerTile every chunk a whole tile
|| (skTiles == tiles
&& itersPerTile % skItersPerWG == 0)) equal cuts, no DP region
The first clause preserves today's acceptance of force-DP-only. The
skItersPerWG != 0 term is load-bearing: the grouped-GEMM callers pass a
default-constructed StreamKSettings, where the modulo would otherwise be
undefined.
Scope it to the packing it describes. StreamK=4 and StreamK=5-dynamic
pack a different ABI and keep the divisibility test, which is redundant
for them because they are separately held to skTiles == 0. The reduction
strategy is now checked before the split arithmetic, since the
arithmetic models the tree packing and the parallel path reinterprets
the same arguments.
This strictly widens acceptance; nothing previously honoured is now
refused. It recovers roughly 10% of the tuned gfx950 Stream-K entries
with no device-side change.
JIRA ID : AIHPBLAS-4252
Four configurations are expressible but were never audited against the
uniform-summation-order guarantee. None of them occurs in the shipped
tuned logic, so no dispatch changes outcome today; each closes a hole a
future kernel or logic file could walk into.
- UseInitialStridesCD, and packed C0/C1 index sets. Both turn off the
coordinate-agnostic store addressing (optSingleColVgpr /
optSharedColVgpr) that keeps an out-of-range lane from influencing an
in-range one on the partials path.
- WaveSplitK, as WorkGroup[2] > 1 with LocalSplitU <= 1. Its
redundant-lane mask covers the D/TD stores, not the workspace partials
store. Gated inside the Stream-K scope on purpose: every shipped
WaveSplitK solution has StreamK=0, and refusing those would buy
nothing. LocalSplitU > 1 and UseSubtileImpl are deliberately untouched.
- MX with DepthU not a multiple of 256. The gfx950 HostPreSwizzle scale
layout packs 32 rows x 8 MX blocks into a 256-byte granule addressed
lane-linearly, and Stream-K shifts the scale SRD by a scalar
StreamKLocalStart*DepthU; a shift inside a granule permutes rows
against K-blocks. Solution validation already forces this for MX fp4
but admits DepthU=128 for MX fp8. Detected from
problemType.mxBlock{A,B}, never from a kernel-name substring: the _MX_
token in several kernel names is F32XdlMathOp=XFloat32, and those f32
kernels ship StreamK=3 with DepthU=16.
Both the throwing gate and the selection-time filter now consult one
shared implementation, so they cannot drift; the filter drops such
solutions quietly instead of failing the launch.
JIRA ID : AIHPBLAS-4255
…tions Three new suites, 27 cases, alongside the 5 pre-existing ones: - RowUniformity/RowUniformityStreamKSplit_pre_checkin, 11 cases. A white-box table over streamKStaticSplit() and streamKStaticSplitRowUniform(), asserting the packed skTiles, SKItersPerWG and extraIters as well as the verdict, so it pins the packing mirror and not only the predicate. Covers both newly admitted cases, both previously admitted ones, four genuinely non-uniform splits, force-DP-only, skFullTiles=0 and a zero grid. - RowUniformityStreamKRejection_pre_checkin, 9 cases. One per rejection against a synthesised solution, plus the over-rejection guards: LocalSplitU > 1 and WaveSplitK without Stream-K must still be admitted, and non-MX DepthU=128 must be admitted, which is the regression test for the f32 kernels whose names carry an _MX_ token. - RowUniformityClauseTwo_pre_checkin, 2 cases. M=4096 N=32 K=10240 bf16 with HIPBLASLT_MATMUL_DESC_SM_COUNT_TARGET=80 and with it unset. Each discovers the resolution per candidate rather than pinning solution indices, asserts the grid is a multiple of the tile count and not equal to it, runs the launch and checks D is bitwise row-uniform, and requires a parallel-reduction resolution to still be refused. Skips with a diagnostic naming the shape and the counts if the library reaches neither regime. The harness gains a per-problem element type and CU budget so the bf16 shape above is reachable. It sets max_workspace_bytes explicitly: the preference defaults to 0 and clamps the request, which would collapse the grid back to the tile count and make the new coverage pass vacuously.
jaopaulolc
force-pushed
the
users/jolabega/streamk-uniform-summation-predicate
branch
from
August 20, 2026 22:53
4805ed6 to
6311476
Compare
jaopaulolc
deleted the
users/jolabega/streamk-uniform-summation-predicate
branch
August 20, 2026 22:57
jaopaulolc
added a commit
that referenced
this pull request
Aug 21, 2026
… output (#10722) JIRA ID : AIHPBLAS-4252, AIHPBLAS-4255 ## Motivation Identical rows of `A` can produce last-bit differences across rows of `D` when tiles accumulate `K` in different orders. This PR adds an opt-in hipBLASLt control for **row uniformity within one run** (not run-to-run determinism). AIHPBLAS-4255 requires StaggerU to be forced off when the mode is on. ## Changes **API** - Descriptor attribute `HIPBLASLT_MATMUL_DESC_UNIFORM_SUMMATION_ORDER_EXT` (105, `int32_t`): `0` off (default), `1` on; other values → `HIPBLAS_STATUS_INVALID_VALUE`. - C++: `GemmPreference::setUniformSummationOrder` / `getUniformSummationOrder`. - Bench: `--uniform_summation_order off|on`. **Guarantee** If the mode is on and `hipblasLtMatmul` succeeds, and every row of `A` is the same vector, then every row of `D` is bitwise identical within that run. **Enforcement** (`ContractionSolution.cpp`) 1. `calculateAutoStaggerU` clamps StaggerU (and related stagger fields) to 0 after env overrides and the prediction-cache write. 2. `checkUniformSummationOrder` fail-closes before dispatch (`UniformSummationOrderError` → `HIPBLAS_STATUS_INVALID_VALUE`). 3. `uniformSummationOrderSupported` drops solutions that are statically unsafe during selection; the launch check remains authoritative. **This PR refuses custom kernels** (`customKernelName` non-empty). Docs, sample, and custom-kernel admission are later stack PRs. Also masks the stagger-mapping bitfield when packing kernel args. Explicit solution-index overrides bypass selection filtering; with the mode on, that path logs a warning instead of hard-rejecting. Bottom of the coherence stack (#10722 → #10765 → #10896 → #10941 → #10942). ## Test plan / result - `row_uniformity_gtest.cpp`: identical `A` rows; bitwise row compares of `D`; `ClampRepaired_*` selects by Tensile metadata. - Aux / `GemmPreference` gtests for default, round-trip, and invalid values. - Suites pass on gfx950 (MI355X). ## Risk Low–medium. Opt-in, off by default, fail-closed. Main shared-path risk is descriptor plumbing (covered by attribute tests).
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.
Folded into #10765. Not part of the live USO stack (#10765 → #10941).
Original work: correct the Stream-K split predicate under uniform summation order (AIHPBLAS-4254). Landed on the #10765 branch; do not restack later PRs on this one.