ROCm gfx1100: fp8 KV cache disables all optimized decode attention kernels (7.5x regression at 16K context)

[ghazni101]

Row: GFX1100-TG200

Problem

--kv-cache-dtype fp8 forces the unoptimized PagedAttnOnline kernel on ROCm, because every optimized decode attention kernel requires k_cache.dtype == kBF16. The fp8 cache is DType::kI8, so all fast paths are skipped.

Measured on kind_tharp (Qwen3.5-4B Q4_K_M, RX 7900 XTX, ROCm 7.14.0, 128-token greedy decode, single request, 2026-08-27):

Context	fp8 KV tok/s	bf16 KV tok/s	Speedup
256	99.94	143.15	1.43x
1024	56.28	129.02	2.29x
4096	20.53	92.08	4.49x
8192	11.08	66.85	6.03x
16384	5.78	43.16	7.47x

The gap widens with context because PagedAttnOnline processes one key at a time with a full-block __syncthreads() reduction per key, while the optimized PagedAttnDecodeGqaF32Q kernel (activated by VT_ATTN_DECODE_GQA4=1) uses warp-strided KV walks with online softmax and no per-key sync.

Root cause

src/vt/rocm/rocm_paged_attn.hip:1878-2268 dispatch:

1. PagedAttnDecodeGqaBf16 (line 2117) — requires k_cache.dtype == kBF16
2. PagedAttnDecodeOptBf16 (line 2147) — requires k_cache.dtype == kBF16
3. PagedAttnDecodeGqaF32Q (line 2202, VT_ATTN_DECODE_GQA4=1) — requires k_cache.dtype == kBF16

All three fail the dtype check for fp8 KV (kI8). The dispatch falls through to PagedAttnOnline (line 2223).

The code acknowledges this at line 2231-2235:

"Only the correctness-grade PagedAttnOnline kernel serves fp8 — the decode-opt bf16 path stages __hip_bfloat16 fragments and a tensor-core fp8 read is a performance brick"

The fp8-kv-cache.md spec (line 293-304) also names this as owed work:

"An fp8 KV cache takes the subject OFF every fast attention kernel this engine has... The fp8 read through the fast kernels is owed below."

Architecture context

Qwen3.5-4B: 32 layers, full_attention_interval=4 (8 full-attn layers), 4 KV heads, head_dim=256, f32 query. The VT_ATTN_DECODE_GQA4=1 flag activates PagedAttnDecodeGqaF32Q for the f32-query + bf16-KV combination — the exact match for this model — but only when KV is bf16.

Fix direction

Widen the PagedAttnDecodeGqaF32Q dtype guard at line 2186-2189 from kBF16 to kBF16 || kI8, and add an fp8 dequant load path inside the kernel (the LoadKv(uint8_t*, ...) helper at line 176 already does Fp8E4M3ToF32Dev(byte) * scale).

Owning row

GFX1100-TG200 campaign (fork issue #5).