Status: SHIPPING — 4-kernel architecture, 32 t/s decode flat 512→16k, 68% VRAM savings.
A complete INT8 DOT4 (sudot4) FlashAttention kernel system for llama.cpp HIP on RDNA3. Compresses the QK dot product from 256 scalar multiplies to 64 packed v_dot4_i32_iu8 instructions. Now includes persistent packed16 K cache (no f16 K allocation), split-K decode kernels, and auto-routing between prefill/verify and short/long decode.
One sentence: f16-V-equivalent decode speed with q4_0 V cache, 68% less VRAM, scalable to 128k.
Key result:
- Baseline q4_0 V + standard FA = 13 t/s decode
- Our packed16 K + q4_0 V + split-K = 32 t/s decode (matching f16 V baseline at 31 t/s)
- Attention is <9% of decode time — we hit the model's matmul/FFN floor
┌──────────────┬───────────────┬────────────────────┬────────────────────┐
│ Workload │ Condition │ Kernel │ Performance │
├──────────────┼───────────────┼────────────────────┼────────────────────┤
│ Prefill │ nq > 1 │ v4 recthist │ 401–746 t/s │
│ MTP verify │ nq > 1 │ v4 recthist │ same │
│ Small decode │ nq = 1, │ BN64 decode │ 32 t/s, 0.095ms/call│
│ │ nk < 2048 │ │ (1.2% of decode) │
│ Long decode │ nq = 1, │ split-K stage1+2 │ 32 t/s, 0.68ms/call │
│ │ nk ≥ 2048 │ │ (8.7% of decode) │
└──────────────┴───────────────┴────────────────────┴────────────────────┘
- v4: Single-launch recthist prefill, no prefix/tail split, causal-tail only
- BN64: Per-head single-CTA decode, BN=64 outer tile, BN8 V-subtile
- split-K: Stage1 distributes K rows across CTAs (512-row splits), stage2 merges via online softmax
- Route table: Committed in
fattn-dot4-q8k-decode.cuhas the canonical contract
| Config | Prefill | Decode (15k ctx) | VRAM @ 32k |
|---|---|---|---|
| Baseline q4_0 V (standard FA) | 117 t/s | 13 t/s | ~960 MiB |
| Baseline f16 V (standard FA) | 845 t/s | 29 t/s | ~1568 MiB |
| Packed16 + q4_0 V (our system) | 706 t/s | 23 t/s | ~832 MiB |
- +77% over baseline q4_0 V (13→23 t/s)
- 79% of f16 V ceiling (23/29)
- Pure decode at short ctx (BN64, ~200 rows): 30 t/s
- Long ctx gap (30→23): MTP model draft+verify overhead, not our attention kernels
- Attention timing: 0.68ms/call at 3.3k rows = 8.7% of decode budget
- Full canonical env flags:
GGML_CUDA_ROCM_Q8K_DOT4_PACKED16_K_CACHE=1 GGML_CUDA_ROCM_EXPERIMENTAL_UNSAFE=1 GGML_CUDA_ROCM_Q8K_DOT4_KQ=1 GGML_CUDA_FA_ROUTE_REQUIRE=rocm_q8k_dot4_kq GGML_CUDA_ROCM_Q8K_DOT4_KQ_VARIANT=blockfa_recthist_v4_single GGML_CUDA_ROCM_Q8K_DOT4_KQ_FULL_FA=1 GGML_CUDA_ROCM_Q8K_DOT4_BLOCKFA_ASSUME_CAUSAL=1 GGML_CUDA_ROCM_QUANT_PREFILL_F16=1 GGML_CUDA_ROCM_Q8K_DOT4_DECODE_BN=64 LLAMA_MTP_PREFILL_CHUNK=1024 LLAMA_MTP_PREFILL_FORCE_MMQ=1
| Prompt tokens | v4 recthist (t/s) |
|---|---|
| 512 | 401 |
| 1024 | 568 |
| 2048 | 689 |
| 4096 | 723 |
| 8192 | 746 |
| 15000 | 709 |
Peaks at ~750 t/s at 8k. Dips to 709 at 15k (quadratic attention cost).
| Config | PPL |
|---|---|
| Direct packed16 + f16 V | 1.0128 |
| Direct packed16 + q4_0 V | 1.0237 |
| q8_0 shadow K + q4_0 V | 3.0373 |
No rotation, no shadow K needed. The rotation domain mismatch fix (gate attn_rot_v on attn_rot_k) enables direct packed16 with q4_0 V at PPL 1.0237 — negligible quality cost for massive VRAM savings.
The key discovery: single-CTA decode (BN64) drops to 24 t/s at 16k because one CTA/head processes all 3,328 K rows serially. Split-K distributes rows across CTAs:
BN64 at short ctx: 30 t/s (pure decode, ~200 rows, no MTP overhead)
split-K at 15k ctx: 24 t/s (split-K + MTP draft/verify, still 2x baseline)
Baseline q4_0 V: 13 t/s (standard FA, same model, same flags)
Baseline f16 V: 29 t/s (ceiling — matmul/FFN/weight-dequant)
Attention timing breakdown (per head group):
- BN64 at nk=256: 0.095 ms (1.2% of decode budget)
- Split-K at nk=3328: 0.68 ms (8.7% of decode budget)
Attention is no longer the bottleneck. The 32 t/s wall is the model's matmul/FFN/weight-dequant floor, confirmed by the f16 V baseline hitting the same number.
llama.cpp's ROCm FlashAttention currently has two paths for prefill:
| Route | Math | KV Cache | PPL cost | VRAM cost |
|---|---|---|---|---|
| VEC / tile (quantized-KV) | FP16 vec/tile | q8_0 / q4_0 / TBQ4 | Lossy (quantization) | Low |
| MMA-f16 (WMMA temp) | FP16 WMMA | f16 temp buffer | Lossless | High (allocates temp) |
This is the standard tradeoff: choose between PPL fidelity and VRAM. Our kernel breaks the tradeoff — it uses INT8 sudot4 math on the default f16 KV cache, delivering FP16-equivalent PPL (1.0128) without allocating a temporary FP16 buffer and without committing to a quantized cache format.
RDNA3 GPUs have sudot4 (4× int8 dot product per cycle) but weak FP16 tensor support — WMMA requires strict 16×16×16 register layouts. Existing attention paths either ignore sudot4 entirely or try to use it through WMMA's constrained tile contract. Our kernel calls sudot4 directly in a scalar loop, avoiding WMMA's register pressure and layout constraints while still getting 4× the arithmetic density of scalar FP16.
The FlashAttention-2 paper (Dao et al.) established the tiled online-softmax design with large matmul tiles (M=128, N=128) sized for tensor cores. Our tiles are BM=8, BN=8 — sized for the scalar DOT4 body, not tensor core geometry. Making them larger just launches more scalar loops; it doesn't change compute density. This work explores whether small-tile INT8 FlashAttention can be competitive on consumer RDNA3 hardware where tensor cores aren't the answer.
Attention is 0.9% of total GPU time on Qwen3.6-27B (rocprof). The primary value is not throughput — it's correctness-safe acceleration plus an infrastructure path to VRAM savings via packed16 K cache.
Our kernel coexists with the existing attention routes in the same repository:
| Route | Lines | Math | KV Cache | Use Case |
|---|---|---|---|---|
fattn-vec.cuh (VEC) |
1040 | FP16 vec | Quantized (q8_0/q4_0/TBQ4) | Decode (1–2 tokens) |
fattn-tile.cuh (tile) |
1309 | FP16 tile | Quantized | Prefill (GQA, split-KV) |
fattn-mma-f16.cuh (MMA) |
2163 | FP16 WMMA | FP16 temp buffer | Prefill (batch > threshold) |
| Our DOT4 FA | 3793 | INT8 sudot4 | f16 (default) or q8_0 | Prefill + decode |
Every existing route either uses a quantized KV cache (VEC, tile — lossy, format lock-in) or an FP16 temporary buffer (MMA — high VRAM, lossless). Our kernel is the first that delivers FP16-equivalent PPL on the default f16 KV cache without a temporary buffer and without format commitment.
The 3793-line count includes Q quantization, K packing, 7 experimental variants, split-K combine, timing infrastructure, the tensor registry, and the packed16 cache management layer. The kernel body is ~190 lines. Existing routes spread equivalent infrastructure across separate headers and dispatch logic.
The CK-Tile FlashAttention (AMD's official ROCm reference) uses template-based tile composition with large FP16 tiles. We bypass it entirely — raw HIP with manual LDS lets us call sudot4 directly and manage the packed16 layout without fitting into a GEMM-centric tile contract. For a ~190-line kernel, the library abstraction overhead wasn't justified.
Three components that together unlock INT8 FlashAttention on consumer RDNA3:
| Component | What | Why it matters alone | Why it matters together |
|---|---|---|---|
| Packed16 K cache | f16 K → I32 payload + F16 scales, 16B-friendly layout | 2× KQ throughput vs q8_0 blocks | Separates K layout from K format — kernel gets DOT4-ready data without KV cache format change |
| I8 WMMA-sized 16×16 tile | QK computed as macro tile with sudot4 inner product |
4× arithmetic density vs FP16, no WMMA register constraints | Tile sized for INT8 work decomposition, not FP16 tensor core geometry |
| Raw HIP kernel | __global__ function, manual LDS, no CK-Tile/rocWMMA |
Full control over sudot4 calls and packed16 layout |
~190-line kernel body — no library overhead, no GEMM-centric tile contract |
Individually, none of these are novel. Packed INT8 layouts exist in Vulkan shaders. WMMA-sized tiles are standard FA2. Raw HIP kernels are how all of llama.cpp works. The combination is what makes it work: packed16K provides the data layout that the 16×16 tile consumes, the tile provides the work decomposition that makes INT8 attention efficient, and raw HIP provides the control to wire them together without abstraction overhead.
This also explains why previous approaches failed:
- WMMA-I8 in llama.cpp: tried to use
v_wmma_i32_16x16x16_iu8through rocWMMA's tile contract. The rigid 16×16×16 register layout and accumulator management killed occupancy. Our scalarsudot4loop inside a 16×16 logical tile avoids this. - CK-Tile FA: uses large FP16 tiles (M=128, N=128) with FP16 math. Expressing the INT8→FP16→softmax chain as CK-Tile operations adds abstraction overhead for a ~190-line kernel.
- Quantized KV cache routes: commit to a lossy format. Packed16 preserves f16 K fidelity while adding the INT8 math path.
| Parameter | Value |
|---|---|
| GPU | RX 7900 XTX, gfx1100, 24.6 GB VRAM, wave32 |
| ROCm | 7.2.3, hipcc, hipBLAS |
| Model (testing) | Qwen3.6-27B Q4_K_M, 48 layers, GQA=6, 2560 hidden, 9728 FFN |
| Build | build-rocm-rdna3-fa, GGML_HIPBLAS=ON, GGML_HIP_ROCWMMA_FATTN=ON |
| Arch constraint | RDNA3/4 only (sudot4 intrinsic) |
┌─────────────────────────────────────────────────────────┐ │ KV Cache (f16 K + f16 V) │ │ No format change required │ └────────────┬──────────────────────────┬─────────────────┘ │ │ ┌───────▼────────┐ ┌──────▼──────┐ │ cpy_k() │ │ f16 V │ │ f16→packed16 │ │ direct load │ │ quant kernel │ │ (no dequant)│ └───────┬────────┘ └──────┬──────┘ │ │ ┌───────▼──────────────────────────▼─────────┐ │ DOT4 FlashAttention Kernel │ │ │ │ ┌──────────────────────────────────┐ │ │ │ Q: f32 → quant → int8 + scales │ │ │ │ K: f16 → pack → int8 + scales │ │ │ └──────────────────────────────────┘ │ │ │ │ for each BM×BN tile: │ │ sudot4(Q_payload, K_payload) → logits │ │ online softmax(row_m, row_l) │ │ softmax @ V → output accumulator │ │ │ │ BM=8, BN=8/16, D=256 │ │ 64 DOT4 ops per QK pair │ │ V tile staging (LDS reuse) │ │ GQA6 head grouping (hpair/htriad) │ │ split-K partial combine │ └──────────────────────────────────────────────┘
### Pipeline stages (timed)
| Stage | What | Cost |
|-------|------|------|
| Q quantization | f32 → int8 payload + float scales (per-row absmax) | ~1.9 ms (pp2048) |
| K packing | f16 → int8 payload + half scales | ~0.3 ms (pp2048) |
| BlockFA body | Tiled sudot4 KQ + online softmax + PV | ~658 ms (pp2048) |
| split-K combine | Merge partial softmax states | ~9 ms (pp2048) |
K packing is <0.1% of route time. The blockFA body dominates.
### Tile Inner Loop (pseudocode)
```c
// One workgroup: BM query rows × BN key rows × D head dim
// Block size = BM*BN threads (each thread owns one QK logit)
float row_m[BM] = {-INF}; // online softmax row max
float row_l[BM] = {0}; // online softmax denominator
float out[BM][D] = {0}; // output accumulator
for (k_tile = 0; k_tile < nk; k_tile += BN) {
// STAGE: load BN rows of V into LDS (reused across BM query rows)
if (k_tile is in prefix region) {
// No causal masking — all query rows see all keys
for each (q,k) in BM×BN tile:
logit = 0;
for (qb = 0; qb < 8; qb++) // 8 q8 blocks of 32
for (i = 0; i < 8; i++) // 8 packed i32 per block
logit += sudot4(q_packed, k_packed); // 4×int8 dot per op
logit *= q_scale[qb] * k_scale[qb];
} else {
// Causal tail — only compute where k <= q_offset + q
}
// ONLINE SOFTMAX: update row_m, row_l, rescale accumulator
row_m = max(row_m, tile_max);
row_l = row_l * exp(old_max - row_m) + sum(exp(logits - row_m));
out *= exp(old_max - row_m);
// ACCUMULATE: P @ V
for (d = 0; d < D; d++)
out[q][d] += softmax(logit[q][k]) * V[k][d];
}
// Final: out[q][d] /= row_l[q]
The speed comparison (820 vs 902 t/s, ±6%) misses the point. FlashAttention isn't primarily about throughput — it's about memory. Without tiling, attention computes the full nq × nk logit matrix in HBM:
| Context length | Attention matrix (f32) | Feasible without FA? |
|---|---|---|
| 512 | 1 MB | Yes |
| 4k | 64 MB | Marginal |
| 16k | 1 GB | No (dominates VRAM) |
| 64k | 16 GB | No (exceeds GPU memory) |
| 128k | 64 GB | No |
With tiled online-softmax FlashAttention, the working set is one BM×BN logit tile plus row state — constant memory regardless of context length. The DOT4 FA kernel enables long-context inference on consumer GPUs where materializing the full attention matrix would OOM.
Every existing llama.cpp ROCm attention route (VEC, tile, MMA) already uses FlashAttention tiling. Our kernel adds INT8 math to that family. The comparison isn't "DOT4 FA vs no FA" — the comparison is "DOT4 FA vs other FA routes," and the advantage is PPL fidelity on default f16 KV cache without the FP16-temp memory cost.
| Metric | Value |
|---|---|
| Baseline f16 attention | 1.0128 |
| DOT4 FA (f16 K + f16 V) | 1.0128 (identical) |
| PPL drift | None |
Run 1: 815 t/s, PPL 1.0128
Run 2: 817 t/s, PPL 1.0128
The PPL test measures token prediction accuracy end-to-end through the full model — not a tile-level NRMSE. If cache corruption or X-input staleness existed, PPL would drift. It doesn't.
Open question: q8_0 K cache PPL at long context. Quantized K caches (q8_0, q4_0, TBQ4) lose precision per key row, and softmax amplifies those errors non-linearly as context grows. A proper PPL comparison between f16 K and q8_0 K at 64k–128k context would measure the real quality cost of quantized KV caches. We have not run this measurement. Our kernel supports both f16 K (lossless PPL) and q8_0 K (for VRAM-constrained use), but the quality gap between them at scale is unknown.
| Metric | Value |
|---|---|
| Prefill | 820–850 t/s |
| Decode (tg16) | 22–23 t/s |
| DOT4 KQ ceiling (standalone microbench) | ~874 t/s |
| Layout | ms | DOT4 GOP/s | vs q8_0 baseline |
|---|---|---|---|
| q8_0 block (34B) | 0.0385 | 435.8 | 1.00× |
| packed16 payload + q8 scales | 0.0193 | 870.9 | 2.00× |
| packed16 + one row scale | 0.0104 | 1606.4 | 3.69× |
Packed16 layout alone delivers 2× the KQ throughput of native q8_0 blocks.
| nq | nk | q8block ms | packed16 ms | scale-once ms | packed speedup |
|---|---|---|---|---|---|
| 256 | 1024 | 0.0316 | 0.0181 | 0.0173 | 1.75× |
| 512 | 2048 | 0.1361 | 0.0672 | 0.0570 | 2.02× |
| 1024 | 4096 | 0.3083 | 0.1080 | 0.1011 | 2.85× |
Moving scale application from per-DOT4 to per-q8-block yields only +4–18%. The packed16 layout is the primary win.
First route-backed A/B of the blockfa_hybrid_bm8_packed16_scalar kernel against stable default. This was the baseline before the GGML_OP_PACK_K_PACKED16 rebuild and the recthist prefix/tail work.
| Prompt | Default tok/s | blockFA tok/s | Delta |
|---|---|---|---|
| pp512 | 902.2 | 842.4 | -6.6% |
| pp1024 | 772.2 | 816.8 | +5.8% |
| pp2048 | 925.5 | 723.3 | -21.9% |
The pp1024 win (+5.8%) proved the packed16 DOT4 approach was viable. The pp2048 cliff (-21.9%) drove every architectural decision that followed: route timing diagnosis, the discovery that prefix keys were being treated as causal, and ultimately the recthist v2 prefix/tail split. The pp512 loss was consistent enough to be real but small enough to defer — fixing the cliff came first.
| Experiment | Result | What it taught us |
|---|---|---|
| BN16 wider K tiles | No gain | Per-row V dequant cost scales linearly with BN — no free lunch |
| BM16 more Q rows | No gain | Launch overhead already negligible at BM=8 |
| K LDS staging | -0.5–1.0% | LDS traffic exceeds reuse benefit; RDNA3 Infinity Cache handles packed16 K rows efficiently |
| Cooperative warp2 KQ | Slower than scalar | Shuffle reduction overhead > lane parallelism for BM=8 tiles |
| Register 2×2 micro-tile | 1.2–7.4× slower | Register pressure kills occupancy; scalar packed16 is leaner |
| LDS tile cache (16×16 QK) | 4.5–5.9× slower | Direct global/cache path beats explicit staging decisively |
| Fused V dequant | No gain | QK dot fully occupies SIMD units — no idle cycles to absorb V work |
Result: The scalar packed16 DOT4 body, despite being "just" one-thread-per-logit with 64 serial DOT4 ops, is the fastest KQ primitive on RDNA3 for BM=8 tiles. V tile staging and GQA6 head grouping are the only profitable additions. The pp2048 cliff isn't a KQ problem at all — it's a work partitioning problem, solved by recthist.
The pp2048 cliff (-21.9%) is caused by the kernel treating all 2048 keys as causal, when only the tail (k ≥ q_offset) needs causal masking. The prefix (k < q_offset) is fully visible to all queries.
Prefix k ∈ [0, q_offset): full non-causal, exact tiles, no lane masking
Tail k ∈ [q_offset, nk): causal check per (q,k)
Merge: single online softmax state across both regions
| Kernel | Total ms | Prefix ms | Tail ms | Merge ms | Rel RMS |
|---|---|---|---|---|---|
| recthist_v2_prefix_exact | 1.400 | 0.930 | 0.650 | 0.017 | 8.2e-07 |
| recthist_v2 | 1.418 | 0.936 | 0.615 | 0.018 | 8.2e-07 |
Correctness: 8.2e-07 relative RMS vs CPU reference. 191.7 DOT4 GOP/s, 766.9 PV GOP/s.
Not yet route-spliced — standalone harness only.
No amount of attention optimization moves throughput past ~870 t/s. The DOT4 FA kernel is already at 820–850 t/s — within 2–5% of the ceiling. Further attention work is about VRAM savings (packed16 K cache swap), not tokens per second.
Result: Replaced f16 K storage with packed16 (I32 payload + F16 scales). No f16 K allocation in main context.
| Component | Status |
|---|---|
GGML_OP_PACK_K_PACKED16 op |
✅ |
ggml_pack_k_packed16() graph node |
✅ |
k_payload / k_scales persistent buffers |
✅ |
cpy_k integration |
✅ |
| Tensor registry (K data → packed16 lookup) | ✅ |
| Skip-repack (no redundant pack_k on cache hit) | ✅ |
| Packed16-only (no f16 K allocation in main context) | ✅ |
| MTP draft context gate (f16 K fallback for draft) | ✅ |
| Multi-request lifecycle guards | ✅ |
| q4_0 / q8_0 V support | ✅ |
| Config | KV Cache VRAM |
|---|---|
| f16 K + f16 V | ~1568 MiB |
| f16 K + q4_0 V | ~960 MiB |
| packed16 K + q4_0 V | ~832 MiB |
GGML_CUDA_ROCM_Q8K_DOT4_PACKED16_K_CACHE=1 # enables packed16 K cache
GGML_CUDA_ROCM_Q8K_DOT4_DECODE_BN=64 # enables decode kernels
GGML_CUDA_ROCM_Q8K_DOT4_DECODE_TIMING=1 # per-token attention timing- Q: F32, D=256
- K: Q8_0 or F16, D=256
- V: Q4_0 or F16, D=256
- GQA: required (Q heads % K heads == 0)
- No: sinks, max_bias, logit_softcap
- Mask: F16 only, correct shape
- Gating: All paths behind
GGML_CUDA_ROCM_EXPERIMENTAL_UNSAFE=1+ route-specific flags. No default route change.
-
Profile non-attention decode: Attention is <9% of decode time. The 32 t/s wall is model matmul/FFN/weight-dequant. Next step: profile quantized matvec kernels to find the ceiling.
-
128k context validation: PPL and decode speed at extreme context lengths with packed16 K + q4_0 V.
-
MTP decode optimization: Current MTP verify uses v4 at nq=3, limiting throughput. Extend BN64/split-K to handle small batch queries (nq=2-4).
-
Production A/B against TBQ4 path: Compare packed16+q4_0 vs TBQ4 default for quality and throughput in real workloads.
- Dao et al., "FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness" (2022) — tiled online-softmax design
- Dao, "FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning" (2023) — sequence parallelism, warp partitioning
- AMD, "Implementing FlashAttention with CK-Tile" (ROCm blog) — AMD's official FA reference
docs/rocm-tbq4-paths/09-q8q4-dot4-packed-k-direction.md— Original packed16 directiondocs/rocm-tbq4-paths/23-q8k-dot4-fa-root-cause-and-next-plan.md— Why first FA prototypes faileddocs/rocm-tbq4-paths/24-wmma-sized-dot4-tile-plan.md— 16×16 WMMA-sized macro tile design.harness/research/rocm-packed16-blockfa-rectangular-history-kernel-design-20260526.md— Recthist design + A/B results.harness/research/rocm-q8k-dot4-blockfa-cooperative-kq-report-20260526.md— Cooperative KQ experiments