RDNA & WMMA — A Distil

01 The RDNA family at a glance

hipfire targets the entire RDNA family with one Rust binary — the matrix-path feature gating is the load-bearing piece of the dispatch layer.

Two structural breaks worth memorizing:

1. RDNA1/2 → RDNA3 is where the matrix unit appears. Before gfx1100, the fastest f16-mac path is v_dot2_f32_f16 (a 2-wide packed-FMA) or v_dot4_i32_i8 (dp4a). After gfx1100, you have v_wmma_* with a 16×16 accumulator per wave.
2. RDNA3 → RDNA4 keeps the wave32 WMMA shape and the 16×16×16 tile, but changes the layout: A/B packing halves from half16_t to half8_t per lane, the builtin gains a _w32_gfx12 suffix, and the C-map moves. LDS stays ≤64 KB per workgroup (≤128 KB per WGP) on both gens. Most kernels still need their own .gfx12.hip variant — for the incompatible builtin/layout, not a wider tile.

02 CDNA vs RDNA — the divergence summary

Why this matters for code: the same .hip file compiles for either family — but a CDNA kernel ported to RDNA hits three landmines: wavefront width (64 → 32 halves the lanes seen by __shfl_*), matrix intrinsic (v_mfma_* doesn't exist on RDNA — you must call v_wmma_*), and memory-tier sizing (96 MB Infinity Cache on the 7900 XTX changes your tiling math vs HBM-direct).

03 WMMA — one instruction, one 16×16 tile

The WMMA family is RDNA3+'s answer to CDNA's MFMA. Same shape — a fused D = A × B + C over a 16×16 tile — but reachable from wave32 (half the lanes) and from a fundamentally consumer-silicon caches.

The canonical RDNA3 builtin (used in hipfire's gemm_f16_wmma.hip):

// WMMA layout (gfx11 RDNA3): __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c)
// computes  D = A x B + C over one 16x16 tile (ISA-defined lane layout).
//
// A is half16_t (16 fp16 values per lane)
// B is half16_t
// C / D are float8_t (8 fp32 accumulators per lane on wave32)
// (gfx12 RDNA4 keeps 16x16x16 but packs A/B as half8_t per lane, uses a
//  _w32_gfx12 builtin, and remaps C — see “What gfx12 WMMA changes” below.)

typedef _Float16 __attribute__((ext_vector_type(16))) half16_t;
typedef float    __attribute__((ext_vector_type(8)))  float8_t;

__launch_bounds__(32, 2)
extern "C" __global__ void gemm_f16_wmma(
    const _Float16* __restrict__ W, const float* __restrict__ X,
    float*          __restrict__ Y, int M, int K, int N)
{
    // Grid: [ceil(M/16), ceil(N/16)],  Block: [32]
    float8_t acc = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
    for (int k0 = 0; k0 < K; k0 += 16) {
        half16_t a_reg, b_reg;
        // ...load 16x16 tiles into a_reg and b_reg...
        acc = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a_reg, b_reg, acc);
    }
    // ...store acc back to Y...
}

What WMMA gives you (RDNA3, gfx1100 / 1151):

• One dense v_wmma_f32_16x16x16_f16 issue performs 16×16×16 = 4096 FMA (8192 FLOPs) over the tile. Throughput is a separate model: GPUOpen lists 512 FLOPS/clk/CU fp16 dense on the 7900 XTX, i.e. ≈ 123 TFLOPs theoretical at boost — a nameplate ceiling, not a measured hipfire number. The packed-dot2 path sits on a much lower model roof, qualitatively far below WMMA.
• Wave32 native — _w32 suffix on the builtin. _w64 WMMA builtins also exist on gfx11, but hipfire's primary path issues wave32.
• The accumulator (float8_t on wave32, float16_t on wave64) lives in the VGPR file, so a fused-GEMM kernel typically lands at 80–120 VGPRs per thread (lane) even without spills — see § 06 for what that does to occupancy.

What gfx12 WMMA (RDNA4, gfx1200+) changes:

• Same 16×16×16 tile, new layout: A/B pack as half8_t (8 fp16 per lane; K splits by lane group), the builtin becomes __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12, and the C-map moves. Most existing gfx11 kernels need their own .gfx12.hip override for the incompatible builtin/layout — not a wider tile. RDNA4 WMMA also adds FP8 mixes; IU4 remains a WMMA dtype, not MFMA.
• gfx1100 and gfx1201 both model at 1536 VGPRs/SIMD (see § 06), so there is no register-file growth story for the 7900 XTX class here — the gfx12 win is layout/throughput, not occupancy.

04 The roofline — three matrix paths, one chip

Because RDNA3 ships both v_wmma_* and v_dot2_f32_f16 (and the scalar fallback), it has three different compute ceilings on the same chip. Which one bounds you is decided by your kernel's matrix path, not by the GPU's nameplate TFLOPs.

• Decode GEMV (batch=1) lives far left on the AI axis. Even with WMMA available, you can't feed a 16×16 matrix unit from a length-1 batch — you fall back to dot2/scalar and end up memory-bound on GDDR or system RAM. Decode is a memory-bandwidth problem on RDNA for the same reason it is on CDNA.
• Batched prefill (B=8) sits mid-AI (AI=32, still memory-bound in this model at 0.960×32≈31 TF). The WMMA kernel path is selected but the memory slope still binds.
• Wide prefill (B=32+) is where WMMA wins. At AI≥128 the intensity crosses the ridge (0.960×128≈123) and the WMMA roof is the only ceiling left.

This is the regime split the fixture-scoped snapshots in § 07 illustrate: batch size and arithmetic intensity decide which model roof binds, not the card's nameplate alone. Nothing on this chart is a measured hipfire throughput.

05 Why CDNA “throws away” graphics and RDNA doesn't

Mirror of slide 40 from the CDNA deck, but inverted:

The result for an LLM workload: an RDNA3 board costs ≈ 1/10 of an MI300X. The matrix-unit throughput per dollar is similar if your kernel actually saturates WMMA. The headline divergence is that one of them runs a display.

06 The wave32 occupancy table

WMMA-fused GEMMs are register-hungry. The 8-float accumulator alone costs 8 VGPRs per thread (lane); with tile-staging and A/B register tiles you quickly land in the 80–120 VGPR range. That has to fit the wave32 occupancy curve — drag the slider:

VGPR-capacity model for gfx1100 and gfx1201: 1536 VGPRs per SIMD, wave32 allocation granule 24, cap 16 waves/SIMD — min(16, floor(1536 / (ceil(vgprs/24) × 24))). LDS, thread/workgroup, and barrier limits are explicitly excluded.

Two rules of thumb that show up repeatedly in the hipfire kernel notes:

• WMMA / MFMA kernels run hot on VGPRs. Common allocation is 80–120 VGPRs per thread with zero spills.
• High theoretical occupancy + low VALUBusy = memory-bound. More occupancy won't help; you need more in-flight HBM/GDDR transactions per wave (multi-quad interleave, half-wave splits, prefetch).
• .private_segment_fixed_size: 0 is the reliable “no spills” indicator from the AMDGPU note section — the vgpr_spill_count field is sometimes elided by the toolchain when zero.

Reading those numbers out of a compiled .hsaco (from the gfx-kernel-metadata skill):

ARCH=gfx1100
# 1. Unbundle the offload container into a real ELF
/opt/rocm/llvm/bin/clang-offload-bundler --type=o --unbundle \
    --input=kernel.hsaco --output=/tmp/kernel.elf \
    --targets=hipv4-amdgcn-amd-amdhsa--$ARCH
# 2. Read AMDGPU notes
/opt/rocm/llvm/bin/llvm-readelf --notes /tmp/kernel.elf
# → .vgpr_count, .group_segment_fixed_size, .private_segment_fixed_size,
#   .wavefront_size: 32 (RDNA) / 64 (CDNA)

07 Measured snapshots — fixture-scoped beta figures

Shared fixtures from src/data/performance.ts (beta branch). Optimization on the hipfire beta branch is continuous. Every figure below is fixture-scoped, and a date is shown only when its source carries a measurement date. These source-published snapshots are not immutable or lasting claims. Numbers update frequently. /docs/benchmarks is the build-time live ledger.

Rows are heterogeneous fixtures. Compare cells only when model, quant/mode, prompt, method, and backend match. No cross-hardware ratio below is presented as an isolated WMMA gain — rows differ in CU count, memory tier, clocks, and backend, not just matrix path.

Qwen3.6 35B-A3B MQ4R · AR · Q8 KV · Ordinary autoregressive decode, single GPU; no MTP, DFlash, reduced-output bench, or manual clock pinning. TG128 = three-run medians. Multi-turn columns from clean eight-turn serving runs. Source: beta README · Qwen3.6 35B-A3B MQ4R (README fixture, undated).

Qwen3.8-27B MQ4V2 product ladder · hiptrx · gfx1201 · Radeon AI PRO R9700 · dated 2026-08-20 · Product-ladder checkpoint on the hiptrx gfx1201 / Radeon AI PRO R9700 fixture. AR decode, prefill, DFlash decode, mean accepted draft length (τ), and bits-per-weight. Source: beta · 2026-08-20 Qwen3.8 MQ-V2 product ladder.

Featured checkpoint figure: Base-tier DFlash 263.3 tok/s (Qwen3.8-27B MQ4V2 Base · gfx1201 R9700 · 2026-08-20). The ladder's AR/prefill/DFlash columns are separate protocols on one dated fixture — the spread across tiers motivates batch/AI regime thinking from § 04, not an isolated WMMA speedup.

08 Memory hierarchy — RDNA edition

Same layered story as CDNA, with consumer-silicon numbers. The absolute bandwidths are smaller across the board, but Infinity Cache (96 MB on the 7900 XTX; 64 MB on the R9700) is a uniquely RDNA tier — it sits between L2 and DRAM and dramatically widens the effective LLM-friendly working set.

 tier             scope         cap          BW          lat
 ────             ─────         ───          ──          ───
 VGPRs            per CU        8 KB/wave    ~100 TB/s   0 c
 LDS              per WGP       ≤64 KB/WG   ~10 TB/s    ~20 c
 L1               per CU        32 KB        ~5 TB/s     ~50 c
 L2               per array     6 MB         ~3 TB/s     ~150 c
 Infinity Cache   device-wide   64–96 MB     ~17 TB/s    ~200 c   ← RDNA-only
 GDDR / sys RAM   device-wide   8–96+ GB     0.5–1.0     ~400 c
        

The standard CDNA-deck rule still holds: every step down the ladder is at least an order of magnitude slower, so a kernel that touches the same byte twice should get the second touch from LDS/L1, not DRAM. The RDNA twist is that Infinity Cache makes the device-wide working set ~96 MB instead of ~6 MB of L2 on the 7900 XTX — a short-context KV slice can live entirely in IC, which helps gfx1100 decode punch above its raw GDDR bandwidth tier.

09 Dispatch — fast paths first, baseline last

rdna-compute::dispatch is the kernel-selection hot path in hipfire. Every GEMM / GEMV / norm / fused op routes through here. The shape is always the same: arch-feature predicates, fast paths first, baseline (fp16-packed, RDNA1) last.

// Simplified WMMA-capability cascade for gemm_qkv_hfq4g256.
// Production may select an MMQ arm first when batch/alignment allow —
// see crates/rdna-compute/src/gemm.rs. Predicate names match arch_caps.rs.
pub fn gemm_qkv_hfq4g256(&self, ...) -> HipResult<()> {
    // if self.arch_caps.has_hfq4_mmq() && dims aligned { mmq(...) }  // may win first
    if self.arch_caps.has_wmma_w32_gfx12() {          // RDNA4: gfx1200/1201
        return self.gemm_qkv_hfq4g256_wmma_gfx12(...);
    }
    if self.arch_caps.has_wmma_w32() {                // RDNA3: gfx11xx, gfx1150/51
        return self.gemm_qkv_hfq4g256_wmma(...);
    }
    if self.arch_caps.has_dot2_f32_f16() {            // RDNA1.1+/RDNA2; false on gfx1010
        return self.gemm_qkv_hfq4g256_dot2(...);
    }
    self.gemm_qkv_hfq4g256_fp16(...)                  // gfx1010/1013 packed-fp16 baseline
}

Two contributor rules from the architecture doc:

1. Predicates are arch-feature checks, not inline arch.starts_with(...) chains. New silicon = update one predicate, not 40 call sites.
2. No unreachable branches. When a new arch absorbs a check that was matched by an older || starts_with("gfxN") clause, drop the redundant clause in the same diff.

The per-arch kernel-file naming convention follows the dispatch:

kernels/src/gemm_qkv_hfq4g256_wmma.hip            # default WMMA (RDNA3)
kernels/src/gemm_qkv_hfq4g256_wmma.gfx12.hip      # gfx12xx override (gfx12 WMMA layout)
kernels/src/gemv_hfq4g256.gfx1030.v4.hip          # chip-specific versioned
kernels/src/gemv_hfq6g256_residual_wave64.hip     # wave64 (CDNA fall-back)

scripts/compile-kernels.sh resolves chip → family → default in that order; .gfx12.hip covers both gfx1200 and gfx1201 with one file.

10 Historical hipfire snapshots (superseded)

The figures in this section are retained for architectural context, not as current throughput claims. For dated, fixture-scoped beta figures see § 07 (src/data/performance.ts).

Lifecycle: historical. These fixture-bound observations do not carry the complete measurement date and binary/model identity required for a current measured claim. See beta docs/BENCHMARKS.md.

Historical 7900 XTX (gfx1100) snapshot, then-default asym3 KV and FlashAttention auto (undated multiplier column removed):

The historical beta table also retained superseded-method DFlash observations: up to 218.6 tok/s on 27B HumanEval/53 and 372.9 tok/s on 9B HumanEval/0. They used asym3 KV and max_tokens=120; results were genre-conditional and are not current baselines. See the historical source table.

Historical 27B AR snapshot across the RDNA family — same model, prompt sweep, asym3 KV, and --no-chatml:

Historical HumanEval N=33 fixture. The source does not carry the full per-row measurement date and binary/model identity now required for a measured claim; the observed 7900 XTX-to-Strix spread was consistent with their different memory-bandwidth tiers under that fixture.

Retained historical APU snapshot — Strix Halo (gfx1151), 27B MQ4 with the 27B-DFlash sidecar and the then-canonical code prompt:

Historical row: 3-cell median, τ=10.0, accept rate 0.67. This deck predates the current complete identity/date manifest requirement, so the comparison remains educational context rather than a current performance claim.

That historical snapshot illustrates the verify-batch opportunity on a system-RAM APU; current performance belongs in a dated, protocol-complete checkpoint.

11 Summary — what to remember

• RDNA wave width is 32, not 64. __shfl_*, occupancy math, WMMA _w32 builtins all key off this.
• WMMA is the RDNA3+ matrix unit. Uses v_wmma_* intrinsics over 16×16×16 tiles on both RDNA3 (gfx1100/1101/1102/1150/1151, half16 A/B) and RDNA4 (gfx1200/1201, half8 gfx12 layout). gfx1010 has no packed DOT and no matrix unit; RDNA1.1/RDNA2 add v_dot4_i32_i8 / v_dot2_f32_f16 packed DOT.
• Beta snapshots are fixture-scoped. The Qwen3.6 MQ4R AR rows and the dated (2026-08-20) Qwen3.8 MQ4V2 ladder in § 07 give per-card, per-protocol context; rows are heterogeneous, so no cross-hardware ratio is claimed as an isolated WMMA gain.
• Historical heterogeneous-cluster snapshots showed large gains. They motivate tier specialization, but are not current end-to-end baselines.
• Dispatch is feature-gated (has_wmma_w32_gfx12 / has_wmma_w32 / has_dot2_f32_f16 in arch_caps.rs), not arch-string-matched. New silicon costs one predicate update, not 40 call sites. The § 09 tree is simplified: production may select an MMQ arm first.
• WMMA fused-GEMMs cost ~80–120 VGPRs per thread. At ≤96 VGPRs the VGPR-capacity model holds full 16/16 waves/SIMD on both gfx1100 and gfx1201 (1536 pool, granule 24); granule-24 cliffs start at 97. LDS/workgroup limits are excluded from that model.
• Infinity Cache is a uniquely RDNA tier, sitting between L2 and DRAM at 96 MB on the 7900 XTX (64 MB on the R9700). It widens the effective LLM working set enough that a 24 GB GDDR6 consumer card punches above its raw bandwidth tier at short context.

12 Primary sources

• AMD GPUOpen — WMMA on RDNA3: 16×16-only tiles, …_16x16x16_… intrinsics, 512 FLOPS/clk/CU fp16 dense.
• AMD GPUOpen — Using matrix core on RDNA4: 16×16 WMMA with gfx12 packing and _w32_gfx12 builtins.
• AMD GPUOpen — Occupancy explained: 7900 XTX VGPR pool, wave slots, allocation granularity.
• AMD — RDNA3 Shader ISA and RDNA4 Instruction Set Architecture: VGPR allocation blocks, LDS sizing, WMMA layout.
• AMD — Radeon RX 7900 XTX and RX 7900 series overview: 24 GB GDDR6, 96 MB Infinity Cache, Navi 31 = gfx1100.
• hipfire beta — crates/rdna-compute/src/arch_caps.rs and crates/rdna-compute/src/gemm.rs: live has_wmma_w32* / has_dot2_f32_f16 predicates and the MMQ-aware dispatch cascade.
• hipfire beta — 2026-08-20 Qwen3.8 MQ-V2 product ladder and beta README fixtures via src/data/performance.ts: all § 07 measured figures.