Repository navigation
Revert "fix(tensilelite): use the 1024-byte LDS granularity on gfx11 … - #11672
Merged
Merged
Conversation
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. |
bnemanich
approved these changes
Sep 3, 2026
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (35.22%) 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 #11672 +/- ##
===========================================
- Coverage 64.11% 64.11% -0.00%
===========================================
Files 2826 2826
Lines 469837 469834 -3
Branches 69389 69389
===========================================
- Hits 301192 301189 -3
Misses 147536 147536
Partials 21109 21109
*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
|
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.
This reverts commit fd61638 - PR #10078
Motivation
hipblaslt build failed on gfx1250 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.