Repository navigation
test(tensilelite): fix stale Stream-K parallel workspace-starved expectations - #11635
Merged
Merged
Conversation
…ctations StreamKLaunchSummaryTest.Sk3ParallelReductionWorkspaceStarvedFallsBackToTree has been failing on develop since 3a558b0 ("feat(hipblaslt): Stream-K extra-iters, grid steering, and parallel admission under USO", #10941). A bisect that rebuilt and reran the test at every step confirms it: the parent commit f4caa56 passes, and 3a558b0 and every descendant fail. #10941 added the helper streamKReconcileReduction(reduction, skGrid, tiles) in ContractionSolution.cpp. It demotes parallel reduction to tree whenever the k-split factor skGrid / tiles falls below 2, and the same commit called it from all three sites that have to agree on the reduction: requiredWorkspaceSize(), resolveStreamKSettings() (which is what solve() launches), and computeStreamKDecisions() (the diagnostic snapshot the test inspects). In the snapshot the call sits immediately before the partials workspace guard. 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 already collapses the grid from 256 down to tiles = 64 for a zero workspace; that clamp is long-standing and #10941 did not touch it. The reconcile therefore sees a split factor of 1 and demotes the reduction to tree before the guard runs, and the guard only reserves when the reduction is parallel or tiles % grid is non-zero, so its body never executes. That is why idealWorkspaceBytes stays 0 and workspaceDPFallbackFired stays false. The launch itself is unchanged. Before and after #10941 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 rather than the workspace-DP fallback is what produced grid = 64. The test was asserting the old attribution, so the expectation was stale and the product logic is deliberately left untouched. Reverting or reordering the reconcile would reintroduce a real divergence between the workspace query and the launch: 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. Changes: - Rename the test to Sk3ParallelWorkspaceStarvedGridCollapsesToTilesAndReconcilesToTree and correct its expectations to the behaviour the code has: 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. - Add Sk3ParallelFixedGridWorkspaceDpFallbackFires to restore the coverage the rename would otherwise drop. It sets skFixedGrid to 128, which is 2 * tiles, with a zero workspace, so the split factor of 2 survives the reconcile and the parallel-side workspace-DP fallback genuinely fires: idealWorkspaceBytes is 8388608, the grid moves from a selected 128 to a final 64, and the summary reports "changedBy = workspaceDP". A fixed grid is the only way left to reach that branch, because 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. - Correct a stale comment on StreamKDecisions::requiredWorkspaceBytes, which described a query-versus-snapshot divergence that the reconcile made unreachable, and the same staleness in a neighbouring test comment. Verified with the tensilelite host unit tests: StreamKLaunchSummaryTest passes 25 of 25, and the full binary reports 445 passed and 2 pre-existing skips across 447 tests in 81 suites, with no failures.
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
yenong-amd
approved these changes
Sep 2, 2026
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (38.61%) 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 #11635 +/- ##
===========================================
- Coverage 63.78% 63.75% -0.03%
===========================================
Files 2825 2825
Lines 469111 469773 +662
Branches 69265 69386 +121
===========================================
+ Hits 299192 299457 +265
- Misses 148872 149256 +384
- Partials 21047 21060 +13
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
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-4633
Motivation
The host unit test
StreamKLaunchSummaryTest.Sk3ParallelReductionWorkspaceStarvedFallsBackToTree, inprojects/hipblaslt/tensilelite/tests/StreamKLaunchSummary_test.cpp, fails ondevelop. It is a CPU-only test: it builds aContractionSolution, 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 parentf4caa56e6eepasses;3a558b0d4d4and 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 reacheddevelopunnoticed.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 / tilesis how many workgroups share one tile; parallel reduction is only meaningful whenFis at least 2.#10941 introduced
streamKReconcileReduction(reduction, skGrid, tiles)inContractionSolution.cpp. It demotes parallel to tree wheneverFdrops 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 whatsolve()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
tilesis 64, and sets the workspace size to 0. Origami'sgrid_k_split_awareselection (shared/origami/src/origami/streamk.cpp) already rejects any grid whose partial tiles would not fit the workspace, collapsing the grid from 256 down totiles= 64. That clamp is long-standing and #10941 did not touch it.The reconcile therefore sees
F = 64 / 64 = 1and demotes the reduction to tree before the guard runs. The guard reserves only when the reduction is parallel ortiles % gridis non-zero, and neither now holds, so its body never executes.idealWorkspaceBytesstays 0 andworkspaceDPFallbackFiredstays 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 == 64and no workspace. Only the attribution moved, and it moved toward the truth, because grid selection producedgrid = 64, not the workspace-DP fallback. Reverting or reordering the reconcile would reintroduce a genuine divergence between the workspace query and the launch, sincerequiredWorkspaceSize()sizes throughrequiredWorkspaceSizeGsu(grid / tiles), which short-circuits to 0 at a split factor of 1, whilecomputeStreamKDecisions()sizes throughpartialTileSize(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:
Renamed and corrected the failing test to
Sk3ParallelWorkspaceStarvedGridCollapsesToTilesAndReconcilesToTree, matching actual behaviour: no workspace-DP fallback, notdpOnly,idealWorkspaceBytesof 0, andchangedBy = none/source = nonein the printed summary. Assertions that were already correct are kept, including thatgetSKReduction()still starts from parallel, the tile and grid equalities, and that both workspace sizings return 0.Added
Sk3ParallelFixedGridWorkspaceDpFallbackFiresto restore the coverage the rename would otherwise delete. It setsskFixedGrid = 128(that is,2 * tiles) with a zero workspace, soF = 2survives the reconcile and the parallel-side workspace-DP fallback genuinely fires:idealWorkspaceBytesis 8388608 bytes (partialTileSize(128)), the grid moves from a selected 128 to a final 64, and the summary reportschangedBy = workspaceDPandsource = workspaceDP(runtime). A fixed grid is the only remaining way to reach that branch: on the dynamic-grid path a parallel launch hasskGrid = F * tiles > tiles, sotiles % skGridis always non-zero, and origami applies the same workspace predicate during selection first, meaning parallel only survives selection when the workspace already fits. feat(hipblaslt): Stream-K extra-iters, grid steering, and parallel admission under USO #10941 silently cost that branch its coverage; this test restores it.Corrected a stale comment on
StreamKDecisions::requiredWorkspaceBytesinContractionSolution.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.mdis a release-notes file keyed to ROCm versions and records user-facing API and feature changes. Recent test-only commits underprojects/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, targettensilelite-tests) and run:./tensilelite-tests --gtest_filter=StreamKLaunchSummaryTest.*../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 (andTENSILE_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 newSk3ParallelFixedGridWorkspaceDpFallbackFires.tensilelite-testsbinary: 447 tests across 81 suites, 445 passed, 0 failed, 2 skipped. Both skips are pre-existing and unrelated to this change (RangeLibraryTest/RangeLibraryTest.SpecificSizes/0andFP4ZeroFrequency/MXDataGenFP4Test.ZeroFrequencyWithinBounds/3).clang-formatreports no changes needed on the modified line ranges of both files.Submission Checklist
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.