/learn · hipfire 0.4.1

How hipfire runs LLMs on AMD GPUs.

Each page explains one idea, shows how hipfire implements it (with links into the v0.4.1 source), and says what it measured in the 0.4.1 release on Strix Halo and the Radeon AI PRO R9700. Nothing here is projected or borrowed from another engine.

native MTP PeaceMaker session cache FWHT quantization sparse attention WMMA · MFMA

inside 0.4.1

Flash-Next · native MTP

Speculative decoding on RDNA

Draft, verify, accept. Greedy and sampled MTP with lossless rejection sampling, penalties on every verify row, n-gram takeovers priced in milliseconds, and an AR floor.

kernels · certification

PeaceMaker: certified kernels

hipfire's own kernel builder. Code objects written without a ROCm assembler, lifted back byte for byte, checked for wait, hazard and LDS bugs, and byte-identical to the hipcc kernels they replace.

prompt reuse

Session cache on a hybrid model

Reusing prefill when part of the model is a recurrence: snapshot chains at chunk boundaries, copy plus delta, cold-exact vs session-exact, and live continuation.

tokenizer

Tokenizer fidelity

One missing regex alternative, \s+(?!\S), changed how whitespace runs split. What it is, how 0.4.1 matches Hugging Face, and how that is tested.

quantization · FWHT

MagnumQuant and FWHT rotation

Why a Walsh-Hadamard rotation makes 4-bit weights behave, the V2 group layout, and how a symmetric grid lets Flash-Next's experts run on int4 WMMA.

attention · memory

QSA sparse attention and VMM KV

How Flash-Next attends to 2,048 selected tokens instead of the whole context, and how the VMM backend commits KV memory only as the context grows.

the silicon underneath

consumer / pro / APU

RDNA & WMMA

wave32, the 16×16×16 WMMA tile on gfx11 and gfx12, occupancy, the memory ladder, and which WMMA instructions 0.4.1 actually issues on Strix Halo and the R9700.

data center

CDNA & MFMA

wave64, MFMA tiles, HBM and the Infinity Fabric package, rocprofv3 counters, and where MFMA sits in hipfire today.

Measured figures on these pages come from the v0.4.1 release notes, with the hardware and protocol named next to each number. Hardware specifications cite AMD.

two GPU families, one engine

Same HIP source, two different machines. The silicon pages walk each row of this table.

	CDNA — Instinct	RDNA — Radeon / Ryzen AI
wavefront	64 threads	32 threads
matrix unit	`v_mfma_*`	`v_wmma_*` (RDNA3 and later)
memory	HBM	GDDR6, or shared LPDDR5X on Strix Halo
LLVM targets hipfire knows	`gfx908`, `gfx940–gfx942` (plus Vega `gfx906`)	`gfx1010` → `gfx1201`
hipfire kernels	portable wave64 path; MFMA through rocBLAS on CDNA3	architecture-tuned WMMA kernels on gfx11 / gfx12