#120 `--kv k8v4`: hybrid KV cache — INT8 K + Hadamard-rotated Q4_0 V (816 B/cell)

描述

# `--kv k8v4`: hybrid KV cache — INT8 K + Hadamard-rotated Q4_0 V (816 B/cell)

Rebased on v0.1.21 (the multi-GPU layer split); all parity checks and the needle
benchmarks below re-run on it.

https://github.com/Niko1221/Strata/issues/116

## What

A fourth `--kv` choice: **K stays INT8** (unrotated — the attention scores stay exactly as
precise as `int8`'s) and **V uses PR #21's rotated q4_0** (kv_q4.hpp). 816 B per cell vs
int8's 1,056 (−23%) and q4_0's 576. The quantization-sensitivity asymmetry is the point:
K enters through the softmax scores (errors redirect attention), V only through a
softmax-weighted average (errors are damped) — the same reasoning behind llama.cpp's
q8_K/q4_V-style hybrids.

## How

- **No new kernels.** The append/gather entry points already select K/V by `blockIdx.z` and
  take separate K/V pointers, so the hybrid composes the existing validated q8/q4 kernels
  with the unused half's lanes folded onto the used pool (a bit-identical duplicate write;
  ~µs of redundant work per token). The fused decode/prefill/verify attention
  (`attn_chunk_kernel`) gets `KV_MODE == 3`: the K loader takes the int8 branch (load8 is
  K-only today), the inline V loop already falls into the q4 branch. Only V and the output
  are Hadamard-rotated; q is not, so `<q, k>` needs no rotation.
- **State**: `QsaState::kv_hybrid` uses `k_q`/`k_scale` + `v_q4`; sizing/init/zero/pools
  handle it. The MTP drafter stays INT8 under this setting (mtp.cpp toggles the format
  globals around its state creation, whatever ring shape it takes) so the kv_stream/kv_ring
  block movers never see the hybrid layout. `--kv-resident` streaming is refused with it.
- **setup.py**: `--kv k8v4` accepted; the streaming RAM math uses 816 B/token.

## Tests

New `kv_hybrid_parity` (built alongside kv_q8_parity/kv_q4_parity): bitwise append/gather
checks through a non-identity page table against host references (INT8 rules with
ties-to-even; rotated q4_0 rules), and mode-3 attention — step and batched forms, the
latter with per-query step arrays as verify uses it — against a host attention over the same
dequantized K/V. All checks pass at machine precision (≤ 1.1e-07); kv_q8_parity and
kv_q4_parity are unchanged.

## Measured (RTX 3090 24 GB, Coder IQ1_M, 118.75k-token prompts, temp 0)

| | int8 | **k8v4** |
|---|---|---|
| Decode @~119k ctx | 75 t/s | **96.6 t/s (+29%)** |
| Cold prefill | 1,170 t/s | 1,171 t/s |
| 26-needle suite, seed 42 | 24/26 | **24/26 (identical tokens)** |
| KV bytes @131k ctx | 1.80 GB | 1.39 GB |

Decode moves ±10% between runs (draft acceptance varies with the generated text), so the
honest band at 119k is 82–97 t/s against int8's 75; the needles reproduce exactly. The gain
is larger than the expert-cache headroom alone would predict because the q4 V side also
halves the attention's V-gather bytes (144 vs 256 B per head per cell).
At `--max-context 204800` on the same 24 GB card: 1,024 t/s prefill, 82.1 t/s decode,
25/26 needles at 197.9k tokens (two seeds), fact-check follow-ups correct.

## Not included (kept out on purpose)

- KV streaming + hybrid (`--kv-resident`): refused for now; the hybrid block layout would
  need its own mover in kv_stream.cu.
- A drafter in hybrid format: the draft layer stays INT8, which also keeps its ring path
  whole-format.

## AI tooling

Work assisted with Claude-code using GLM-5.3