Repository navigation
feat(tensilelite): Enable PrefetchGL2 with GSU - #11851
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. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #11851 +/- ##
===========================================
- Coverage 70.55% 70.55% -0.01%
===========================================
Files 2810 2810
Lines 462487 462532 +45
Branches 68090 68094 +4
===========================================
+ Hits 326286 326294 +8
- Misses 112661 112696 +35
- Partials 23540 23542 +2
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
(cherry picked from commit c581f26)
Add GlobalSplitU 2 to the gfx1250 TDM-multicast PrefetchGL2 fork so the contiguous-chunk (GSUC) layout gets benchmarked, and drop the stale comments that said PrefetchGL2 rejects GSU > 1. Co-authored-by: Cursor <cursoragent@cursor.com>
…ject 6-bit float The prefetch stream derives its K chunk from the GSU and GSUSumIdx SGPRs at run time, so it follows a host-supplied GSU as well as the tuned one; drop the SupportUserGSU force-off. Separately, a 256-byte prefetch granule is not a whole number of 6-bit elements, so reject 6-bit float rather than emit the wrong address math. Widen the multicast benchmark to the [1,1] and [2,2] cluster shapes. Co-authored-by: Cursor <cursoragent@cursor.com>
The default launch does not give a cluster's workgroups a shared GSU group. The WG-id decode sets WorkGroup1 = cluster_y*nwg_y + wg_y, so a cluster owns consecutive raw y values, and GSUOn.graWorkGroup then takes GSUSumIdx = wg1 % GSU with GlobalSplitUWorkGroupMappingRoundRobin off by default. Cluster peers therefore land on different K chunks and their divided WorkGroup1 collapses onto one tile, while the harness feeds the group index from grid z and keeps them distinct -- so pairing the two asserted a fan-out production does not build. Move the GSU configs to ClusterDim [1,1] and record the reasoning. MT 768 on gsu3_contiguous_rem carries the non-POT ncc and multi-load geometry the cluster used to supply; the fan-out itself stays covered by the cluster configs. Co-authored-by: Cursor <cursoragent@cursor.com>
With GlobalSplitUWorkGroupMappingRoundRobin set, GSUOn.graWorkGroup splits the raw y index as GSUSumIdx = wg1 / NumWorkGroups1, making the group the slow axis, so a cluster's workgroups share one group and span distinct tiles. That makes GSU + ClusterDim a launch the prefetch has to handle, which the default mapping does not produce. Add four configs for it. They are not a second copy of the chunk math: the per-iteration increment is built from the cluster-folded leading extent, so the chunk offset startIter*inc is itself scaled by the cluster and the folded block has to translate as one piece. Covers both chunk layouts, the widest [4,4] fan-out with an uneven split, non-POT cluster extents, and MX scales where each tensor scales the shared start iteration by its own increment. Rework the mapping note to describe both arrangements rather than only the default one. Co-authored-by: Cursor <cursoragent@cursor.com>
boringmorning
force-pushed
the
users/boringmorning/gl2-gsu
branch
from
September 11, 2026 04:36
e3e03ba to
39bff5e
Compare
jichangjichang
approved these changes
Sep 21, 2026
This was referenced Sep 21, 2026
tony-davis
added a commit
that referenced
this pull request
Sep 21, 2026
Reverts #11851 Fixes #12384 ## Motivation `TensileLite coverage / Linux` has been red on `develop` since #11851 landed, and `Component CI Summary` is a required check, so every PR that triggers the tensilelite lane is blocked. The gate reports one regressed file: ```text coverage_ratchet: FAIL - 1 file(s) regressed: baseline current delta file 94.98% 84.65% -10.33 Tensile/Components/GL2Prefetch.py ``` #11851 did ship tests (`test_gl2_prefetch_offset.py` and `test_prefetchgl2_streamk_guard.py`), but the new GSU prefetch paths outran them, so the file measures 10.33 points below the floor #11939 set for it on Sep 11. #12384 has the full bisect: last green on `3082a49c`, red on #11851's own merge commit and on every tensilelite-touching push since. ## Technical Details Straight revert of #11851, no hand edits. It removes the GSU-with-PrefetchGL2 enablement from `GL2Prefetch.py`, `KernelWriter.py`, `KernelWriterAssembly.py`, `Solution.py` and `ValidParameters.py`, plus the two unit tests and the `tdm_multicast_gfx1250.yaml` test config that came with it. `GL2Prefetch.py` at this PR's head is byte-identical to the file on `3082a49c`, the last commit where the coverage lane was green (blob `4256c495c` in both). So this clears the gate without touching `coverage-baseline.json`, and the 94.98% floor stands as-is. This also clears the red on bump PR #12364 and anything else currently tripping the lane. ## Test Plan - `TensileLite coverage / Linux` on this PR. It re-measures `GL2Prefetch.py` from source against the committed baseline, so a green run is the direct check that the floor is satisfied again. - The rest of the tensilelite and hipblaslt CI as normal for a codegen change. This removes a kernel-generation path, so the functional suites are the check that nothing depended on it. ## Test Result Pending this PR's own CI run. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level Low (2). This restores the exact `GL2Prefetch.py` content that was on `develop` for the two weeks before #11851, so the path it removes has only been live since Sep 21. The cost is feature loss rather than instability: anyone depending on PrefetchGL2 with GSU loses it until a re-land. No baseline or CI configuration changes.
Contributor
TheRock Submodule Bump ActivityNewest first
|
evedovelli
pushed a commit
that referenced
this pull request
Sep 23, 2026
Reverts #11851 Fixes #12384 ## Motivation `TensileLite coverage / Linux` has been red on `develop` since #11851 landed, and `Component CI Summary` is a required check, so every PR that triggers the tensilelite lane is blocked. The gate reports one regressed file: ```text coverage_ratchet: FAIL - 1 file(s) regressed: baseline current delta file 94.98% 84.65% -10.33 Tensile/Components/GL2Prefetch.py ``` #11851 did ship tests (`test_gl2_prefetch_offset.py` and `test_prefetchgl2_streamk_guard.py`), but the new GSU prefetch paths outran them, so the file measures 10.33 points below the floor #11939 set for it on Sep 11. #12384 has the full bisect: last green on `3082a49c`, red on #11851's own merge commit and on every tensilelite-touching push since. ## Technical Details Straight revert of #11851, no hand edits. It removes the GSU-with-PrefetchGL2 enablement from `GL2Prefetch.py`, `KernelWriter.py`, `KernelWriterAssembly.py`, `Solution.py` and `ValidParameters.py`, plus the two unit tests and the `tdm_multicast_gfx1250.yaml` test config that came with it. `GL2Prefetch.py` at this PR's head is byte-identical to the file on `3082a49c`, the last commit where the coverage lane was green (blob `4256c495c` in both). So this clears the gate without touching `coverage-baseline.json`, and the 94.98% floor stands as-is. This also clears the red on bump PR #12364 and anything else currently tripping the lane. ## Test Plan - `TensileLite coverage / Linux` on this PR. It re-measures `GL2Prefetch.py` from source against the committed baseline, so a green run is the direct check that the floor is satisfied again. - The rest of the tensilelite and hipblaslt CI as normal for a codegen change. This removes a kernel-generation path, so the functional suites are the check that nothing depended on it. ## Test Result Pending this PR's own CI run. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level Low (2). This restores the exact `GL2Prefetch.py` content that was on `develop` for the two weeks before #11851, so the path it removes has only been live since Sep 21. The cost is feature loss rather than instability: anyone depending on PrefetchGL2 with GSU loses it until a re-land. No baseline or CI configuration changes.
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-4771
Motivation
PrefetchGL2(the gfx1250global_prefetch_b8GL2 warm-up) was only usable on a narrow slice of the solution space:GlobalSplitU > 1was rejected outright.This PR lifts the restriction so the prefetch stream follows the same K order the real load stream walks. When cluster is enabled,GSUWGMRRis forced on since the cluster owns ClusterDim[1] consecutive raw y values, and the default mapping makes the group the fast axis of that y (GSUSumIdx = wg1 % GSU), which spreads peers across chunks and folds their divided WorkGroup1 onto fewer tiles. Round-robin makes it the slow axis instead (GSUSumIdx = wg1 / NumWorkGroups1) so a cluster shares one group and spans distinct tiles.Technical Details
GlobalSplitU (
Components/GL2Prefetch.py,SolutionStructs/Solution.py)calculateGSUIterOffsetdecodes bit 15 (GSUC) of the packedGSUkernel argument once perkernel and produces the unroll iteration this workgroup's chunk starts at:
GSUSumIdxforinterleaved chunks, or the accumulated iterations of the lower groups (via the existing
calculateLoopNumIterOffsetGsu) for contiguous ones.own per-iteration increment in
applyGSUChunk, which then widensGL2PrefetchInc{tc}to thechunk stride (
GSUinterleaved,1contiguous). The chunk offset is applied before the PGRpre-skip, since it must consume the unscaled increment that the pre-skip then reuses at chunk
stride.
Solution.pyis dropped. ForceGSUWGMRRon when cluster is enabled.Test Plan
Test Result
2 streamK yamls failed, not related to this PR. All other tests passed
[31mFAILED�[0m Tensile/Tests/common/test_config.py::�[1mtest_config[Tensile/Tests/common/streamk/gfx1250/core/sk_hybrid_quick.yaml]�[0m - subprocess.CalledProcessError: Command '['/home/marhuang/rocm-libraries/pro...
�[31mFAILED�[0m Tensile/Tests/common/test_config.py::�[1mtest_config[Tensile/Tests/common/streamk/gfx1250/core/sk_dynamic_sgemm_quick.yaml]�[0m - subprocess.CalledProcessError: Command '['/home/marhuang/rocm-libraries/pro...
Submission Checklist