Repository navigation
feat(hipblaslt): Stream-K extra-iters, grid steering, and parallel admission under USO - #10941
Merged
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. |
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-per-tile-extra-iters
branch
from
August 18, 2026 19:51
3df250f to
1b78569
Compare
This was referenced Aug 18, 2026
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #10941 +/- ##
===========================================
- Coverage 69.25% 69.25% -0.00%
===========================================
Files 2820 2820
Lines 467904 467890 -14
Branches 69025 69030 +5
===========================================
- Hits 324030 324017 -13
+ Misses 120320 120300 -20
- Partials 23554 23573 +19
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
jaopaulolc
force-pushed
the
users/jolabega/streamk-per-tile-extra-iters
branch
from
August 19, 2026 16:48
f3b6566 to
281ae13
Compare
jaopaulolc
added a commit
that referenced
this pull request
Aug 19, 2026
…/jolabega/streamk-coherence-grid-steering Rebuild the stack after #10941 was rebased onto the rewritten predicate branch.
jaopaulolc
force-pushed
the
users/jolabega/streamk-per-tile-extra-iters
branch
from
August 19, 2026 20:38
281ae13 to
e80b3c3
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-per-tile-extra-iters
branch
from
August 19, 2026 20:47
e80b3c3 to
4525035
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-per-tile-extra-iters
branch
from
August 20, 2026 15:15
4525035 to
73a80dc
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-per-tile-extra-iters
branch
from
August 20, 2026 16:06
73a80dc to
e70fe99
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-per-tile-extra-iters
branch
from
August 20, 2026 17:37
e70fe99 to
13f051f
Compare
jaopaulolc
force-pushed
the
users/jolabega/streamk-per-tile-extra-iters
branch
from
August 28, 2026 16:06
f4d080b to
de1d347
Compare
caio96
approved these changes
Aug 28, 2026
…ggerU The CompiledInStaggerU clause in uniformSummationOrderLaunchObstacle() refused any solution with `!internalArgsSupport.staggerU && staggerU != 0`. That predicate does not mean what it reads as. SupportCustomStaggerU only controls whether the host writes the packed StaggerU field into internalArg0 bits [31:16]; when it is false the field is simply left 0. It says nothing about where the kernel gets its stagger from. For a generated kernel the answer is: only from that packed field. The solution's StaggerU survives code generation solely as an SGPR-pool sizing input (KernelWriterAssembly.py, the one remaining read of kernel["StaggerU"]); rocisa has no StaggerU concept at all. So the clamp added in calculateAutoStaggerU() does reach these kernels, and refusing them costs uniform summation order a large slice of the tuned solution space for no safety benefit. Evidence, all reproducible on CPU: - Emitting a shipped disqualified solution five ways -- as shipped, with derived parameters recomputed, with StaggerU=0, with SupportCustomStaggerU=true, and with both -- yields byte-identical assembly. Reproduced on a CustomMainLoopSchedule solution and on a non-CMS one. - Disassembling the shipped code objects and running the existing wrap-site and packed-decode detectors over every kernel symbol: 1487 symbols across four libraries, 1045 with an in-loop stagger wrap site, 1045 of those carrying the canonical packed unpack (s_and_b32 against 0x1f00 / 0xe000 / 0xff). Zero kernels stagger without reading the packed argument. Only a handwritten custom kernel can carry a stagger no host clamp can undo, because those are frozen assembly rather than generator output. That is what the clause now tests, via the existing handwrittenCustomKernel() predicate (non-empty customKernel.name with customKernel.generated == false). Handwritten custom kernels are already refused unconditionally at the top of the same function, so the narrowed clause cannot fire today; it is kept, and documented as such, alongside the equally defensive ResolvedStaggerU check, so that relaxing the CustomKernel clause to admit individually vetted kernels cannot silently readmit a compiled-in stagger. The premise is now pinned by a test rather than by review. The custom-kernel StaggerU checks gain a generated-kernel arm: it takes shipped tuning logic containing solutions of exactly the refused shape, trims each to a single-solution logic file, runs the code generator, and holds the emitted assembly to the same clampCannotReach predicate the handwritten kernels are held to -- asserting the kernel does stagger, and that its wrap sites unpack the runtime value. If the generator ever learns to bake a literal StaggerU into a wrap site, this fails and the clause has to widen again.
1 task done
…y tests Narrowing the uniform-summation-order StaggerU refusal to handwritten custom kernels changed which solutions the predicate admits, but the row-uniformity tests still encoded the old rule, so RowUniformityStreamKRejection_pre_checkin.CompiledInStaggerU failed. That case synthesises a generated solution with SupportCustomStaggerU false and a declared StaggerU of 16, and asserted a refusal. The refusal used to come from the clause keyed on internalArgsSupport.staggerU, which the predicate no longer reads: a generated kernel takes StaggerU only from the packed internal argument, so the clamp reaches it and the solution is now admitted. The case drives a mock AMDGPU and a synthesised solution rather than the real device, so it is arch-independent and failed the same way everywhere it ran. Split it into the two halves of the rule it now has to describe. A generated kernel with a compiled-in StaggerU must be admitted, which is the point of the narrowing, and the resolved-StaggerU assertion stays to show the clamp really does settle the launch at zero. A handwritten custom kernel carrying one must still be refused; that case pins the outcome rather than a particular clause, because the CustomKernel clause refuses handwritten kernels before the StaggerU clause is reached and neither is visible through the public predicate. The comment says so rather than leaving the attribution implied. checkClampRepairedSolutions used the same stale rule as its oracle, asserting HIPBLAS_STATUS_INVALID_VALUE for every enumerated candidate declaring SupportCustomStaggerU false. Those candidates now flow into the same bitwise row-uniformity comparison as any other staggering solution, which is strictly stronger: a stagger the clamp genuinely could not reach would surface as a row mismatch instead of being asserted away.
…eamk-per-tile-extra-iters
tony-davis
pushed a commit
that referenced
this pull request
Aug 29, 2026
…ers (#11442) JIRA ID : AIHPBLAS-4531 GitHub ID: #11397 ## Motivation `Tensile/Tests/unit/test_PlaceholderMerge.py::test_logic_yaml_sibling_device_names_consistent` is currently **failing on `develop`**, turning CI red for unrelated hipBLASLt PRs (observed on #10941 among others). This PR fixes the underlying data divergence so the shared CI lane goes green again. The test asserts that library-logic YAMLs sharing a basename within one architecture directory all declare the same `DeviceNames` header. Four gfx1250 `Equality/` basenames declared `- [Device 73f0, Device 0073, Device 75c1]` while their `GridBased/` and `Origami/` siblings declared `[Device 73f0]`. Provenance: the 3-ID list was introduced by `3a08f1b4d94` (#10250, a tuning PR whose description never mentions device IDs, admin-merged with partial CI) and never propagated to the sibling trees — 208 of the 212 files in each affected arch directory use `[Device 73f0]`. The divergence has existed since then but was invisible to CI. `134466bcb4a` (#11274) later re-nested the tree (`asm_full/gfx1250/Equality/` -> `asm_full/gfx1250/gfx1250/Equality/`), which caused the test's `_iter_arch_dirs()` to start descending into that directory and **surfaced** the pre-existing drift. #11274 is not the cause — the divergence is verifiably present at `134466bcb4a^`. ## Technical Details **Commit 1 — normalize the data (8 files).** Four basenames x the `gfx1250` and `gfx1250v0` trees = 8 `Equality/` YAMLs narrowed from `- [Device 73f0, Device 0073, Device 75c1]` to `- [Device 73f0]`, matching their siblings. Narrowing 8 files was chosen over widening the 416 siblings to two additional IDs of unknown provenance (see the open question below). **Why this is functionally safe.** `DeviceNames` entries are PCIe chip IDs consumed to emit a runtime `PciChipId` hardware predicate — but only where `supportsChipIdPredicate()` is true, which is `gfx950` and nothing else: - `tensilelite/Tensile/Common/Architectures.py:180-186` - mirrored in C++ at `tensilelite/include/Tensile/AMDGPUPredicates.hpp:53-56` On gfx1250 the field is parsed and then dropped: no `PciChipId` predicate is emitted and no `_ID` filename suffix is generated. It is inert for kernel selection on this architecture, so no GPU sees a different kernel set as a result of this change. The value does remain live in the Python solution de-duplication key (`Tensile/SolutionStructs/Solution.py:6193-6205`, which is not arch-gated) — which is exactly why keeping it consistent across siblings matters. **Commit 2 — close a test blind spot (separate concern, 1 file).** `_DEVICE_NAMES_RE` in `test_PlaceholderMerge.py` matched only the positional list form (`- [Device ...]`) and silently skipped the mapping form (`DeviceNames: [Device ...]`) used by Origami logic files. The regex now matches both. This is a genuine coverage gap, not cosmetic: one of the four basenames (`*_BBS_BH_Bias_HA_S_SAV_UserArgs.yaml`) diverges only against a mapping-form Origami sibling and was therefore invisible to the old test. Parse coverage over the `asm_full/` tree goes from **1840/1936 to 1936/1936** files. ## Test Plan Regression coverage is the pre-existing shared-CI unit test itself — this is the test that is red on `develop`, so the fix is validated by the same lane that caught it (`test_PlaceholderMerge.py` matches the `test_*` pattern picked up by the Libraries PR Bot unit-test check). Run locally with `rocisa` built via `tlbuild` into `build_tmp` and executed through a `tlrun`-style `PYTHONPATH` (nothing pip-installed): ```bash pytest projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PlaceholderMerge.py pytest projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PlaceholderMerge.py::test_logic_yaml_sibling_device_names_consistent ``` Control run: the same test executed against the pre-fix YAML content, to confirm it actually detects the condition rather than passing vacuously. ## Test Result - `test_logic_yaml_sibling_device_names_consistent` — **PASSED**. - Full `test_PlaceholderMerge.py` — **22 passed**. - Control run against pre-fix data — **failed as expected**, with `AssertionError: Divergent sibling DeviceNames` listing both `*_BBS_*` pairs. Confirms the assertion is live and the widened regex is what surfaces the second pair. - Post-rebase scan of the whole `asm_full/` tree (replicating `_iter_arch_dirs` + `_read_device_names` + the grouping assert): **1936 YAML files walked, 1936 with `DeviceNames` parsed, 0 violations.** Re-run after rebasing onto current `origin/develop` to confirm no newly-landed YAMLs reintroduce divergence. - Branch rebased onto current `origin/develop` (clean, no conflicts) per the `hipblaslt-pr-quality` stale-base rule for validator-vs-validated-data changes. - Automated GitHub checks are pending. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level **Low (1/5).** Eight data files change a field that is provably inert on the only architecture they apply to (`supportsChipIdPredicate()` is gfx950-only), plus one widened regex in a host-side unit test. No kernel source, no build output, no host library code, no public API. The change makes 8 files agree with the 208 already-consistent files in the same directory.
…eamk-per-tile-extra-iters
…eamk-per-tile-extra-iters
…eamk-per-tile-extra-iters Three files conflicted, all from develop-side StreamK work that landed in this window: #10036 (StreamK launch summary diagnostic), #10951 (multi-stream Stream-K Synchronizer fix) and #10925 (fused all-to-all epilogue). Debug.hpp / Debug.cpp: both sides claimed TENSILE_DB bit 0x200000. #10036 is merged and its bit is documented and exercised by StreamKLaunchSummary_test, so it keeps 0x200000; the uniform-summation-order no-solution report moves to 0x400000, the next free bit. Both diagnostics remain independently switchable. ContractionSolution.cpp, getSKGridImpl(): #10036 added out-params recording which post-selection fallback rewrote the grid, and captures the selected grid between the CU/override ladder and the tree-fixup bounds clamp -- the same point the uniform-summation-order F-star snap is inserted at. The snap now runs first and the capture follows it, because the snap is grid selection under uniform summation order rather than a fallback: it has no changedBy attribution, so capturing ahead of it would report selected != final with changedBy = none, the unattributed rewrite the out-params exist to prevent. #10951's flag-region clamp stays last, immediately before the return, so the launch summary's grid ordering is unchanged. ContractionSolution.cpp, solve(): #10036 extracted the reduction / grid / workspace-DP-fallback logic into computeStreamKDecisions(); this branch had extracted the same logic into resolveStreamKSettings(). resolveStreamKSettings() remains the source of truth for what is launched, since it additionally applies streamKReconcileReduction() so requiredWorkspaceSize() and the launch agree on the reduction. computeStreamKDecisions() now applies that same reconciliation at the same point, before the workspace-fit fallback, so the snapshot cannot report a fallback the launch did not take; solve() re-asserts the final reduction alongside the final grid before printing.
ThanHenderson
approved these changes
Sep 1, 2026
ThanHenderson
left a comment
ThanHenderson
left a comment
Contributor
There was a problem hiding this comment.
lgtm after Caio's comments are addressed
The uniform-summation-order grid snap in getSKGridImpl() picks a grid of tiles * F, and the StreamKFlagElements clamp that follows it rewrites any such grid above the bound to StreamKFlagElements. That constant is not in general a multiple of the tile count, so the rewritten grid packs into a ragged static split and checkUniformSummationOrder then refuses the solution outright with StaticSplitNotRowUniform. Kernels that were selectable under uniform summation order stopped being selectable. Fold the bound into the snap's F-star search instead, as one more admissibility condition alongside F | I, MinItersPerCU and the workspace fit. The search then picks the largest uniform grid that already sits inside the flag region and the clamp has nothing left to rewrite. The clamp itself is unchanged and still runs last, so the flag-region invariant is enforced in exactly the same place and on exactly the same launches as before. The new condition carries the same three predicates the clamp uses (non-atomic, not ForceDPOnly, not parallel reduction), so a launch that never reaches the flag protocol is not constrained by the region size. The clamp's fourth predicate, tiles % skGrid != 0, holds for every F >= 2 candidate and is omitted. The all-full grid at F == 1 divides the tiles evenly, has no partial tiles and takes no flag region, so it stays admissible however large the tile count is and remains the floor of the search. A tile count above the bound is therefore not a case that has to fail closed. Enumerating the packer and gate predicates over grids, tile counts, iteration counts, workspace limits and the three launch flags: of the configurations the bound made unselectable, all are selectable again, and no configuration is admitted that was not admitted before the clamp existed. The grid chosen differs only where the previously chosen grid was itself outside the flag region. Cover the three regimes around the bound in row_uniformity_gtest: the bound binds and a smaller uniform grid exists, the bound binds and only all-full remains, and the launch never reaches the flags. Add the packer arithmetic at the clamped, bounded and all-full grids to the split table, and pin that a clamped dynamic-queue grid that no longer divides the tile count is refused at selection rather than launched.
…eamk-per-tile-extra-iters
…olution The generated-kernel arm of the custom-kernel StaggerU check writes a single-solution copy of a shipped tuning-logic file and runs the emitter over it. Its dict-schema branch dropped the ExactLogic key outright, on the grounds that the exact-match table indexes into the solution list the copy replaces. Dropping the key is not equivalent to emptying the table. prepareLibraryLogicDict() reads data["ExactLogic"] unconditionally for Equality, GridBased and Range logic, so a document without it fails to parse with KeyError: 'ExactLogic' before the emitter is ever reached. That branch had not run until now: the gfx950 Equality logic it points at was in the positional schema, which takes the other branch and empties the same table in place at index 7. The files have since been rewritten in the dict schema, and both parametrizations of the check now fail. Empty the table instead of removing it, matching what the positional branch does. The table has to be emptied either way, since its entries reference solution indices the trimmed document no longer has.
…eamk-per-tile-extra-iters
…ideTheSnap The control leg of FlagRegionBoundBindsInsideTheSnap expected the fixed grid of 3200 to survive untouched with uniform summation order off. It does not: the post-selection flag-region clamp is the last write in getSKGrid and fires on this shape, because 3200 exceeds the 2048-entry region and tiles % 3200 != 0 leaves partial tiles that take the flag protocol. The launch is handed 2048. Expect 2048 instead. The leg still pins mode-off apart from mode-on (1600) by exact equality, and it now names the value the clamped split at the end of the same test re-derives: 2048 is not a multiple of the tile count, which is the ragged shape selecting under the bound avoids. No behaviour change; the grid selection logic is unchanged.
1 task done
jaopaulolc
added a commit
that referenced
this pull request
Sep 3, 2026
…ctations (#11635) JIRA ID : AIHPBLAS-4633 ## Motivation The host unit test `StreamKLaunchSummaryTest.Sk3ParallelReductionWorkspaceStarvedFallsBackToTree`, in `projects/hipblaslt/tensilelite/tests/StreamKLaunchSummary_test.cpp`, fails on `develop`. It is a CPU-only test: it builds a `ContractionSolution`, asks it which Stream-K launch it would choose for a given problem, and inspects the resulting decision record. No GPU is involved. A bisect that rebuilt and reran the test at every step identifies `3a558b0d4d4`, "feat(hipblaslt): Stream-K extra-iters, grid steering, and parallel admission under USO" (#10941). Its parent `f4caa56e6ee` passes; `3a558b0d4d4` and every descendant fail. The test's expectations, not the product code, are what became wrong. #10941 changed which stage of the decision pipeline gets credit for the chosen grid, and the test asserted the older attribution. The Jenkins job that runs these host unit tests, `mci/rocm-libraries/tensilelite-unit-codecov(hipblaslt)`, is not a required check for merging, so a failure in it does not block a pull request. That is how the regression reached `develop` unnoticed. ## Technical Details Stream-K spreads a matrix multiply across a fixed grid of workgroups rather than one workgroup per output tile, so several workgroups can contribute partial sums to the same tile and those partials must be combined afterwards. Two combining strategies exist here. *Parallel* reduction writes partial tiles to a scratch workspace and reduces them in a second pass, so it requires that workspace. *Tree* reduction combines in place and requires none. The k-split factor `F = grid / tiles` is how many workgroups share one tile; parallel reduction is only meaningful when `F` is at least 2. #10941 introduced `streamKReconcileReduction(reduction, skGrid, tiles)` in `ContractionSolution.cpp`. It demotes parallel to tree whenever `F` drops below 2, and the same commit called it from all three sites that must agree on the reduction: - `requiredWorkspaceSize()`, which callers use to size the workspace allocation; - `resolveStreamKSettings()`, which is what `solve()` launches; - `computeStreamKDecisions()`, the diagnostic snapshot this test inspects. In the snapshot the call sits immediately before the guard that decides whether a partials workspace is needed. The test drives a 256x4096x4096 problem with a 128x128 macro tile on a 256-CU gfx950 device, so `tiles` is 64, and sets the workspace size to 0. Origami's `grid_k_split_aware` selection (`shared/origami/src/origami/streamk.cpp`) already rejects any grid whose partial tiles would not fit the workspace, collapsing the grid from 256 down to `tiles` = 64. That clamp is long-standing and #10941 did not touch it. The reconcile therefore sees `F = 64 / 64 = 1` and demotes the reduction to tree *before* the guard runs. The guard reserves only when the reduction is parallel or `tiles % grid` is non-zero, and neither now holds, so its body never executes. `idealWorkspaceBytes` stays 0 and `workspaceDPFallbackFired` stays false, contradicting the old assertions. **Why the fix belongs in the test.** The launch is unchanged by #10941: before and after, this configuration launches tree reduction with `grid == tiles == 64` and no workspace. Only the attribution moved, and it moved toward the truth, because grid *selection* produced `grid = 64`, not the workspace-DP fallback. Reverting or reordering the reconcile would reintroduce a genuine divergence between the workspace query and the launch, since `requiredWorkspaceSize()` sizes through `requiredWorkspaceSizeGsu(grid / tiles)`, which short-circuits to 0 at a split factor of 1, while `computeStreamKDecisions()` sizes through `partialTileSize(grid)`, which does not. With the reconcile ahead of both, that triple reaches neither formula and the two cannot disagree about whether a workspace is needed. All production logic is therefore left untouched; the only non-test edit is a comment. The change has three parts: 1. **Renamed and corrected the failing test** to `Sk3ParallelWorkspaceStarvedGridCollapsesToTilesAndReconcilesToTree`, matching actual behaviour: no workspace-DP fallback, not `dpOnly`, `idealWorkspaceBytes` of 0, and `changedBy = none` / `source = none` in the printed summary. Assertions that were already correct are kept, including that `getSKReduction()` still starts from parallel, the tile and grid equalities, and that both workspace sizings return 0. 2. **Added `Sk3ParallelFixedGridWorkspaceDpFallbackFires`** to restore the coverage the rename would otherwise delete. It sets `skFixedGrid = 128` (that is, `2 * tiles`) with a zero workspace, so `F = 2` survives the reconcile and the parallel-side workspace-DP fallback genuinely fires: `idealWorkspaceBytes` is 8388608 bytes (`partialTileSize(128)`), the grid moves from a selected 128 to a final 64, and the summary reports `changedBy = workspaceDP` and `source = workspaceDP(runtime)`. A fixed grid is the only remaining way to reach that branch: on the dynamic-grid path a parallel launch has `skGrid = F * tiles > tiles`, so `tiles % skGrid` is always non-zero, and origami applies the same workspace predicate during selection first, meaning parallel only survives selection when the workspace already fits. #10941 silently cost that branch its coverage; this test restores it. 3. **Corrected a stale comment** on `StreamKDecisions::requiredWorkspaceBytes` in `ContractionSolution.hpp`, which documented a query-versus-snapshot divergence (using this same 256x4096x4096 example) that the reconcile made unreachable, plus the same staleness in a neighbouring test comment. No changelog entry: `projects/hipblaslt/CHANGELOG.md` is a release-notes file keyed to ROCm versions and records user-facing API and feature changes. Recent test-only commits under `projects/hipblaslt/tensilelite/tests/` do not touch it, and this change alters no public behaviour. ## Test Plan Build the tensilelite C++ host-library gtest binary (CMake `TENSILELITE_BUILD_TESTING=ON`, target `tensilelite-tests`) and run: 1. The affected suite on its own: `./tensilelite-tests --gtest_filter=StreamKLaunchSummaryTest.*`. 2. The full binary, `./tensilelite-tests`, to confirm no collateral damage. Both tests read `TENSILE_STREAMK_DATA_PARALLEL`, which forces data-parallel mode and would make the scenarios vacuous, so that variable (and `TENSILE_DB`) must be unset; each test asserts this up front. No GPU or device library is required. ## Test Result Both runs pass. - `StreamKLaunchSummaryTest`: 25 of 25 pass, including the renamed test and the new `Sk3ParallelFixedGridWorkspaceDpFallbackFires`. - Full `tensilelite-tests` binary: 447 tests across 81 suites, 445 passed, 0 failed, 2 skipped. Both skips are pre-existing and unrelated to this change (`RangeLibraryTest/RangeLibraryTest.SpecificSizes/0` and `FP4ZeroFrequency/MXDataGenFP4Test.ZeroFrequencyWithinBounds/3`). `clang-format` reports no changes needed on the modified line ranges of both files. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level None. The change touches no production logic. It edits one test file and one comment in `ContractionSolution.hpp`; no declaration, signature, or executable statement in the header is modified, so no runtime behaviour and no generated kernel can change. The worst case is a test that is itself wrong, which the full-suite run above bounds.
Contributor
TheRock Submodule Bump ActivityNewest first
|
This was referenced Sep 10, 2026
jaopaulolc
added a commit
that referenced
this pull request
Sep 11, 2026
…ites #10941 did not introduce the launch-side forcing in computeStreamKDecisions(); it copied it into the extracted resolveStreamKSettings(). Removing it is out of scope for a revert of #10941's unguarded changes and breaks USO-off parity with f4caa56. Only the requiredWorkspaceSize() mirror, which is net-new in #10941, stays removed.
jaopaulolc
added a commit
that referenced
this pull request
Sep 11, 2026
streamKReconcileReduction()'s threshold predates uniform summation order: dc47acc turned the same triple into a hard throw. #10941 converted that throw into a graceful demotion, which is a fix in its own right, so gating it on the mode re-exposed the throw for non-USO SK3 / SK5-static launches that reach parallel with F < 2 and sufficient workspace. Drop the guard parameter and update the three call sites. The throw in solve() stays as defense in depth.
jaopaulolc
added a commit
that referenced
this pull request
Sep 11, 2026
…ernel disjunct Two of my earlier conclusions were wrong, both caught by row_uniformity_gtest: * Re-adding streamKPerQueueStrideBytes * streamKBakedQueueCount to the dynamic-queue workspace turns an exactly-sized allocation into a spurious DP fallback. Nothing addresses those bytes; #10941 removed them deliberately and TileSchedulingOnModeOffPinsBaselineLaunchGrid pins the removal with the mode off. Removed from all three sizing sites again. * getSKReduction()'s handwrittenCustomKernel() disjunct is not redundant: a solution can carry customKernel.name with sizeMapping.customKernelName empty, which ForceTreeUnderUniformSummationOrderCustomKernel exercises in both modes. Restored unconditionally.
jaopaulolc
added a commit
that referenced
this pull request
Sep 11, 2026
…Settings solve() already resolves streamK5EffectiveDynamic() for its dynamic-queue predicate, and resolveStreamKSettings() then recomputed the identical value for the same (problem, hardware) -- a third call per SK5 solve() where the pre-#10941 baseline made two. The helper is not free: with a CU-count hint and streamKTileSchedulingMode == 0 it runs origami::streamk::select_hybrid_mode. Pass the value out of band instead, via an optional `bool const*` hint parameter, mirroring the sk5EffectiveDynamic pointer getSKGridImpl already takes for exactly this purpose. Null (every other caller) recomputes as before, so behaviour is unchanged.
jaopaulolc
added a commit
that referenced
this pull request
Sep 11, 2026
Comment-only pass over the USO guard changes, plus a leaner unit test. - Rewrite test_streamk_uso_gate.py at 294 lines (was 423). Coverage of the 44 original assertions is preserved; the file was mutation-tested against five mutations of the source under test, all still caught. - Collapse four over-long SGPR-budget comments in KernelWriter.py and StreamK.py to point at the StreamK.py header rather than restating the bit-29 protocol at each site. - Drop the stale "(+ the per-XCD work-queue region on the dynamic path)" clause from the partials-workspace comment in ContractionSolution.cpp: the body sizes idealWorkspace as partialTileSize(grid) with no work-queue term, so the parenthetical contradicted the rule this PR fixes. - Apply a prose pass over 73 comment blocks in 12 files, including factual corrections that name #10941 as the origin of each guarded behavior change. The five selection-site comments are deliberately not merged: each names a different pre-#10941 predicate conjunction. - Use the symbol-free SPDX copyright form in the two added test files, per AGENTS.md.
jaopaulolc
added a commit
that referenced
this pull request
Sep 11, 2026
Fix eight comment inaccuracies found by an accuracy audit: invert the baseline attribution on the solve() dynamic-queue guard (the pre-#10941 baseline did call computeStreamKDecisions() there), drop the stale single-source claim on computeStreamKDecisions(), correct the magicNumberAlg2 shift bound (p <= 64, so the shift field is bits 0..5), say pre-#10941 rather than pre-USO/before-the-feature where USO itself predates #10941, give each selectionPredicate arm its own rationale, qualify the requiredWorkspaceSize() reconcile note with skFixedGrid == 0 instead of USO, and refresh two StreamK.py line citations. Also correct two stale TENSILE_DB bit references for the no-solution tally (0x400000, not 0x200000). Comments only; no executable change.
AlexBrownAMD
pushed a commit
that referenced
this pull request
Sep 15, 2026
…flag (#11943) JIRA ID : AIHPBLAS-4253 ## Motivation Uniform Summation Order (USO, `HIPBLASLT_MATMUL_DESC_UNIFORM_SUMMATION_ORDER_EXT`) is host-side runtime state that defaults **OFF** at every layer. The USO PR stack was supposed to be inert when the mode is off. It is not: it changes kernel **selection**, **grid sizing** and other **launch parameters** with USO disabled. Two independent defects: **1. Device side — the kernel inferred the mode from grid arithmetic.** The Stream-K K-split mapping diverged on `skTiles != 0 && skGrid % skTiles == 0` rather than on an explicit signal. With USO off, any launch that happened to satisfy that geometric coincidence still took the new "per-tile extra iterations" mapping instead of the historical global "first-E" mapping. Reproduced on hardware (see below). The unconditional mapping also costs a **fixed, additive per-call** amount that does not scale with problem size; it tracks workgroup waves. Measured drift-free (the guard ported directly onto #10941, so unrelated `develop` commits are out of the comparison; 25 shapes x 8 interleaved reps, in-run A/A control +0.138%): **#10941 alone costs +1.660%** geomean against the pre-#10941 baseline, 95% CI [+0.985%, +2.390%], and **the guard cuts that to +1.060%**, CI [+0.688%, +1.469%] — roughly **a third** of the cost removed. Selecting on bit 29 restores the pre-existing mapping with the flag off. **2. Host side — a queue-region workspace addend split the workspace query from the launch.** `resolveStreamKSettings()` — the single source of truth for the launched grid — sized its workspace-sufficiency threshold on the dynamic-queue path as ``` partialTileSize(grid) + streamKPerQueueStrideBytes(hw) * streamKBakedQueueCount(hw) ``` PR #10941 added a **third copy** of that addend, in the `resolveStreamKSettings()` it introduced. PR #11515 (`99878a06a3c`) then deleted it from `requiredWorkspaceSize()` and `computeStreamKDecisions()` but left the `resolveStreamKSettings()` copy — the path `solve()` actually launches from — still adding it (still present on `develop` at `ContractionSolution.cpp:4868`). The workspace-size query and the actual launch therefore report different byte counts, and a caller that allocated exactly what the library asked for got the workspace-DP fallback anyway (`reduction = tree; grid = tiles`), inflating the launch grid. The addend reserves bytes nothing ever addresses. The per-XCD work-queue counters it nominally sizes do not live in the workspace at all: they sit in a separate 8 MiB allocation addressed off `AddressFlags` (`StreamK.py:540-576`), present identically in the pre-#10941 baseline. The workspace region is written solely through `SrdWS`, with slot stride `MacroTile0 * MacroTile1 * bpeCinternal` over slots `[0, skGrid)` (`StreamK.py:1858-1891`), so `partialTileSize(skGrid)` is exactly the highest address touched. USO is never consulted anywhere on that path, which is why it fires with USO off. Measured on 207 Stream-K problems with `--streamk_tile_scheduling on` and USO off: **62 of 207** inflate, e.g. 768 → 3456, 512 → 4782, 768 → 6000, ranging from 1.3x up to **43.9x** (512 → 22475). ## Technical Details ### Device side (`tensilelite/Tensile/Components/StreamK.py`, `KernelWriter.py`) Bit 29 of the `MagicShiftItersPerTile` transport word now carries the mode explicitly. `emitUsoBranchToGlobal` emits `s_bitcmp1_b32 sN, 29` followed by `s_cbranch_scc0` to the pre-existing global-mapping label, at the three K-split divergence sites. The bit is tested **in place** — never extracted, never cleared, never held: the `PGL1/SKFDPO0/PGR1/SIA0` gfx1250 configuration sits at 105 of 106 SGPRs, so holding the mode in an SGPR was not affordable. Net codegen delta: **-3 instructions** on the non-VGPR path. `SupportStreamKPerTileExtraIters` is re-documented and re-derived as a pure **capability** flag ("this kernel's asm carries both mappings and honors bit 29"), computed in `Solution.py` from `StreamK in (3, 5) and not customKernel`, overriding whatever the YAML said — only the generator knows what it just emitted. Custom (hand-written) kernels can no longer claim it. ### Host side (`tensilelite/src/ContractionSolution.cpp`, `.hpp`) - Two packer sites OR in `0x20000000u`, each gated on `internalArgsSupport.perTileExtraIters && problem.getParams().uniformSummationOrder()`. The wire-format asserts widen from `0xC0000000u`/`0x40000000u` to `0xE0000000u`/`0x60000000u`. - `streamKStaticSplitRowUniform()` and `streamKWorkgroupIterRange()` gain `perTileCapable` + `uniformSummationOrder`; the clause becomes `perTileActive = perTileCapable && uniformSummationOrder`. - The F-star grid snap in `getSKGridImpl()` is gated the same way. - **The queue-region workspace addend is removed on all paths.** This closes the query/launch split that #11515 created, on the smaller side, and conforms to the approach taken in PR #11951. With USO off the allocation is **1024 bytes smaller** on 235 of 889 traced shapes under a `--sm_count_target 224` CU hint (all of them `SK5 -> dynamic(SK4)` with tree reduction) and unchanged on the rest; at the default CU count no shape resolves to that sub-mode, so the change is invisible without the hint. Re-adding the addend was tried and measured harmful: it turns an exactly-sized allocation into a spurious data-parallel fallback. - **Deliberate deviation from strict baseline parity:** with USO off, a Stream-K parallel reduction whose split factor `F = skGrid / tiles` is below 2 is **demoted to tree reduction** rather than throwing. The pre-#10941 baseline threw `std::runtime_error` on exactly this condition (`f4caa56e6ee`, lines 3559-3562, introduced by `dc47acc3628` in January 2026 to fix a customer-visible CPX-mode failure), surfacing to callers as `HIPBLAS_STATUS_INTERNAL_ERROR`. #10941 converted that into a graceful demotion, which is a correctness improvement, so this branch keeps it **unconditional**. Guarding it would reintroduce a query/launch workspace disagreement at USO off. - `solve()` no longer calls `computeStreamKDecisions()` on the hot path; the diagnostic snapshot is built lazily only under `TENSILE_DB` bit `0x200000`. ### Selection path (`SolutionLibrary.hpp`, `SingleSolutionLibrary.hpp`, `PredictionLibrary.hpp`, `MLPClassificationLibrary.hpp`, `GranularitySelectionLibrary.hpp`, `tensile_host.cpp`) The stack had widened the candidate filter at six call sites — adding `taskPredicate` conjuncts, `problem.setGroupedGemm(true)`, and the `softwarePredicate()` wrapper. None of those are inert. Each site now restores the exact pre-USO conjunction when `uniformSummationOrder()` is false and takes the widened path only when it is true. `streamKDynamicQueueSupported()` is deliberately left **unconditional** — it excludes dynamic-queue Stream-K kernels on non-power-of-two-XCD devices regardless of USO, which is a correctness filter, not a USO policy. One unconditional fix rides along and is called out in the code: `MLPClassificationLibrary.hpp` had `while(it != it + numToSort)`, which re-evaluates `it + numToSort` against the current `it` and so degenerates to `numToSort != 0` with no end bound — it walked off the vector as soon as any ranked kernel was rejected. Now bounded by a snapshotted `it_end`. This is a correctness fix, not part of the guard. ### Regression test `projects/hipblaslt/clients/tests/src/row_uniformity_gtest.cpp` — `RowUniformityGridSteering_pre_checkin.TileSchedulingOnModeOffPinsBaselineLaunchGrid` **There is a usability trap here and the test exists to pin it.** In the *default* regime, the unguarded build and the fixed build agree on **1216 of 1216** problems, so anyone validating this guard with a plain `hipblaslt-bench` sweep concludes it is a no-op. The host-side regression is only reachable with `--streamk_tile_scheduling on`, because `streamK5EffectiveDynamic()` keys off `problem.getParams().streamKTileSchedulingMode()` — a **separate** public API attribute (`HIPBLASLT_MATMUL_DESC_STREAMK_TILE_SCHEDULING_EXT`) independent of USO and also defaulting off. With both off, every launch stays on the static sub-mode and the touched code is never reached. The test builds a synthetic SK5 solution on a synthetic 256-CU / 8-XCD gfx950 device, sets tile scheduling ON, asserts USO is OFF (anti-vacuity), allocates **exactly** `requiredWorkspaceSize()` bytes, calls `ContractionSolution::solve()`, and pins the launch parameters to **literals**, not to self-consistency: - `tiles == 1248`, selected grid `== 256`, `tiles % grid != 0`, so the guard is reachable and the test is not vacuous; `numWorkGroups == (256, 1, 1)`, `numWorkItems.x == 65536`, `numWorkGroups.x != tiles` - `requiredWorkspaceSize()` bytes + 1 changes nothing (the threshold is a plain inequality); − 1 **must** still fall back to `grid == tiles`, so the test also fails if the guard is removed outright rather than merely mis-sized `solve()` is the instrument on purpose: `resolveStreamKSettings()` is private, and `computeStreamKDecisions()` does **not** model the launch threshold, so a test written against it would pass on the regressed build. ## Test Plan Three prebuilt arms, all `TENSILELITE_LOGIC_FILTER=gfx950/*/*`, 271 `.co` / 273 `.dat.zlib`: | Arm | Worktree | Commit | |---|---|---| | BASE | `uso-baseline` | `1326d629f07` (pre-USO) | | CUR | `uso-revert` | `a111ebfad16` (unguarded HEAD) | | FIX | this branch | `a111ebfad16` + patch | 1. **G1** — with USO OFF, FIX must equal BASE. 2. **G2** — with USO ON, FIX must equal CUR. 3. **Selection identity** — with USO off, every launch-decision field must match the pre-#10941 baseline (`f4caa56e6ee`) across the traced shape population. 4. **Codegen diff** — object-level opcode census, FIX vs CUR at the same commit. 5. **Falsification of the new regression test** — it must FAIL on the unguarded behaviour. 6. Test suites: `tensilelite-tests`, StreamK gtest subset, Python codegen StreamK/USO subset, `hipblaslt-test '*RowUniformity*'`. All binaries `ldd`-verified to resolve `libhipblaslt.so.1` into their own build tree, never `/opt/rocm`. ## Test Result ### Selection identity with USO off — 889 / 889 The 24 traced decision fields are winning kernel name, Stream-K mode, grid, tiles, skTiles, skSplit, totalItems, partials, reduction strategy, `isDynamic`, allocated / required / ideal / given workspace, and every fallback flag. With USO off, all 24 match the pre-#10941 baseline (`f4caa56e6ee`) on **889 / 889** shapes at the default CU count. Under a `--sm_count_target 224` CU hint, **654 / 889** match exactly and the remaining 235 differ in workspace bytes alone — always 1024 bytes smaller than baseline, all of them `SK5 -> dynamic(SK4)` with tree reduction, the intended effect of removing the queue-region addend. The winning kernel, grid, tiles and reduction strategy are identical on all 889 shapes under both configurations. This is the acceptance criterion for the guard, and it is met. ### G1 — USO off: FIX == BASE - **0 of 120** problems differ in `norm_error`, the client's scalar error vs its CPU reference — not an output-buffer comparison, so an equal scalar is strong evidence, not proof, of equal output bits. A bit-exact check is not expressible today: `--dump_matrix` writes 7 significant digits where f32 needs 9, masking the low-mantissa differences at issue. - In the dynamic tile-scheduling regime, **207/207** launch-parameter-identical on the fix's own native library, including all 62 inflating problems. End-to-end corroboration on four representative problems, all three arms, both against the baseline Tensile library (isolating the host) and against each arm's own library: | Problem | BASE grid | CUR grid | FIX grid | |---|---|---|---| | M=48 N=884736 K=16 NN | 768 | **3456** | 768 | | M=256 N=458984 K=96 NN | 512 | **4782** | 512 | | M=160 N=107005 K=128 TN | 768 | **1254** | 768 | | M=9984 N=2048 K=128 NN | 256 | **507** | 256 | BASE and FIX agree in *every field* of the parsed launch record (solution name, every kernel argument, occupancy, WGM, stagger, grid) on all four. All three arms report the same Stream-K mode decision (`reason=api-on -> effective=dynamic`), so the divergence is purely grid sizing, not mode selection. ### G2 — USO on: FIX == CUR Across **1216** problems, 1183 differ and the **only** field that ever differs is `MagicShiftItersPerTile`, by exactly `+2^29` (536870912), on all 1183 without exception. The 33 identical problems are exactly the `SK=0` ones, which carry no Stream-K transport word. No solution name, family, grid, skGrid, skTiles, GSU, WGM, StaggerU, occupancy, workspace or Stream-K mode difference anywhere. ### Hardware demonstration of the device-side leak M=1 N=9216 K=512, batch 1, `f32_r` in/out, `c_xf32_r` compute, NN: | Build | USO | norm_error | |---|---|---| | BASE | off | 1.8864e-06 | | CUR (unguarded) | off | **1.88837e-06** ← the USO-ON number, leaking | | FIX | off | 1.8864e-06 | | FIX | on | 1.88837e-06 | The host launch record is **identical in every parsed field** across all three arms, which is what proves the leak is device-side. ### Per-call cost | Comparison | geomean delta | 95% CI | |---|---|---| | this branch vs its own parent (`a111ebfad16`) | **-0.571%** | [-0.946%, -0.212%] | | guard vs pre-#10941 baseline, drift-free | +1.060% | [+0.688%, +1.469%] | | #10941 alone vs pre-#10941 baseline, drift-free | +1.660% | [+0.985%, +2.390%] | Against its own parent the guard **recovers** time and adds none. The residual against the baseline is a fixed per-launch cost attributable to ~820 bytes of additional kernel code that #10941 emits into every Stream-K kernel regardless of the guard. Removing that would require a library reship and is out of scope for this PR. ### Codegen 194 code objects differ between CUR and FIX (same commit, so attributable to the patch alone). `s_bitcmp1_b32` and `s_cbranch_scc0` change in **194/194** and are always exactly paired. `v_cndmask_b32_e32` churns in 130 objects in **both** directions (−19 to +7, mode ±1) — the assembler re-scheduling around the inserted branch, not new computation. **No other opcode differs anywhere.** ### Regression test falsification The test passes on the fixed build and fails when only the deleted `if(dynamicQueuePath) idealWorkspace += perQueueStride * bakedQueueCount;` hunk is restored, which isolates the defect to that single line. Under the unguarded behaviour the launch grid reverts to `tiles` (1248) instead of the CU-bounded 256, and `numWorkItems.x` to 319488 instead of 65536. ### Suites | Suite | Result | |---|---| | `tensilelite-tests` | **447 passed**, 0 failed, 2 skipped | | StreamK gtest subset | **57/57** | | Python codegen StreamK/USO subset | **420 passed**, 0 failed | | `hipblaslt-test '*RowUniformity*'` | **63/63**, 0 failures (62 pre-existing + the new one) | ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level **Medium.** For lower risk: the change is a restoration, not new behaviour, and every result above is a parity result — baseline-identical with USO off, HEAD-identical but for bit 29 with USO on, with codegen a net instruction reduction. For higher risk: it touches six selection call sites and the single source of truth for the Stream-K launch grid — code that decides which kernel runs and how it is shaped. Grouped-GEMM and the low-workspace regime are unvalidated, and the low-workspace regime is exactly where the workspace-threshold change lives. The `F < 2` demotion is a known, deliberate deviation from the baseline's throwing behaviour.
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-4253, AIHPBLAS-4254
Depends on #10765.
Summary
Under
HIPBLASLT_MATMUL_DESC_UNIFORM_SUMMATION_ORDER_EXT, every static Stream-K tile must use the same K-iteration fold. This PR maps leftovers inside each tile and steers grid/reduction onto a launch the #10765 gate will accept.Per-tile extra iters. When
grid = tiles * F(F ≥ 2) andFdoes not divideitersPerTile, leftovers used to go to the first workgroups in global order, so tiles got different split-length signatures. #10765 refusesextraIters != 0in that regime. WhenskGrid % skTiles == 0, this PR redistributes leftovers inside each tile so every tile has the same number of splits and matchingPartialIdxlengths. Device mapping is inStreamK.pyfor SK3 / SK5-static and matching fixup peers; it reuses caller SGPRs (no extra named checkOuts) and keeps the current tile index across the SK3 DP-to-SK extra-iters path.Capability bit.
InternalArgsSupport::perTileExtraItersdefaults to false for existing and custom binaries. Newly generated StreamK 3/5 set it true.streamKStaticSplitRowUniformthen admits all-partial launches with packedextraIters != 0when the bit is set andskGrid % tiles == 0. Host helperstreamKWorkgroupIterRangemirrors the device formula for tests.Grid and reduction steering. Parallel reduction is admitted when the launch is non-atomic, uses static two-tile packing,
grid % tiles == 0, andF ≥ 2(streamKParallelReductionRowUniform). Any other parallel pick under the mode is forced to tree. When both tree and parallel are legal, keep admitted parallel. After the usual CU grid chain and before the tree magic-division guard, snap the grid never-upward onto an admissible uniform size; mixedg0 < tilessnaps to all-full (grid = tiles).requiredWorkspaceSizeuses the same SK4 / SK5-dynamic tree forcing assolve().Identifiers use uniform summation order, not “coherence”. Handwritten custom kernels in steering fixtures are tagged with
customKernel.name/generatedso they hithandwrittenCustomKernel().No public API change. Mixed
tiles > gridDP/SK band repacking is out of scope.Test plan
./clients/hipblaslt-test --gtest_filter='*RowUniformity*'*PerTileExtraIters*)GridDividesTilesto all-full, parallel snap skipsF|ICuCount_testSK3 parallel / USO steering fixturesRisk
Medium. Touches Stream-K codegen and the fail-closed gate. Mitigated by defaulting
perTileExtraItersto false for existing binaries, shared split/parallel helpers, never-upward snap, and gtest coverage. Mode remains opt-in.