Repository navigation
fix(tensilelite): use the 1024-byte LDS granularity on gfx11 - #10078
Merged
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
mgehre-amd
force-pushed
the
users/mgehre/maxoccupancy-range-64
branch
3 times, most recently
from
August 13, 2026 09:14
85b7055 to
a80d66f
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #10078 +/- ##
===========================================
- Coverage 63.78% 63.78% -0.00%
===========================================
Files 2825 2825
Lines 469128 469124 -4
Branches 69270 69270
===========================================
- Hits 299226 299223 -3
Misses 148858 148858
+ Partials 21044 21043 -1
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
mgehre-amd
force-pushed
the
users/mgehre/maxoccupancy-range-64
branch
from
September 1, 2026 16:06
a80d66f to
6dd0548
Compare
Base automatically changed from
users/mgehre/maxoccupancy-range-64
to
develop
September 2, 2026 06:10
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
getLdsLimitedOccupancy rounded every target's LDS to a 256-byte granule. gfx11 allocates LDS in 1024-byte granules, so occupancy was over-estimated wherever a kernel's LDS fell between two 1024-byte boundaries: the model fitted a workgroup the hardware cannot place. The granule is now the archCaps["LdsGranularity"] hardware capability, 1024 on gfx11 and 256 elsewhere, which getOccupancy passes to getLdsLimitedOccupancy. Changes: - getLdsLimitedOccupancy takes the granule as a required argument rather than defaulting it, so no caller can silently assume 256; the two gfx950 unit-test callers pass it explicitly. Only gfx11 changes behaviour. - The updateOccupancyFromMaxVgpr test stub builds archCaps by hand, so it gains the new key too; it mirrors the real cap (1024 for the gfx11 case the stub models, 256 for CDNA) rather than hardcoding one granule. - No characterization golden covers gfx11 codegen (the _codegen data is gfx90a, gfx908, gfx942, gfx1201, gfx1250), so none needed re-recording. Confirmed by running Tests/unit before and after: the snapshot tally is unchanged. - Touches rocisa C++: a rocisa reinstall is needed to pick up the new cap.
mgehre-amd
force-pushed
the
users/mgehre/gfx11-lds-granule
branch
from
September 2, 2026 16:23
d599c32 to
617f656
Compare
mgehre-amd
marked this pull request as ready for review
September 2, 2026 16:27
|
🎉 All checks passed! This PR is ready for review. |
mgehre-amd
changed the title
KKyang
approved these changes
Sep 3, 2026
tony-davis
pushed a commit
that referenced
this pull request
Sep 3, 2026
#11672) This reverts commit fd61638 - PR #10078 ## Motivation hipblaslt build failed with this PR. ## Technical Details ldsGranularity exception error observed with this PR. ## Test Plan Build hipblalslt. Running hipblaslt CI tests. ## Test Result Build was successful after revert. Co-authored-by: Babak <bpoursar@amd.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
mgehre-amd
added a commit
that referenced
this pull request
Sep 9, 2026
archCaps has two independent producers: rocisa's hardware_caps.hpp and stinkytofu's python_module/src/HardwareCaps.cpp, reached through rocisa_stinkytofu_adaptor when ROCISA_BACKEND=stinkytofu. #10078 added LdsGranularity to the first only, so any build served by the stinkytofu provider got a fully populated archCaps missing exactly that key, and getOccupancy died with KeyError: 'LdsGranularity' in globalWriteElementBatch. That is what got #10078 reverted in #11672. The reverted build's own capability dump is the evidence: it lists DeviceLDS but neither CacheLineBytes nor NumXCD, and those two exist only in rocisa's initArchCaps (added 2026-07-24, well before that build). The table is a union over every probed ISA, so their absence identifies stinkytofu as the provider. No stale rocisa binary is involved -- a clean build fails the same way. Because the stinkytofu provider probes comgr against the target ISA name and ignores the host GPU, this reproduces on any machine with -DGPU_TARGETS=gfx1250; the hardware is not needed. Changes: - The expression matches rocisa's exactly, so both providers stay in lock-step; verified by linking initArchCaps and dumping it (gfx11 -> 1024, gfx942/gfx950/ gfx1030/gfx1201/gfx1250 -> 256). - Tests pin it on both sides of the adaptor: presence in the adaptor's dynamic caps test, and the gfx1250 value in the comgr cap table test, so a future cap added to one provider and not the other fails here rather than in codegen.
mgehre-amd
added a commit
that referenced
this pull request
Sep 9, 2026
…fu table regCaps has two independent producers: rocisa's hardware_caps.hpp and stinkytofu's python_module/src/HardwareCaps.cpp, reached through rocisa_stinkytofu_adaptor when ROCISA_BACKEND=stinkytofu. The preceding commit raised PhysicalMaxSgpr to 2048 on gfx11 in the first only; this brings the second in line so the two cannot disagree about the same key. Changes: - The expression matches rocisa's exactly, so both providers stay in lock-step. This is the same drift class that got #10078 reverted in #11672 (a cap added to one provider and not the other). - No behavioural effect today: STINKYTOFU_ALL_ARCHS is Gfx1250/Gfx1250v0 only (cmake/StinkytofuArchList.cmake), so HardwareCaps::query returns an empty result for every gfx11 triple and the new branch is unreachable until a gfx11 backend exists. The change is for consistency, not to fix an observed failure. - Tests can only pin the non-gfx11 branch on this side, for the same reason: test_comgr.py asserts the gfx1250 value stays 800 so the expression cannot be folded back to a flat constant, and the adaptor's dynamic caps test asserts the key is present. The gfx11 value stays covered on the rocisa side in Tensile/Tests/unit/test_occupancy.py. - Not addressed here: stinkytofu also reports MaxWavesPerSimd = 10 for gfx11 where rocisa reports 16. That is a separate pre-existing divergence and changing it would alter occupancy on the stinkytofu path.
mgehre-amd
added a commit
that referenced
this pull request
Sep 10, 2026
…fu table regCaps has two independent producers: rocisa's hardware_caps.hpp and stinkytofu's python_module/src/HardwareCaps.cpp, reached through rocisa_stinkytofu_adaptor when ROCISA_BACKEND=stinkytofu. The preceding commit raised PhysicalMaxSgpr to 2048 on gfx11 in the first only; this brings the second in line so the two cannot disagree about the same key. Changes: - The expression matches rocisa's exactly, so both providers stay in lock-step. This is the same drift class that got #10078 reverted in #11672 (a cap added to one provider and not the other). - No behavioural effect today: STINKYTOFU_ALL_ARCHS is Gfx1250/Gfx1250v0 only (cmake/StinkytofuArchList.cmake), so HardwareCaps::query returns an empty result for every gfx11 triple and the new branch is unreachable until a gfx11 backend exists. The change is for consistency, not to fix an observed failure. - Tests can only pin the non-gfx11 branch on this side, for the same reason: test_comgr.py asserts the gfx1250 value stays 800 so the expression cannot be folded back to a flat constant, and the adaptor's dynamic caps test asserts the key is present. The gfx11 value stays covered on the rocisa side in Tensile/Tests/unit/test_occupancy.py. - Not addressed here: stinkytofu also reports MaxWavesPerSimd = 10 for gfx11 where rocisa reports 16. That is a separate pre-existing divergence and changing it would alter occupancy on the stinkytofu path.
mgehre-amd
added a commit
that referenced
this pull request
Sep 11, 2026
…fu table regCaps has two independent producers: rocisa's hardware_caps.hpp and stinkytofu's python_module/src/HardwareCaps.cpp, reached through rocisa_stinkytofu_adaptor when ROCISA_BACKEND=stinkytofu. The preceding commit raised PhysicalMaxSgpr to 1696 on gfx11 in the first only; this brings the second in line so the two cannot disagree about the same key. Changes: - The expression matches rocisa's exactly, so both providers stay in lock-step. This is the same drift class that got #10078 reverted in #11672 (a cap added to one provider and not the other). - No behavioural effect today: STINKYTOFU_ALL_ARCHS is Gfx1250/Gfx1250v0 only (cmake/StinkytofuArchList.cmake), so HardwareCaps::query returns an empty result for every gfx11 triple and the new branch is unreachable until a gfx11 backend exists. The change is for consistency, not to fix an observed failure. - Tests can only pin the non-gfx11 branch on this side, for the same reason: test_comgr.py asserts the gfx1250 value stays 800 so the expression cannot be folded back to a flat constant, and the adaptor's dynamic caps test asserts the key is present. The gfx11 value stays covered on the rocisa side in Tensile/Tests/unit/test_occupancy.py. - Not addressed here: stinkytofu also reports MaxWavesPerSimd = 10 for gfx11 where rocisa reports 16. That is a separate pre-existing divergence and changing it would alter occupancy on the stinkytofu path.
mgehre-amd
added a commit
that referenced
this pull request
Sep 11, 2026
…fu table regCaps has two independent producers: rocisa's hardware_caps.hpp and stinkytofu's python_module/src/HardwareCaps.cpp, reached through rocisa_stinkytofu_adaptor when ROCISA_BACKEND=stinkytofu. The preceding commit raised PhysicalMaxSgpr to 1696 on gfx11 in the first only; this brings the second in line so the two cannot disagree about the same key. Changes: - The expression matches rocisa's exactly, so both providers stay in lock-step. This is the same drift class that got #10078 reverted in #11672 (a cap added to one provider and not the other). - No behavioural effect today: STINKYTOFU_ALL_ARCHS is Gfx1250/Gfx1250v0 only (cmake/StinkytofuArchList.cmake), so HardwareCaps::query returns an empty result for every gfx11 triple and the new branch is unreachable until a gfx11 backend exists. The change is for consistency, not to fix an observed failure. - Tests can only pin the non-gfx11 branch on this side, for the same reason: test_comgr.py asserts the gfx1250 value stays 800 so the expression cannot be folded back to a flat constant, and the adaptor's dynamic caps test asserts the key is present. The gfx11 value stays covered on the rocisa side in Tensile/Tests/unit/test_occupancy.py. - Not addressed here: stinkytofu also reports MaxWavesPerSimd = 10 for gfx11 where rocisa reports 16. That is a separate pre-existing divergence and changing it would alter occupancy on the stinkytofu path.
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.
KernelWriterAssembly.getLdsLimitedOccupancyrounded every target's LDS usage to a 256-byte granule:gfx11 allocates LDS in 1024-byte granules, so occupancy was over-estimated wherever a kernel's LDS fell strictly between two 1024-byte boundaries: the model fitted a workgroup the hardware cannot place.
The granule becomes a hardware capability,
archCaps["LdsGranularity"](1024 on gfx11, 256 everywhere else), whichgetOccupancypasses togetLdsLimitedOccupancy.Only gfx11 changes behaviour. Follow up PRs should address this for gfx10 and gfx12, where the LDS granularity is also 1024.
Validation
Rebuild all gfx11 kernels. 19 changes their assembly (they are occupancy limited due to LDS with the correct rounding; so they used additional VGRPs). All of those were benchmarked and their A/B perf was within noise.
Tensile/Tests/unitoccupancy lane (test_occupancy.py,test_occupancy_buildtime.py,test_occupancy_hip.py,test_custom_kernel_occupancy.py,test_PrefetchAcrossPersistent.py): 144 passed, 11 skipped.Tensile/Tests/unit/characterization/_codegen: 254 passed, 60 snapshots passedpytest -m unit -n4 Tensile/Tests/unitwas run on this branch and on its base commit and compared test-by-testFixes #10049