RDNA & WMMA

01 The RDNA family at a glance

hipfire targets the whole RDNA family with one Rust binary; feature predicates over the matrix path are the load-bearing piece of its dispatch layer.

From AMD's ROCm GPU hardware specification tables (VGPR file per WGP ÷ 4 SIMDs ÷ 128 bytes per wave32 VGPR). *RDNA1 value from the table's gfx1012 entry, the same generation. Every RDNA part listed has 128 KiB of LDS per WGP; a single workgroup can use up to 64 KiB.

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 next to the A/B fragments and any staging, so WMMA GEMMs are register-hungry. A real 0.4.1 example: the gfx1201 HC-down tile uses 178 VGPRs per lane with no scratch. § 06 shows 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 inputs (v_wmma_f32_16x16x16_fp8_fp8) and a 4-bit integer form with twice the K per issue (v_wmma_i32_16x16x32_iu4), which 0.4.1 issues for Flash-Next's expert GEMMs on the R9700 (§ 07).
• gfx1100, gfx1151 and gfx1201 all have 1536 VGPRs per SIMD (see § 01 and § 06), so the gfx12 change is layout and data types, not register capacity.

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 that 0.4.1's routing follows (§ 07): batch size and arithmetic intensity decide which 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

The mirror image of the CDNA page's transistor-budget section:

The result for an LLM workload: fewer CUs and much less memory bandwidth than an Instinct part, but the same kind of 16×16 matrix tile per wave, on a card that sits in a workstation or a laptop-class APU. Decode is bandwidth-bound on both; what RDNA changes is how much bandwidth and memory you have to work with.

06 The wave32 occupancy table

WMMA GEMMs are register-hungry: the 8-float accumulator alone costs 8 VGPRs per lane, and A/B fragments, rings and staging come on top. The count 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.

A worked 0.4.1 example and two rules of thumb:

• The gfx1201 HC-down tile (qwen4_wmma_lds_160_80_32_80_k32_p_s4, five 32×80 waves) uses 178 VGPRs: by the model above, 8 waves per SIMD. The bigger lever was LDS: the kernel it replaced needed 24,576 B per workgroup, which capped a WGP at five workgroups; the new tile needs 15,360 B, so eight fit in a WGP's 128 KiB. Same accumulation order, byte-identical output.
• High theoretical occupancy + low VALUBusy = memory-bound. More occupancy won't help; you need more in-flight memory transactions per wave (interleave, 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 How hipfire 0.4.1 uses WMMA

The § 04 regime split shows up directly in 0.4.1's routing: decode and few-row verify stay on GEMV-style kernels, and wide prefill chunks go to WMMA. For Flash-Next on the two release cards the line is drawn at 512 rows per prefill chunk; the chunk itself is 8,192 rows on Strix Halo and 4,096 on the R9700.

The IU4 route is the most RDNA-specific of these: once the weights sit on a symmetric 4-bit grid and the activations are quantized to 4 bits, an expert GEMM is an integer WMMA with one scale at the end (MagnumQuant § 04). On gfx12 the IU4 instruction takes K = 32 per issue instead of 16. Those expert kernels, the gathered QSA attention and the MQ6 trunk GEMMs are built with PeaceMaker.

Outside Flash-Next, the dense HFQ4 Q/K/V GEMM on RDNA3 has one more WMMA arm: from 128 rows, an MMQ path quantizes activations to Q8_1 and runs integer WMMA, ahead of the F16 WMMA kernel (§ 09).

What it buys you

08 Memory hierarchy — RDNA edition

Same layered story as CDNA, with consumer-silicon capacities. For decode the tier that matters is the bottom one: every active weight is read once per token, so the rate at which weights can be streamed sets the decode ceiling.

Cache and register capacities from AMD's ROCm GPU specification tables; DRAM size and peak bandwidth from AMD's product pages.

The standard rule still holds: every step down the ladder is much slower, so a kernel that touches the same byte twice should get the second touch from registers or LDS, not DRAM. That is what the PM MQ6 twin on the PeaceMaker page does: it stages activations in LDS and keeps B fragments in an 8-deep register ring. A large MoE like Flash-Next (a 125 GB weights file) does not fit in any cache, or even in a discrete card's VRAM: every prefill chunk re-streams the routed experts, and on the R9700 part of the experts live in host memory.

09 Dispatch — fast paths first, baseline last

rdna-compute's Gpu methods are the kernel-selection hot path in hipfire: GEMM, GEMV, attention, norm and MoE calls route through them. The shape is always the same: arch-feature predicates from arch_caps, fast paths first, baseline last. Here is the real prefill (more than one row) cascade of the dense HFQ4 Q/K/V GEMM in 0.4.1:

// gemm_qkv_hfq4g256, prefill (rows > 1), trimmed from crates/rdna-compute/src/gemm.rs.
// The CDNA3 rocBLAS and gfx906 wave64 arms that come first are covered on the CDNA page.
if self.arch_caps.should_use_mmq(rows) {            // gfx11: rows >= 128 (gfx906: >= 8)
    return mmq_set_prequant(...);                   // Q8_1 activations, integer WMMA
}
if self.arch_caps.has_hfq4_mmq() && m % 128 == 0 {  // RDNA1.1 / RDNA2 (sdot4)
    return self.gemm_qkv_hfq4g256_mmq(...);
}
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 / 3.5
    return self.gemm_qkv_hfq4g256_wmma(...);
}
if self.arch_caps.has_dot2_f32_f16() {              // gfx1011/1012/1030-1032
    return self.gemm_qkv_hfq4g256_dot2(...);
}
self.gemm_qkv_hfq4g256_fp16(...)                    // gfx1010/1013: v_pk_fma_f16

Two rules from hipfire's 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)

Which source a module uses is decided in Rust (kernels.rs includes the .gfx12.hip override where one exists) and recorded in an exact source/flags registry; scripts/compile-kernels.sh builds the per-arch kernel packs from that registry, and the runtime verifies each installed object's SHA-256 before loading it.

10 Summary — what to remember

• RDNA wave width is 32, not 64. __shfl_*, occupancy math and the _w32 WMMA builtins all key off this.
• WMMA is the RDNA3+ matrix unit: v_wmma_* over 16×16×16 tiles on RDNA3 / 3.5 (half16 A/B) and RDNA4 (half8 gfx12 layout, plus FP8 inputs and a K = 32 IU4 form). gfx1010 has no packed DOT and no matrix unit; RDNA1.1 / RDNA2 add v_dot4_i32_i8 / v_dot2_f32_f16.
• 0.4.1 draws the WMMA line at 512-row prefill chunks for Flash-Next: IU4 experts with A4 activations, gathered F16 QSA attention, MQ6 trunk projections and gfx1201 HC GEMMs. Decode and verify stay on GEMV-style kernels.
• Dispatch is feature-gated (should_use_mmq, has_hfq4_mmq, has_wmma_w32_gfx12, has_wmma_w32, has_dot2_f32_f16 in arch_caps.rs), not arch-string-matched.
• WMMA kernels are register- and LDS-hungry. gfx1100, gfx1151 and gfx1201 all have 1536 VGPRs per SIMD; the 0.4.1 gfx1201 HC-down tile (178 VGPRs, 15,360 B LDS) is a worked example of trading tile shape for occupancy.
• Infinity Cache sits between L2 and DRAM: 96 MiB on the 7900 XTX, 64 MiB on the R9700, 32 MiB on Strix Halo.

11 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 — ROCm GPU hardware specifications: CU counts, LDS, L2, Infinity Cache and VGPR file sizes in § 01 and § 08.
• AMD — Radeon RX 7900 XTX, Radeon AI PRO R9700, Ryzen AI Max+ 395: memory size and bandwidth.
• hipfire v0.4.1 — crates/rdna-compute/src/arch_caps.rs and crates/rdna-compute/src/gemm.rs: predicates and the dispatch cascade.
• hipfire v0.4.1 — CHANGELOG (Flash-Next WMMA routes, the HC-down tile) and the release notes (all measured figures on this page).