Repository navigation
Revert "feat(tensilelite): Enable PrefetchGL2 with GSU" - #12385
Merged
Merged
Conversation
This reverts commit 9b702dc.
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
One or more issues must be addressed before approval.
Get a fresh assessment by requesting another Copilot review.
Review effort: Lite
Findings: 2
Open (3)
What changed in this PR
Reverts PrefetchGL2’s GSU support in TensileLite, restoring rejection of unsupported GSU configurations and removing related tests and code paths.
Changes:
- Restores PrefetchGL2 GSU validation and removes GSU address-generation logic.
- Removes GSU-specific offset test coverage.
- Adjusts gfx1250 benchmark parameters and documentation.
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_prefetchgl2_streamk_guard.py | Updated as part of this pull request. |
| projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_gl2_prefetch_offset.py | Updated as part of this pull request. |
| projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/tdm_multicast_gfx1250.yaml | Updated as part of this pull request. |
| projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py | Updated as part of this pull request. |
| projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py | Updated as part of this pull request. |
| projects/hipblaslt/tensilelite/Tensile/KernelWriter.py | Updated as part of this pull request. |
| projects/hipblaslt/tensilelite/Tensile/Components/GL2Prefetch.py | Updated as part of this pull request. |
| projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py | Updated as part of this pull request. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+6516
to
+6517
| if state["GlobalSplitU"] > 1 or state["GlobalSplitU"] == -1: | ||
| reject(state, printRejectionReason, "Currently PrefetchGL2 does not support GSU") |
| - HalfPLR: [2] | ||
| - StoreRemapVectorWidth: [0] | ||
| - GlobalSplitU: [1,2,3] | ||
| - GlobalSplitU: [1,3] |
| # or Stream-K other than DP-first (StreamK==3). GSU is supported; with a | ||
| # workgroup cluster it forces GlobalSplitUWorkGroupMappingRoundRobin on so the | ||
| # cluster's peers share a K chunk. May remove these limitations in the future. | ||
| # Currently we do not support GSU, StaggerU, StreamK and general batch. May remove these limitations in the future. |
✅ 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. |
talumbau
approved these changes
Sep 21, 2026
Collaborator
|
#11851 should not have landed anyway due to patch coverage report anyway: 
That's just too low to land. |
tony-davis
enabled auto-merge (squash)
September 21, 2026 18:40
tony-davis
disabled auto-merge
September 21, 2026 19:21
1 task done
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.
Reverts #11851
Fixes #12384
Motivation
TensileLite coverage / Linuxhas been red ondevelopsince #11851 landed, andComponent CI Summaryis a required check, so every PR that triggers the tensilelite lane is blocked. The gate reports one regressed file:#11851 did ship tests (
test_gl2_prefetch_offset.pyandtest_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 on3082a49c, 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.pyandValidParameters.py, plus the two unit tests and thetdm_multicast_gfx1250.yamltest config that came with it.GL2Prefetch.pyat this PR's head is byte-identical to the file on3082a49c, the last commit where the coverage lane was green (blob4256c495cin both). So this clears the gate without touchingcoverage-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 / Linuxon this PR. It re-measuresGL2Prefetch.pyfrom source against the committed baseline, so a green run is the direct check that the floor is satisfied again.Test Result
Pending this PR's own CI run.
Submission Checklist
Risk level
Low (2). This restores the exact
GL2Prefetch.pycontent that was ondevelopfor 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.