Occupancy Math on the AMD MI355X (CDNA4): A From-First-Principles Guide · Shekhar Pandey

Ask a GPU kernel engineer how their kernel is doing and occupancy comes up within a sentence or two. It’s the number everyone quotes and the dial everyone reaches for — and, in my experience, the metric people understand least. Most treat it as an opaque percentage the profiler hands back. It isn’t. Occupancy is fully derivable by hand from a kernel’s resource usage and a handful of fixed hardware limits, and being able to do that derivation changes how you tune. TL;DR. On MI355X, occupancy — the fraction of a SIMD’s wavefront slots your kernel keeps filled — is set by whichever of four resource limiters runs out first: VGPRs, SGPRs, LDS, or workgroup/barrier slots. Each is just a fixed hardware budget divided by what your kernel spends, so you can compute the ceiling by hand from the binary: occupancy = min(those four limiters). The VGPR file is 512 per lane, shared by regular and accumulator registers (not a separate AccVGPR pool). And maximizing occupancy is usually the wrong goal: in a measured MXFP8 MFMA sweep below, the matrix core stays at ~97% of peak even as occupancy falls to a fraction of full — its throughput tracks matrix-engine utilization, not how full the SIMD is. This is a from-first-principles guide to occupancy math on the AMD Instinct MI355X (CDNA4, gfx950). We’ll build it from the silicon up: what the hardware budget actually is, which four resources cap how many wavefronts can go resident, and how to compute a kernel’s occupancy ceiling on paper and then confirm it with rocprofv3. The worked examples lean on MXFP8 GEMM tiles — the kind of kernel where these numbers decide whether the kernel is fast. The post is in three parts: Part 1 — The MI355X architecture. The CU, the four SIMDs, the wavefront, the Matrix Core, and the memory hierarchy that feeds them — the resources the math counts. Part 2 — Occupancy math. The definition, the four limiters (VGPRs, SGPRs, LDS, and workgroup/barrier slots), the per-SIMD vs per-CU split that trips everyone up, worked examples, and how to measure occupancy for real.

Part 3 — Better performance at lower occupancy. The twist at the end: once you can compute the ceiling, why reaching for it is often the wrong move. Little’s Law, ILP versus occupancy, and a microbenchmark where the matrix core stays saturated even as occupancy collapses. Occupancy is worth understanding precisely — both so you can fix the kernels that are genuinely occupancy-limited, and so you can recognize the ones that aren’t. Part 1 — The MI355X architecture Occupancy is, start to finish, a story about how a fixed pool of resources gets divided among the work you launch. So before any of the math lands, you need a clear mental model of the chip — and in particular, which resources are private and which are shared. We’ll build it top-down and keep returning to that one distinction, because it’s the hinge the whole calculation turns on. The chip, top down The MI355X is CDNA4, ISA target gfx950. At the top level it’s eight Accelerator Complex Dies (XCDs) stitched together over fourth-gen Infinity Fabric, totaling 256 Compute Units — 32 per XCD — clocked up to 2.4 GHz. Wrapped around the dies are 288 GB of HBM3E at 8 TB/s, fronted by 256 MB of last-level Infinity Cache. Each XCD carries its own L2 slice, which is why scheduling that keeps a tile’s traffic on one die (XCD-aware swizzling) pays off — but that’s a locality story, not an occupancy one. For occupancy, the only unit you reason about is the Compute Unit. Everything above it governs how data moves; occupancy is decided one CU at a time. So that’s where we zoom in. Inside a Compute Unit A CU is four SIMD units plus shared infrastructure. Each SIMD is 64 lanes wide and owns a private register file. The pieces that matter: VGPRs — a 512-entry-per-lane vector register file, private to the SIMD. Allocated per wave. This is usually the resource that caps occupancy. AccVGPRs — the matrix accumulators, carved from that same 512 file.

MFMA instructions can accumulate here. On CDNA4 the file is split between regular VGPRs (≤256/wave) and accumulation VGPRs (≤256/wave), but the two share one 512-entry budget — a wave’s regular plus accumulator count is what’s measured against 512 (ISA §3.6.4: “up to 512 total VGPRs, 256 of each type… the number of each type is flexible”). This is not the separate physical ACC file of MI100/MI200; CDNA3 unified them and CDNA4 keeps it that way. In practice the compiler fills regular VGPRs first: most of the MXFP8 GEMM tiles I profiled run with zero AccVGPRs — the accumulator sits in regular VGPRs — and only the largest tiles spill into the accumulator pool. Either way it’s one budget, and that’s the fact Part 3 turns on. SGPRs — ~800 per SIMD. Scalar registers, allocated per wave. Rarely the binding limit. LDS — 160 KB, shared across all four SIMDs. One physical scratchpad per CU. It’s the cooperation mechanism for a workgroup, and on CDNA4 it’s 2.5× the size of CDNA3’s 64 KB. Here’s the distinction to burn in: the register files are per-SIMD; the LDS is per-CU. A wave that lands on SIMD 0 cannot touch SIMD 1’s registers, but every wave in a workgroup — wherever it lands — sees the same LDS. That single asymmetry is why the four occupancy limiters in Part 2 don’t share a denominator. Threads, lanes, and wavefronts The hardware doesn’t execute threads one at a time. It executes wavefronts: bundles of 64 threads that run in lockstep on a SIMD’s 64 lanes. AMD’s wavefront is exactly NVIDIA’s warp, just 64 wide instead of 32 — and “wave” and “wavefront” are the same thing. A 256-thread workgroup is therefore four waves, not 256 independent units of scheduling. A SIMD holds up to 8 wavefronts resident at once.

Resident means their registers are live and reserved; only one wave issues per cycle, and the scheduler hides latency by switching to a different ready wave whenever the current one stalls. Occupancy is just the ratio of resident waves to that maximum of 8 — counted per SIMD, capped at 32 per CU. Crucially, occupancy says nothing about whether those waves are doing useful work; it only counts how many are parked. Hold that thought for Part 3. If you think in CUDA: a Rosetta stone AMD’s vocabulary maps almost one-to-one onto NVIDIA’s. If your instincts are CUDA-shaped, read this first and the rest of the post translates itself: AMD (CDNA4)NVIDIANotesCompute Unit (CU)SM256 on MI355XSIMD — 4 per CUSM sub-partition — 4 per SM64 lanes eachWavefront — 64 threadsWarp — 32 threadsAMD waves are 2× wideVGPR / AccVGPRregistersone shared 512/lane fileLDS — 160 KB/CUshared memoryper-CU scratchpadMatrix Core / v_mfma_*Tensor Core / mmawaves resident per SIMD (≤8)warps resident per SMthe occupancy ratioworkgroup · s_barrierthread block · __syncthreads() The one place the analogy frays: NVIDIA has no separate “accumulator register” concept — on CDNA4 the AccVGPRs are simply part of the same register file, which is exactly the subtlety Parts 2 and 3 hinge on. The Matrix Core The reason any of this exists on an AI accelerator is the Matrix Core — the MFMA engine. CDNA4 overhauled it with native FP8, FP6, and FP4 support and roughly doubled per-CU matrix throughput versus CDNA3. The MI355X peaks at about 5 PFLOPs of MXFP8 and 10 PFLOPs of MXFP6/FP4 dense matrix throughput, with structured sparsity pushing FP4 past 20 PFLOPs.

The instruction you’ll meet in MXFP8 GEMM kernels is the scaled-MFMA family — v_mfma_scale_f32_16x16x128_f8f6f4 or its 32x32x64 sibling, depending on the tile — which folds per-block microscaling directly into the matrix op. The mechanical detail that connects back to occupancy: MFMA reads its operands from the VGPR file and accumulates back into registers (regular or accumulator VGPRs — the same 512-entry file). The matrix engine is fed by registers — nothing slower can keep it busy at peak. That’s exactly the tension the memory hierarchy makes concrete. The memory hierarchy From fastest and smallest to slowest and largest: the register file (VGPR/AccVGPR, per-SIMD) → LDS (160 KB/CU) → L1 → L2 (per XCD) → 256 MB Infinity Cache → 288 GB of HBM3E at 8 TB/s. Each step down trades bandwidth for capacity. The gap that matters is the very first one. Only the register file delivers operands fast enough to sustain the Matrix Core at peak; LDS, despite being on-chip and despite CDNA4’s generous 160 KB, is meaningfully slower and prone to bank conflicts. Keeping the hot accumulation register-resident — not in LDS — is what lets a kernel hit the matrix peak, and it’s the seed of the argument that closes the post. CDNA3 → CDNA4, in one box What actually changed for kernel authors moving from MI300X/MI325X to MI355X: LDS: 64 KB → 160 KB per CU. The single biggest occupancy-relevant change; LDS-bound tiles get dramatically more headroom. Matrix Core: ~2× per-CU throughput, plus native FP6/FP4 (FP6 runs at FP4 rates). VGPRs: unchanged at 512/SIMD. Despite a common assumption, CDNA4 did not double the vector register file. If you’ve heard otherwise, that’s the myth to drop before doing the math. With the resources and their boundaries in hand, we can finally count them.

That’s Part 2. Part 2 — Occupancy math Occupancy has a one-line definition: it’s the number of wavefronts resident on a SIMD divided by the maximum that SIMD can hold (8). Report it per SIMD or scale it to the CU (max 32 waves) — same ratio either way. The number you get is set by whichever resource runs out first, and there are four candidates. The limiters Each limiter has the same shape — a fixed hardware budget divided by what one unit of your kernel consumes, floored to a whole number — but they count different things: VGPRs hold every per-thread value live at once — loaded operands, loop variables, and the accumulator tile — with regular and accumulator registers drawing on one 512-entry file. Bigger tiles and deeper unrolling cost more. SGPRs hold values uniform across a wave — base addresses, loop bounds, predicates. Usually cheap. LDS holds the workgroup’s shared staging buffers: the tiles you cooperatively load before feeding the matrix core. Workgroup slots and barriers are spent once per resident workgroup, no matter how large or small it is. As formulas: VGPR limit: floor( 512 / total-VGPRs-per-lane ) -> waves / SIMD (total = regular + accumulator) SGPR limit: floor( ~800 / SGPRs-per-wave ) -> waves / SIMD LDS limit: floor( 160 KB / LDS-per-workgroup ) -> workgroups / CU WG limit: max workgroups resident / CU (hard cap + barrier slots) Your occupancy is the minimum of the four, then clamped by the hardware caps (8 waves/SIMD, 32 waves/CU). In practice VGPRs or LDS bind; SGPRs almost never do unless you have heavy scalar address math. The fourth limiter is workgroup-level. A CU can hold only a fixed number of resident workgroups, enforced in part by barrier resources — every workgroup of more than one wavefront needs a hardware barrier to implement __syncthreads / s_barrier, and a CU has a limited pool of them.