VGPR pressure caps the fused RDNA4 prefill GEMM at nine waves

Yesterday’s post ended on a number that did not add up. Fusing the Q4_K dequant step into ZINC’s RDNA4 prefill GEMM cut weight DRAM traffic by about eight times, and a first-order roofline said that should be worth a 2.7x ceiling on the gate-and-up matmul. The measured prefill gain on Qwen3.5-9B on the RX 9070 XT was 1.95x, from 219 to roughly 430 tok/s. That is a good result, but it is not 2.7x, and the difference is not rounding. Something is stopping the fused kernel from reaching the compute roof it earned the right to sit on.

The something is occupancy. A kernel that is no longer waiting on memory bandwidth can still be waiting on memory latency, and the way a GPU hides latency is by having many independent waves in flight so it always has other work to run while one wave stalls on a load. Fusing dequant into the GEMM did not come for free in register terms. The kernel now has to unpack four-bit weights into registers and hold the WMMA accumulator tile there across the whole inner loop, and that pushed its register footprint high enough to starve the very latency hiding it needs.

This post is about the second tax, the one that shows up right after you fix the first. It is a smaller, older story than a bandwidth cliff, and it is the reason a bandwidth-optimal kernel can still leave a third of the card’s matrix throughput on the floor.

Why occupancy is the thing that pays for latency

Start with the mechanism, because the numbers only make sense once the mechanism is clear. A RDNA4 SIMD runs one wave of 32 lanes at a time, but it can hold many waves resident and switch between them cycle to cycle. When the active wave issues a load from VRAM and has to wait a few hundred cycles for the data, the SIMD does not stall. It runs a different resident wave. The more resident waves there are, the more independent work is available to fill those gaps, and the closer the SIMD gets to being busy every cycle. This is occupancy, and it is the single biggest lever on whether a compute-bound kernel actually reaches its compute roof.

The catch is that resident waves are not free. Every wave needs its own registers, and they all come out of one fixed register file per SIMD. If each wave asks for more registers, fewer waves fit, and occupancy falls. This is not new to RDNA4. The classic version of the problem was written up years ago by Sebastian Aaltonen, who found that on older AMD hardware dropping a shader from 40 to 32 registers per thread doubled the number of resident thread groups and delivered a 50 percent speedup with no other change. Same silicon, same arithmetic, just more waves to hide latency behind. The register file is the budget, and occupancy is what you buy with it.

RDNA4 gives us exact numbers to work with. The RX 9070 XT and R9700 carry a 192 KB register file per SIMD. In wave32 mode each vector register is 1024 bits wide, so 192 KB holds 1536 vector general-purpose registers, and each SIMD has sixteen wave slots. Divide 1536 registers across sixteen waves and you get 96 registers per wave. That is the line: a kernel that uses 96 or fewer VGPRs runs all sixteen waves at full occupancy, and every register past 96 costs you waves.

What the fused GEMM actually spends registers on

The staged kernel from the last post was register-cheap precisely because it was two dumb passes. The first pass read Q4_K and wrote fp16 scratch. The second pass was a plain fp16 GEMM that read the scratch back. Neither pass held much live state, so the kernel sat around 72 VGPRs and ran all sixteen waves. It just happened to be memory-bound, so the spare occupancy bought nothing.

Fusing the two passes changes the register picture entirely. Now a single kernel has to, inside one loop iteration, load a Q4_K block, unpack its packed four-bit weights and its scale and min into usable values, load the matching activation fragment, and feed both into the WMMA matrix intrinsic, whose accumulator sits in registers and is read and rewritten on every step. The RDNA4 WMMA layout is deliberately lean, with each lane holding eight elements of a 16x16 tile, and AMD notes the RDNA4 layout was simplified specifically to reduce VGPR pressure versus RDNA3. Even so, a prefill tile accumulates several of those 16x16 fragments at once so it can reuse each weight block across many tokens, and the accumulators are all live for the duration of the K loop. Add the unpack temporaries and the per-block scales, and Radeon GPU Analyzer put the first honest version of the fused kernel at roughly 168 VGPRs per wave.

Read the staircase and 168 VGPRs is not a gentle penalty. Because these parts allocate registers in 24-register blocks, 168 rounds to a seven-block allocation, and 1536 registers divided into seven-block chunks leaves room for nine waves. Nine of sixteen slots is 56 percent occupancy, and it is almost exactly the picture Chips and Cheese caught in a completely different workload, where a RDNA4 raytracing shader was pinned to nine of sixteen threads by VGPR usage. Nine resident waves is not enough independent work to keep the SIMD busy while weight loads and WMMA results are in flight, so the fused kernel spends part of every stall idle. That idle time is the gap between the roofline’s 2.7x and the measured 1.95x.

The recoverable registers versus the stubborn ones

Not all 168 of those registers deserve to be there, and telling the two kinds apart is the whole job. The stubborn registers are the WMMA accumulators. The C and D matrices have to live in VGPRs across the entire K loop by definition, and a wide prefill tile deliberately holds several fragments so it can amortize each weight load over many tokens. Shrinking that tile would cut registers, but it would also cut the reuse that made the fused kernel worth building, so those registers are load-bearing.

The recoverable registers are everything that does not actually vary per lane. A Q4_K block’s scale and min are the same for all 32 lanes reading that block, which makes them exactly the kind of wave-invariant data that belongs in the Local Data Share or in scalar registers rather than in a vector register replicated across every lane. Aaltonen’s old advice applies unchanged: move group-shared values out of VGPRs and the register budget frees up for occupancy. Packing the fp16 scales two to a register and tightening the accumulator live ranges by tiling the K loop into shorter segments does the rest. Together those take the kernel from about 168 VGPRs to about 120, which is a five-block allocation and twelve resident waves.

Qwen3.5-9B prefill GEMM, RX 9070 XT	VGPR/wave	waves	occupancy	prefill
staged fp16 (memory-bound)	~72	16	100%	219 tok/s
fused, first cut	~168	9	56%	~430 tok/s
fused, scales to LDS + tight live ranges	~120	12	75%	~500 tok/s
llama.cpp pp512 reference	—	—	—	973 tok/s

The jump from 9 to 12 waves is worth roughly another 15 percent of prefill on the 9B, taking it from about 430 to about 500 tok/s in the profiling runs. That is a smaller win than the fusion itself, and it should be, because latency hiding has diminishing returns: the step from 56 to 75 percent occupancy closes most of the stall gap, and the last few waves up to full occupancy would buy less and cost the accumulator tile. The honest read of the table is that the fused kernel traded a bandwidth problem for an occupancy problem, and the occupancy problem is the cheaper of the two to chip away at.

The hardware already has a better answer

There is a cleaner fix than counting registers by hand, and it is sitting in the RDNA4 silicon unused on this path. RDNA4 introduced dynamic VGPR allocation, a mode where a wave launches with a minimal register allocation and asks for more at runtime with an s_alloc_vgpr instruction, then frees them again when it leaves the hungry code. The occupancy question inverts: the driver sets how many waves run per SIMD directly, and a wave only holds its peak register count during the short window it actually needs it. A GEMM built this way could keep sixteen low-register waves resident and spike up to the accumulator-heavy allocation only inside the multiply loop, getting both the occupancy and the wide tile.

The reason ZINC cannot use it yet is entirely about plumbing, not silicon. Dynamic VGPR mode is restricted to wave32 compute shaders, it is gated behind a chip-wide control register the driver has to set up, and allocation requests can fail and force a wave to busy-wait, which brings its own deadlock-avoidance machinery. So far it has only been seen in AMD’s own indirect-mode raytracing shaders, and Chips and Cheese notes the obvious next step directly: generic compute could benefit too, once the feature is exposed through the toolchains. It is not in the Vulkan compute path today. Nvidia shipped its own version of this idea, setmaxnreg, back in Hopper, and some of its GEMM libraries already lean on it, which is a fair sign of where AMD’s compute stack is heading.

So the near-term move is the boring one, and the long-term move is waiting on a driver. ZINC prefills Qwen3.5-9B on the 9070 XT at roughly 500 tok/s once the wave-invariant scales come out of the vector registers, against llama.cpp’s 973, and the gap is down to about 1.9x. The lesson that carries past this kernel is that the two bottlenecks are a matched pair. First you stop moving bytes you do not need to move, and the moment you do, the registers you spent to avoid moving them become the thing that caps you. On a machine where the register file is the real budget, a fused kernel is a bet that the occupancy you give up is cheaper than the bandwidth you save, and on RDNA4 the arithmetic of that bet is a staircase with a step every 24 registers.