Repository navigation
fix(tensilelite): fix VGPR occupancy calculation on gfx11 - #9708
Open
mgehre-amd wants to merge 8 commits into
Open
mgehre-amd wants to merge 8 commits into
mgehre-amd wants to merge 8 commits into
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
mgehre-amd
commented
Jul 22, 2026
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (48.93%) 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 #9708 +/- ##
========================================
Coverage 75.95% 75.95%
========================================
Files 2880 2880
Lines 479025 479059 +34
Branches 71958 71967 +9
========================================
+ Hits 363828 363862 +34
- Misses 90338 90340 +2
+ Partials 24859 24857 -2
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
mgehre-amd
force-pushed
the
matthias.tensilelite-per-simd-vgpr-occupancy
branch
3 times, most recently
from
July 23, 2026 11:40
4028a44 to
4db4853
Compare
mgehre-amd
marked this pull request as ready for review
July 23, 2026 12:34
mgehre-amd
force-pushed
the
matthias.tensilelite-per-simd-vgpr-occupancy
branch
from
July 23, 2026 16:18
1b887e5 to
3651192
Compare
This was referenced Jul 24, 2026
mgehre-amd
marked this pull request as draft
July 24, 2026 12:49
mgehre-amd
changed the base branch from
develop
to
users/mgehre/maxoccupancy-range-64
July 28, 2026 13:16
mgehre-amd
force-pushed
the
matthias.tensilelite-per-simd-vgpr-occupancy
branch
2 times, most recently
from
July 28, 2026 14:06
52d71eb to
813013c
Compare
mgehre-amd
force-pushed
the
matthias.tensilelite-per-simd-vgpr-occupancy
branch
from
August 12, 2026 09:31
813013c to
07a7156
Compare
Base automatically changed from
users/mgehre/maxoccupancy-range-64
to
develop
September 2, 2026 06:10
Contributor
|
This pull request has been inactive for 25 days and will be marked as stale. If you would like to keep this PR open, please:
This PR will be automatically closed in 5 days if no further activity occurs. |
mgehre-amd
changed the base branch from
develop
to
users/mgehre/tensilelite-gfx11-wgp-lds-pool
October 1, 2026 09:57
mgehre-amd
force-pushed
the
matthias.tensilelite-per-simd-vgpr-occupancy
branch
from
October 21, 2026 14:41
7292f0d to
5190e4e
Compare
Code Coverage Report for gfx950Summary
Artifacts
Commit Hashes |
Contributor
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #9708 — use the per-SIMD VGPR file for gfx11 occupancy (tensilelite) — 🟢 First-round: no closer look needed — on a real gfx1151 the new model tracks the hardware far better than the old one, and it equals LLVM's own occupancy in every case we measured; stacked on #12363
Utility: 🟢 Positive — fixes a large gfx11 occupancy undercount that mis-sized store batching, held to today's kernel occupancy by the logic-file changes, with gfx12 and gfx1250 left for follow-ups. Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted First-round review: a local qwen3.6 model made the independent first pass; the orchestrator (Claude Opus 5.5) verified it — the threads against the code, the stack against its base branch, the logic-file diff, HIP's occupancy API on a gfx1151 against the old and new models, LLVM's occupancy remarks, and the CI logs — and owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection. |
Base automatically changed from
users/mgehre/tensilelite-gfx11-wgp-lds-pool
to
develop
October 22, 2026 10:47
getVgprOccupancy/getMaxRegsForOccupancy assumed a flat 512-VGPR per-SIMD file (MaxVgpr*2) for every doubleVgpr arch. On gfx11 the real per-SIMD file is 1536 (gfx1100/gfx1101/gfx1151) or 1024, so occupancy was badly under-counted, which inflated the auto NumElementsPerBatchStore VGPR budget in setOccupancy and dropped achieved occupancy on store-heavy kernels. Make regCaps["PhysicalMaxVgpr"] the per-SIMD VGPR file (set per arch in hardware_caps.hpp) and use it as the occupancy denominator: PhysicalMaxVgpr in wave32/unified (doubleVgpr) mode, PhysicalMaxVgpr//2 in wave64. Allocation granularity follows the 64-block invariant (file // 64): 24/16 on gfx11, unchanged 4/8 on gfx8-10. Changes: - PhysicalMaxVgpr now means the per-SIMD file, not the per-wave arch+acc envelope: 1536/1024 on gfx11; gfx1250 stays 1024 (its wave32 file, == MaxVgpr); 512 elsewhere. - gfx8/gfx9/gfx10 occupancy is unchanged (PhysicalMaxVgpr == old MaxVgpr*2, and wave64 PhysicalMaxVgpr//2 == old MaxVgpr; granularity identical). gfx1250 is corrected from an over-counted 2048 to its true 1024 per-SIMD file. Only gfx11 changes; its shipped wave64 fp32 kernels assemble bit-identically (verified), so this is codegen-neutral for them and only corrects reported occupancy. - Added unit tests for the per-SIMD occupancy path, gfx11 24-VGPR granularity, and the per-arch PhysicalMaxVgpr values.
getVgprOccupancy applies the per-SIMD PhysicalMaxVgpr denominator only for gfx11; gfx1250 keeps the legacy MaxVgpr*2 model (its physically-correct 1024 per-SIMD file is deferred to separate benchmarking). This keeps gfx8/gfx9/ gfx10/gfx12 and gfx1250 byte-identical to base; only gfx11 changes. Adds test_gfx1250_keeps_legacy_maxvgpr_occupancy and restores the version field to the occupancy test stub (getVgprOccupancy now reads states.version).
…size getVgprOccupancy computed the waves-per-SIMD that one workgroup occupies as ceil(max(numThreads, 256) / 256). That divisor hardcodes wave64: 256 threads span 4 waves over 4 SIMDs only when a wave is 64 lanes wide. On a wave32 gfx11 kernel the same 256 threads are 8 waves, so the workgroup occupies 2 wave slots per SIMD, not 1, and every wave32 workgroup larger than 128 threads was counted below its real wave-slot cost. The divisor is now WavefrontSize * 4 (4 SIMDs per CU on CDNA, per WGP on RDNA), which reproduces the old value on wave64 and corrects wave32. gfx12 is held on the legacy formula. It is wave32 as well, so the corrected divisor would move its occupancy too, and that has not been benchmarked; a TODO marks it. Changes: - Solution.py is a comment only: the MaxOccupancy LDS floor (DeviceLDS//MaxOccupancy) is a per-workgroup budget and is not affected by this change. - The earlier gfx11 128 KB WGP LDS pool correction that this commit used to carry now comes from the base commit instead, so the before/after kernel counts measured for the combined change no longer describe this commit alone and have been dropped rather than re-attributed.
doubleVgpr reads as a wave64 or AGPR flag, but KernelWriter sets it for ArchAccUnifiedRegs parts *or* any wave32 kernel, so it is True for 2521 of the 2537 shipped gfx1151 kernels. The gfx11 occupancy tests pinned it to False and so exercised the wave64 arm, which production never takes on those kernels. Pass the production value instead, and name the rule where the flag selects the VGPR file size so the next reader does not have to re-derive it. Changes: - The four gfx11 getOccupancy tests use 8 VGPRs, so the VGPR term is slack under either flag and no expected value moves. This buys consistency with production, not extra coverage; the gfx11 tests that do exercise the VGPR path already passed doubleVgpr=True. - The buildtime stub now derives the flag from its own wavefront size rather than reusing arch_acc_unified, so its gfx11 stand-in stops contradicting itself. - The ternary selecting totalVgprs is reordered to lead with the common case rather than the negation; the two arms are unchanged in value. - doubleVgpr=False is left in place where it is deliberate: the wave64 arm test, the pre-MFMA accvgprs=0 test, and the gfx908 half of the CDNA parametrisations.
The occupancy-model corrections earlier in this branch raise the occupancy the model grants, and on hardware that lands as a real increase: across all eight gfx11 targets 3148 of 13022 shipped solutions gain residency and none lose it. Pin the ones that gain back to the occupancy they have on develop, so the model corrections land occupancy-neutral and the gains can be adopted deliberately, per family, once they have been benchmarked. MaxOccupancy pads a kernel's LDS up to pool // MaxOccupancy, which getLdsLimitedOccupancy then rounds up to the 1024-byte gfx11 granule. For each raised solution the pin is the largest MaxOccupancy whose rounded floor lands exactly on the develop value, so none is pushed below it. 3108 solutions across 273 logic files: gfx1100 118, gfx1101 117, gfx1102 537, gfx1103 537, gfx1150 527, gfx1151 197, gfx1152 537, gfx1153 538. Changes: - Occupancy is computed with the hardware model min(16, file // ceil_granule(vgpr)) against the 128 KB WGP LDS pool, not with hipModuleOccupancyMaxActiveBlocksPerMultiprocessor. That API assumes a power-of-two VGPR allocation granule and rounds by masking, so on gfx1151 (24-register granule) it is not even monotone in the VGPR count -- it reports a *higher* occupancy for a kernel declaring *more* registers. The model is validated against resident-workgroup counts on gfx1151 at 24/24 points, where that API is right at only 12/24. - Solutions are matched to logic-file entries by their rank under the library's global solution index, not by libraryLogicIndex. That field is neither a position in the YAML list -- some libraries number sparsely, e.g. 0..21 for a 17-entry file -- nor unique, since a library that merges several logic files gets a 0-based run from each. Every alignment is verified by comparing the macro tile of each YAML entry against the generated kernel's MT token, and all eight targets map with nothing left over. - Verified by rebuilding all eight targets and re-deriving: 12994 of 13022 solutions sit exactly at their develop occupancy, none below. Seven of the eight targets are exact. - The 28 that remain above are gfx1151 solutions wanting occupancy 20 or 40, which no MaxOccupancy can express: 131072/20 rounds up to 7168 and yields 18, 131072/40 rounds up to 4096 and yields 32. - Only MaxOccupancy values change; no kernel is otherwise retuned. Rebased onto develop, which had meanwhile replaced the solutions this commit pinned in ten of the files: - The eight gfx1150 HHS/BBS TN+NN GridBased libraries were retuned, each 64-73 solutions replaced by 100 new ones. The old per-solution pins describe solutions that no longer exist, so they are dropped and re-derived against the new set: 204 of the 800 solutions gain residency (8->9 x108, 4->5 x60, 5->7 x28, 5->6 x8), all of them expressible, pinned at MaxOccupancy 9/5/6. Re-verified by rebuilding gfx1150: all 1435 solutions whose kernel could be identified now sit exactly at their develop occupancy, none above, none below. That covers the thirty gfx1150 files this rebase did not touch as well, so the pins derived before the rebase still land after it. - The two navi31 HHS NN GridBased libraries were deleted in favour of gfx1100 Origami prediction libraries. Their pins go with them, which is why gfx1100 drops from 120 to 118. - The thirteen gfx1150 solutions left out of that re-verification are HSS/BSS entries whose recorded KernelNameMin no longer identifies a unique generated kernel; they are unpinned either way.
getVgprOccupancy now derives the waves-per-SIMD cost of a workgroup from kernel["WavefrontSize"], but test_occupancy_hip.py's _make_writer builds states with only archCaps/regCaps/version. Every case it parametrises is gfx9, which takes that path, so on a runner that has one of those GPUs test_hip_occupancy_matches_tensile raised AttributeError before it compared any occupancy value. The failure is invisible on other hosts because the test skips when the GPU is absent. Pass the wave size the kernel is actually assembled with, matching the stubs in test_occupancy.py and test_occupancy_buildtime.py. Changes: - The wave size is a parameter defaulting to 64 rather than a literal, so the stub and _build_minimal_kernel_asm cannot drift apart silently; both the parametrised test and test_gfx950_occupancy_not_overcounted use it. - Added test_writer_stub_supports_getoccupancy, which runs getOccupancy through the stub for each target in OCCUPANCY_CASES. It needs no GPU, so unlike the test it guards it actually executes on every runner; without the fix it reproduces the AttributeError.
mgehre-amd
force-pushed
the
matthias.tensilelite-per-simd-vgpr-occupancy
branch
from
October 22, 2026 10:47
5190e4e to
6f1368c
Compare
CongMa13
reviewed
Oct 22, 2026
CongMa13
left a comment
CongMa13
left a comment
Contributor
There was a problem hiding this comment.
Please address the inline occupancy-limit issue below. It can allow gfx11 wave32 code generation to allocate beyond the VGPR budget used by the new occupancy calculation.
mgehre-amd
removed
documentation
project: rocrand
ci:hipsparselt-fast
project: hip-kernel-provider
rocke: platform
setOccupancy and the GSU path limit register-pool growth to PhysicalMaxVgpr // occupancy. On gfx11 occupancy now counts resident workgroups, and each workgroup takes numThreads / (wavefront * 4) waves on every SIMD, so dividing the whole file by it lets codegen allocate several times the per-wave budget the occupancy was computed for. A gfx1151 wave32 256-thread kernel at occupancy 8 could grow to 192 VGPRs, which leaves room for 4 workgroups, not 8. The division is also wrong for wave64, where the file is half of PhysicalMaxVgpr, and it ignores the 16/24-VGPR allocation granule. On gfx11 the VGPR limit is now the maxVgprs that getMaxRegsForOccupancy already returns at both call sites, i.e. the largest per-wave count that keeps the modeled occupancy. Changes: - The two duplicated limit blocks become one setOccupancyLimits helper. Its early return for gfx12 / gfx117x is the same test as the existing guard, negated; those targets still get no limit. - Other architectures, and the SGPR limit everywhere, keep PhysicalMaxSgpr / PhysicalMaxVgpr // occupancy, so their codegen does not change. gfx11 has no SGPR file limit to tighten. - When even the current allocation fits no workgroup (e.g. wave32, 1024 threads, over 192 VGPRs on a 1536 file), getMaxRegsForOccupancy returns (vgprs, 1); the limit then stays at the old non-binding value instead of freezing the pool at its current size. - A tighter limit makes growth past it throw, which drops the kernel at generation time instead of silently lowering its residency. Shipped gfx11 solutions still have to be rebuilt to confirm none is rejected. - getVgprOccupancy's `multiplier` is renamed wavesPerSimdPerWorkgroup: it is the waves one workgroup places on each SIMD, not waves per workgroup. Its `maxOccupancy` becomes maxWorkgroups, since occupancy is a workgroup count. - Occupancy deliberately stays a workgroup count (as on develop): the launch grid (cuCount * CUOccupancy), Origami, the LDS term, the MaxOccupancy floor and the SIA2 guard all need workgroups; only this limit needs a per-wave figure, and a wave count would still need the granule-aware maxVgprs. - CDNA has the same over-allowance for >256-thread workgroups (develop's gfx942 limit lets a 512-thread kernel at 4 workgroups grow to 2). That is left as a TODO because it changes CDNA codegen and needs its own testing. Tested: new test_occupancy cases grow the pool to its limit across gfx1100/1150/1151, wave32/64, 128/256 threads and four VGPR/LDS mixes and require the modeled occupancy to hold. With the previous limit patched back in, 40 of the 52 new tests fail. TensileLite unit lane: 10464 passed, 0 failed.
Capping the gfx11 VGPR pool at the modeled budget changes how many VGPRs the kernels allocate, so the MaxOccupancy pins that held them at their develop occupancy no longer land: 221 solutions across the eight gfx11 targets ended up above the hardware occupancy they had before the 128 KB WGP LDS pool change, none below. Re-pin those 221 so every kernel is back at that occupancy. The reference is develop just before the LDS-pool change (540442e), not today's develop: that change and the occupancy model fixes are meant to land occupancy-neutral, and gains get adopted per family later. Changes: - 59 logic files, MaxOccupancy values only; no kernel is otherwise retuned. - Some pins tighten because the capped pool now allocates fewer VGPRs: MO9 -> MO5 on gfx1102/1103/1150/1152/1153, MO11 -> MO6, MO21 -> MO10 and MO64 -> MO11 on gfx1151. - Files the earlier pinning never reached get their first pin: the gfx1100 Origami libraries added since, and the *_HAH and the logic files whose names do not spell their kernels' problem type (e.g. navi32 Alik_Bljk HHS_UserArgs builds HHS_BH kernels). - The pins were derived by matching every logic solution to its built kernel. Logic files' KernelNameMin is not a reliable key: non-Aux files carry their Aux twin's names, and some files have none. Matching used name tokens with the file's own problem type, cross-checked against the built libraries' own solution-to-kernel records (one run of libraryLogicIndex per merged logic file); the two never disagreed. Verified by building all eight gfx11 targets at the reference and with these pins and scoring every matched kernel's occupancy from its code object (per-SIMD VGPR file and allocation granule, waves per SIMD per workgroup, 128 KB WGP LDS pool): 13133 of 13172 solutions are exact, none above or below. The 39 not matched are gfx1151's int4 w4a16 libraries and one Equality library with non-standard entries; this change leaves them as develop has them. All 11057 kernels were generated in both builds.
mgehre-amd
added a commit
that referenced
this pull request
Nov 2, 2026
…8309) Fixes #14455 ## Motivation On RDNA3/3.5 (gfx11) the VGPR file has 4 banks (`bank = vgpr % 4`). A WMMA reads its A, B and C operands in lockstep; when two of them share a bank the read serializes and the WMMA takes 34 instead of 32 cycles. TensileLite's register allocator currently *forces* the A operand onto the same bank as the C accumulator, so every WMMA in a gfx11 GEMM pays this penalty. ## Technical Details For WMMA kernels on gfx11 (`HasWMMA_V1` and `EnableMatrixInstruction`), the allocator now places A on a bank different from C (B is already placed on an odd bank). The change is a few lines in `KernelWriter.py`. - No new solution parameter: it is on by default. Kernel names do not change, so no characterization goldens change on any architecture. - Non-WMMA gfx11 kernels (f32 VALU, dot2) keep the existing placement. Applying the rule to them as well gave mixed f32 results (-3%..+8% depending on the tile), so they are left unchanged. - gfx12 (WMMA V2) and WMMA V3 targets are unchanged. ## Test Plan - Built the full gfx1151 device library with and without the change and compared all 3056 kernels (VGPR count, LDS, waves/SIMD) and every `v_wmma` operand bank in the disassembly. - `hipblaslt-bench` A/B: one bench binary, device library swapped via `HIPBLASLT_TENSILE_LIBPATH` (only TensileLite Python differs), shader clock pinned at 2100 MHz, `--adaptive --rotating 512 --flush --use_gpu_timer`, 3 interleaved reps with alternating order, 36 configs (bf16/f16 x TN/NN x 9 shapes from 512^3 to 8192^2x4096). - TensileLite unit lane (`tensilelite/Tests/unit`). - New CPU-only characterization test `test_wmma_bank_distinct_gfx1151_char.py`: asserts every emitted `v_wmma` of a gfx1151 bf16 MT96x128x64 kernel reads A, B and C from three distinct banks. ## Test Result - Resources (gfx1151, 3056 kernels): 6 kernels change waves/SIMD under the current occupancy model. Four bf16/f16 NN bias MT128x32x32 kernels go from 10 to 9 (VGPRs 128->168, from the store-batch heuristic spending the freed budget); two int8 kernels gain. The affected bf16 kernel measured neutral (-0.01% to +0.13% on the 4 shapes that select it). With #9708 (per-SIMD VGPR occupancy model) applied, no kernel changes waves/SIMD. - WMMAs with an A/C bank collision: 13,410 -> 0. WMMAs with any operand-bank collision: 100% -> 11.5%. - Performance (gfx1151, based on develop a0a98e0, same solution selected in every case): geomean **+1.21%**, all 36 configs faster, best +2.29% (bf16 TN 4096x2048x11008), worst +0.14% (bf16 NN 2048^3). Large compute-bound bf16/f16 shapes gain +1.4% to +2.3%. - Unit lane: 11,319 passed, all 796 golden snapshots unchanged. 8 failures, none touching this change: 6 in `test_arch_steppings.py` (gfx1250 Softmax/LayerNorm/AMax generators, `No module named 'rocisa.code'` in a subprocess, local test environment) and 2 in `test_canonical_imports.py`, which flags legacy `Tensile.*` imports in existing develop test files (`test_stinky_register_allocation.py`, `test_r3_declared_count_occupancy_gfx1250_char.py`, `test_register_overflow_deferral.py`). - The new characterization test passes, and fails with the allocator change reverted. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. Register numbering changes only for gfx11 WMMA kernels; instruction sequences are otherwise identical. Verified occupancy-neutral across the shipped gfx1151 library.
mgehre-amd
changed the title
This branch has not been deployed
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.
Summary
Fix gfx11 VGPR occupancy estimates used to budget registers for automatic store batching. Underestimating occupancy can allow store batching to consume too many VGPRs and reduce kernel residency.
PhysicalMaxVgpr(1536 for gfx1100/gfx1101/gfx1151 and 1024 for other gfx11 targets, expressed in wave32 units) instead ofMaxVgpr(registers addressable by one wave) for occupancy calculations. They just happen to have the same value on gfx9.gfx12 retains its existing workgroup multiplier; gfx1250 also retains its existing VGPR capacity and allocation model.
The gfx11 logic files adjust
MaxOccupancyto preserve existing kernel occupancy where possible. Follow-up tuning will make use of the fixed occupancy model.Validation
ISSUE ID : ROCm/TheRock#6771
Partially addresses the issue above for gfx11. gfx12 fixes are deferred to a follow up PR.