Repository navigation
fix: [hipblaslt] AIHPBLAS-4228 Fix Stream-K launch limits and add unit tests - #12204
Merged
Merged
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. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #12204 +/- ##
===========================================
+ Coverage 70.54% 70.55% +0.01%
===========================================
Files 2812 2812
Lines 463054 463065 +11
Branches 68187 68189 +2
===========================================
+ Hits 326661 326705 +44
+ Misses 112749 112708 -41
- Partials 23644 23652 +8
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
jaopaulolc
approved these changes
Sep 18, 2026
yenong-amd
changed the title
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
tony-davis
added a commit
that referenced
this pull request
Sep 25, 2026
…nch guard #12204, merged after this PR was opened, changed the tree-fixup fallback in getSKGridImpl(). Once the tile count reaches (UINT32_MAX + 1) / threadsPerWorkGroup, it now launches cuCount * occupancy instead of one workgroup per tile. That was the main path StreamKWorkgroupNumberCheck existed to keep selection away from. With it fixed, the predicate rejects every Stream-K solution at or above that tile count, including ones that now run correctly whenever the workspace fits, so it throws away working candidates. Remove the predicate, its serialization and Python emission, its gtests, the characterization test for the isStreamK arm, and the DECISIONS.md entry. This also drops the float-precision test that review found never crossed the boundary. The launch guard in SolutionAdapter::launchKernel stays and is now the whole fix. Paths that still hand Stream-K one workgroup per output tile (the insufficient-workspace fallback in resolveStreamKSettings(), K == 0, the data-parallel debug knob) and the skFixedGrid / skGridMultiplier overrides fail with hipErrorInvalidValue instead of launching a truncated grid. Its comments now say that instead of pointing at the removed predicate. A selection-time check that resolves the grid that will actually launch, workspace fallback included, is the better fix on the selection side. It is left to a follow-up; see the review discussion on #12283. tensilelite-tests on gfx950 (MI355X): 500 ran, 497 passed, 3 skipped (the same three noted in the PR description). ROCM-31290 Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
added a commit
that referenced
this pull request
Sep 28, 2026
…2 bits (#12283) JIRA ID : ROCM-31290 ## Motivation hipBLASLt launches each GEMM kernel with a grid, the number of parallel work items the hardware should run. `KernelInvocation` holds that grid in a 64-bit type, but both launch entry points take 32-bit grid parameters. `hipExtModuleLaunchKernel` declares its `globalWorkSize` arguments as `unsigned int`, and the cluster path's `gridDim` fields are the same width. A grid too large to express was narrowed at the call rather than refused. That produces two outcomes. A work-item count of exactly 2^32 narrows to zero, and HIP rejects the launch with an error. A count above 2^32 narrows to its low 32 bits, a small non-zero number, so the launch succeeds while only a fraction of the intended workgroups run. The output tiles the missing workgroups would have written keep whatever the buffer held before, and `hipblasLtMatmul` returns success. Stream-K kernels reach this most easily. They normally size their grid to the compute-unit count, but several paths hand them one workgroup per output tile instead, and solution selection's `LaunchLimits` check exempts Stream-K. ## Technical Details `SolutionAdapter::launchKernel` now validates the grid and returns `hipErrorInvalidValue` when it does not fit, with a diagnostic naming the offending dimensions. The check runs first, before the code object is looked up, since that is the last point where the 64-bit values are still intact. The two launch paths narrow different quantities, so each is bounded against the one it actually passes. `hipDrvLaunchKernelEx` takes the grid in workgroups and is bounded on `numWorkGroups`. `hipExtModuleLaunchKernel` takes work items and is bounded on `numWorkItems`. A single shared check would wrongly reject a cluster launch whose workgroup count fits even though the work-item product does not. #12204 already fixed the most common way to get there, the tree-fixup fallback in `getSKGridImpl()`, which now launches `cuCount * occupancy` once the tile count reaches the thread-aware cap. The paths that remain include the insufficient-workspace fallback in `resolveStreamKSettings()`, K == 0, the data-parallel debug knob, and the `skFixedGrid` / `skGridMultiplier` overrides. The guard turns all of them from silent truncation into a loud error. Earlier revisions of this PR also added a selection-time predicate, `StreamKWorkgroupNumberCheck`. It was removed in review. After #12204 it rejected Stream-K solutions that now run correctly, and a check that resolves the grid that will actually launch is a better fit for the selection side. That is left to a follow-up. This changes host-side launch validation only. No kernel code or solution selection changes. ## Test Plan - `tensilelite-tests`, `tensilelite/tests/LaunchDimGuard_test.cpp`. Six cases that drive `SolutionAdapter::launchKernel` directly with a synthetic `KernelInvocation`, so no GPU work is dispatched and no kernel needs to exist. They cover the 2^32 boundary, the silent case above it, the largest grid that must still pass, an oversized y dimension, and both directions of the cluster path's workgroups-versus-work-items distinction. - Full `tensilelite-tests` suite on gfx950 to check nothing else moved. Not covered: rocblaslt's rocRoller custom kernels call `hipExtModuleLaunchKernel` directly and never reach `SolutionAdapter`. That path builds its grid from `uint32_t` tile counts, so its failure mode is arithmetic overflow rather than narrowing, which is tracked separately as ROCM-31230. ## Test Result `tensilelite-tests` on an MI355X (gfx950), MessagePack configuration: 500 ran, 497 passed, 3 skipped. The three skips are unrelated and also skip on `develop`: `RangeLibraryTest.SpecificSizes/0`, `MXDataGenFP4Test.ZeroFrequencyWithinBounds/3`, and `DataInitializationReferenceRecompute.Gfx1250ReusesPreswizzledMXAcrossSolutions`. For context on impact, a hardware sweep against the TheRock `10.1.0a20260822` library (which predates #12204) found 39 Stream-K algorithms at M = 19,398,656, N = 256 that returned success with part of the output never written, 22 of them leaving exactly 2^32 elements unwritten. #12204 fixes most of those at that shape. The sweep was not rerun against a current library. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5). The guard is a few integer comparisons at the top of the launch path, and it only changes behavior for grids that could not have run correctly. Solution selection is unchanged. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: tony-davis <7938412+tony-davis@users.noreply.github.com>
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-4228
Motivation
hipBLASLt launches Stream-K kernels (work split along the inner dimension K) through TensileLite’s host solution code. For very large problems the old tree-fixup path set the launch grid to the full tile count. When that count is at or above 2^24 work-items (HIP’s 32-bit work-item limit, also the tree-fixup 24-bit dividend bound), the launch fails. One failing size is M=524288, N=98304, K=128 with a 32×96 tile.
Technical Details
In
tensilelite/src/ContractionSolution.cpp, the tree-fixup fallback now computes the maximum tile count as(UINT32_MAX+1) / threadsPerWorkGroup. If tiles reach that cap, the grid iscuCount * occupancy(occupancy at least 1, still capped at 3 on the Force-DP-only path). Below the tile cap, the previous data-parallel fallback (skGrid = tiles) is unchanged.workGroupSizeis set on Stream-K launch-summary fixtures so the new bound is defined.Test Plan
SKLaunchGridLimitsTest/SKLaunchGridLimitsParamTestintensilelite/tests/CuCount_test.cpp(tile-cap, origami path, below-threshold DP fallback, bench macro-tiles).tensilelite-unit-codecov,preliminary,precheckin(hipblaslt); gfx94X hipBLASLt multi-arch shards.Test Result
tensilelite-unit-codecov, hipblaslt preliminary/precheckin, and gfx94X hipBLASLt shards passed. hipBLASLt ASAN gfx90a quick failed (container/test not classified).
codecov(hipblaslt)Math CI failed. gfx950 device results are author-reported only.Submission Checklist
Risk level
Medium (3/5): changes default Stream-K launch size for huge tile counts. Host unit tests cover the cap; ASAN and hipBLASLt codecov are not green; gfx950 is not in the PR’s GPU CI.