TurboQuant KV Cache Optimization

Gemma 4's K=V shared projections interact differently with KV cache quantization due to correlated K/V errors

Processing 2 elements at once with half2 and ternary sign reduces address count

Threshold choice across 5 orders of magnitude does not affect quality

Flipping computation to iterate over centroids and accumulate matching Q elements avoids per-element lookup

Storing first N tokens at fp16 improves PPL (sink tokens get disproportionate attention)

KV cache operations are negligible fraction of total decode compute at batch_size=1

Establish baseline PPL values for TBQ types on Qwen3.5-9B

turbo3 maintains retrieval accuracy with multiple distractor needles through 32K context

Fully branchless FMA chain with zero memory access beats constant LUT

Gemma-3 V cache broken because V un-rotation missing from iSWA build_attn overload

Fusing activation quantization into the weight GEMM kernel eliminates separate q8_1 quantize kernel launch + L2 traffic

turbo3 with FWHT rotation + norm correction matches q8_0 quality

QJL correction in turbo4 is unnecessary overhead

Alpha applied at decode time (scaling dequantized V) vs encode time (baked into fp16 norm before quantization) may give different results, with decode-time enabling context-adaptive deployment

Better byte extraction with 8-entry LUT improves over baseline

Per-channel scaling before FWHT reduces head_dim=128 quality gap by aligning channel variances

Subtract per-element mean before FWHT to improve quantization

turbo3 KLD tracks bit rate, not implementation quality

Capturing decode step as CUDA Graph eliminates per-kernel launch overhead (hundreds of launches per token)

Asymmetric q8_0/tbq3 will maintain <5% prefill gap and decode parity across diverse model architectures

The fused kernel's end-to-end slowdown is amplified/masked by non-attention compute (FFN, norms, embeddings). By comparing f16 KV (zero dequant overhead) vs TBQ3 baseline vs TBQ3 fused, we can isolate the attention-only time and understand what fraction of total time is available for optimization.

Since the SRHT (H, diag(r)) is fixed across all tokens, we can pre-rotate Q with forward SRHT once, compute K scores directly against centroids (no K butterfly), accumulate V in compressed domain (no V butterfly), and apply ONE inverse SRHT at the end. This eliminates 14 butterfly stages per KV token from the inner loop, reducing per-token compute by ~8x.

Using absolute query positions (nkv-nq+q_start+q_len) for causal skip correctly prunes Q batches without affecting PPL, saving 38-47% compute during full prefill

Adding head_dim=256 support to the fused TBQ3 attention kernel enables it to work with Qwen3.5-27B and Gemma-3 models. The kernel needs per-block SRHT (two independent 128-element butterflies for D=256) and FLOATS_PER_LANE=D/WARP_SIZE generalization.

Fusing TBQ3 dequant (inverse SRHT) directly into a FlashAttention-style online softmax kernel eliminates all intermediate buffers (k_tmp, v_tmp, S) while producing identical results

Bulk-dequanting V to fp16 before MMA launch (instead of fusing V dequant into the tile loader) will close the pp8192 gap because V can use the standard cp.async.cg pipeline

The fused kernel's Q addressing was swapped — it used q_head * (nq * D) instead of q_idx * (nh_q * D). After ggml_permute(0,2,1,3), Q physical layout is [nq, nh_q, D] (token-major) not [nh_q, nq, D] (head-major). Fixing this should make PPL match baseline.