[bsorensen110]

KV cache: FP16 sliding window + attention sinks, and an optional 2-bit K / 4-bit V cache scheme (--kv q2k) [EXPERIMENTAL]

Status: experimental. Opt-in only (--kv q2k, --kv-window N, --kv-k-gain B); default behavior is unchanged when the flags are off. It needs more testing before it can be considered production-ready: the behavioral and long-context runs below are single-model, single-document measurements, and the window-1024 anomaly is unexplained.

Summary

Two composable KV-cache features behind opt-in flags, plus a correction that makes them work together:

1. --kv-window N (--kv-sink S) — an FP16 overlay over the first S cells (attention sinks) and the most recent N cells of the quantized KV bulk. Those cells read exact; the rest stays quantized. Costs 304 KiB per layer at N=128, S=8 including 16 slack slots (3.56 MiB for 12 layers).
2. --kv q2k — a cache format storing K at 2 bits with per-channel scales and V at 4 bits per token (q4_0): 448 B/cell vs q4_0's 576 B (−22%), 4.6× smaller than FP16 (2048 B).
3. --kv-k-gain B (q2/q2k) — a K reconstruction gain: the stored K scale is B × the rule's, the codes are unchanged. It undoes the 2-bit codebook's score shrink, which is what made the window hurt q2k. Zero extra bytes.

Recommended with q2k: --kv-window 256 --kv-sink 8 --kv-k-gain 1.10 — decode-path KL −36.8% vs plain q2k (8 of 9 texts) for +6.6 MiB.

Correction to the previous revision (run-to-run variance). The last revision said greedy decode is not run-to-run reproducible "under QSA top-k tie-breaks and atomic partial-group KV writes" and reported decode-path numbers as 5-run means. That diagnosis was wrong, and the cause was our measurement setup: we did not have the prefill CPU share turned off. For prompt chunks under 1024 tokens the prefill hands part of the MoE expert work to the CPU, and that split changes from run to run — a 730-token prompt's logits differed by up to 6.8 between two identical runs. With STRATA_PREFILL_CPU_SHARE=0 two runs are bit-identical (129 teacher-forced decode rows, max |diff| 0; a second forced FP16 run scores KL 0.0000 against the first on all nine texts below). Every number in this revision was measured that way, and the 512-token logits table was re-measured (it moved by up to 2.5 points of top-1). The earlier decode-path comparison also used a greedy reference, which speculatively drafts multi-token windows and so does not produce rows comparable to the arms'; the decode numbers below use a teacher-forced FP16 reference instead.

Every directional claim below carries the number that pins it; the gate (kv_q2_parity) reproduces the attention-error tables from a real KV dump.

Why this shape

The bit-allocation rule came from the real cache, not the synthetic one

Measured with the gate's attention-output probe (relative RMS of attention output vs an FP16 cache) on a real stored KV cache dumped from a live serve (--dump-kv; 864 cells, 2 KV heads × 256 dims, rotated as stored), layer 3 unless stated:

arm	B/cell/head	attention error
q4_0 (reference)	288	0.888%
K 2-bit per-channel + V 4-bit (q2k)	224	1.911%
K 4-bit + V 2-bit per-channel	224	3.171%
K 2-bit + V 2-bit per-channel (symmetric q2)	160	3.508%
K per-channel + V per-token, both 2-bit (KIVI's placement)	160	5.188%

At equal bytes (224 B), the bits belong on V: K2c+V4 beats K4+V2c on every real layer measured (table below). On the synthetic gate cache the ranking is the reverse (K looks more sensitive) — synthetic data has no channel structure, which is why no ranking here is taken from it.

The mechanism: V's quantization error enters the output directly through softmax(qKᵀ)·V, while much of K's error is common-mode across cells and cancels in the softmax normalization (docs/kv-q2.md). Both published intuitions survive, scoped:

• K's outliers are channel-aligned (KIVI's premise) — true on real layers (K channel RMS max/median 5.8× on L3 up to 74.9× on L47), and honored by the per-channel K scale, not by more bits.
• V is the tensor the output feels — so V gets 4 bits while K takes 2.

Why per-channel K and per-token V

Per-channel K scales (block_q2c_0: one fp16 scale per dimension per 32-cell group, a 4-level Lloyd-Max codebook) absorb K's channel outliers. V reuses the existing q4_0 path (per-token blocks), so the format adds no new V quantizer. KIVI's 2-bit per-token V loses on this cache: the Hadamard rotation leaves V's channel structure in place (V channel RMS max/median 7.9–14.5×), so 2-bit V would want per-channel scales too — at 4 bits per token, V carries q4_0's own error.

Why the FP16 window — and why q2k needs the gain to use it

The window is a separate, format-agnostic device. It only acts at decode: prompt-mode (prefill) logits never read the overlay, which is why the previous revision measured "no logit precision" — that was a blind spot of the measurement, not a property of the window.

• At decode it helps q4_0: −19.6% KL (7 of 9 texts).
• With q2k alone it hurts: +10.3% (4 of 9 better). The 4-level Lloyd-Max K codebook is MSE-optimal, so its reconstruction is shrunk: on real queries q2k's attention scores come out at 0.92–0.97× the exact ones (q4_0: 0.997–0.998). Next to exact window cells that is a bias — the exact cells out-score the shrunk bulk, and the softmax over-weights the window by +1.7..+5.9% of its mass on 11 of 12 layers.
• --kv-k-gain stores the K scale B × larger while choosing the codes against the un-gained scale, so no value moves to another level and only the reconstruction grows. With B = 1.10 the window turns into a gain: −29.3% at 128 cells, −36.8% at 256 (vs plain q2k).
• Behaviorally the window recovers symmetric q2's short-answer loss (q2 2/3 → q2 + window 3/3); q2k matches the q4_0 reference on the recall probe without it (3/3).

Measurements

Gate (device kernels vs host reference)

• kv_q2_parity: synthetic cache 40/40 checks (incl. the K-gain pin: stored scale = B × the rule, bit for bit, codes chosen against the un-gained scale); real cache layers 3/11/23/47 36/36 each. The q2k device arm (real append + gather kernels) lands on the host row: L3 1.910% device vs 1.911% host.
• kv_window_parity: 10/10 — sink and window cells read bit-exactly equal to an FP16 pool; cells outside read the bulk pool bit-for-bit; kv_window_restore reproduces the bulk dequantization.
• kv_q4_parity, kv_q8_parity, kv_stream_parity, kv_hybrid_parity: pass (existing formats untouched).

Attention error by layer (real cache)

layer	K channel spread	q4_0	q2k (K2c+V4)	K4+V2c	q2 (K2c+V2c)
L3	5.8×	0.888%	1.911%	3.171%	3.508%
L11	5.9×	0.740%	2.002%	2.589%	3.471%
L23	14.2×	0.941%	2.211%	2.768%	3.655%
L47	74.9×	1.040%	1.645%	2.663%	2.948%

The outlier layer (L47) is where q2k comes closest to q4_0.

Attention error under the engine's append patterns (real cache)

The row above appends every cell in one call — whole 32-cell groups — so it never runs the partial-group scale rule. The engine does, every request: speculative verify appends 2–5-cell windows that cross group boundaries, decode appends into the group the prompt ended in, and after a reused conversation prefix the 2048-cell prefill chunks fall mid-group. kv_q2_parity now replays those patterns on the same K/V and readers (attention output vs FP16). Three scale rules for a group first written by a partial batch: old = the pre-review kernel (sum / 32), /n = RMS of the n cells, shrink = shipped (full-group RMS, the previous group's mean square standing in for the cells not yet written):

layer	q4_0	q2k one-shot	spec verify (3-cell windows) old / /n / shrink	chunked prefill (100-cell chunks, stress) old / /n / shrink	prompt + decode
L3	0.888%	1.910%	1.978% / 1.912% / 1.892%	3.631% / 2.076% / 2.186%	1.910%
L11	0.740%	2.002%	2.030% / 2.011% / 2.012%	3.617% / 2.727% / 2.626%	2.002%
L23	0.941%	2.211%	2.895% / 2.667% / 2.245%	3.767% / 2.509% / 2.161%	2.211%
L47	1.040%	1.648%	1.457% / 1.799% / 1.720%	1.438% / 1.346% / 1.803%	1.648%

• Speculative verify (what every decode does at each group boundary): the excess over one-shot, summed over the four layers, is +8.0% with /n and +1.3% with shrink.
• Chunked prefill is a stress case (real chunks are 2048 cells, aligned from position 0): shrink is mixed by layer there (1.4% worse summed than /n) and 30% better than the pre-review kernel.
• Prompt + decode lands exactly on the one-shot row (pinned). On the extreme-outlier layer 47 an undersized scale acts as clipping, which is why the pre-review kernel measures better there.

Logits (512-token prompts, prompt mode, vs FP16: top-1 agreement / mean KL) — re-measured with the CPU share off

text	q4_0	q2	q2k
code	90.4% / 0.060	74.6% / 0.329	77.1% / 0.226
document	91.6% / 0.030	81.2% / 0.203	83.8% / 0.133
chat	92.4% / 0.027	81.2% / 0.177	83.4% / 0.144

q2k's KL is 3.8–5.3× q4_0's at the prompt level: it is a cheaper format, not a lossless one. (The previous revision's table, measured with the CPU share on, read 2.5–5×.)

Decode path (teacher-forced decode, 9 texts, vs FP16)

Every arm is fed the same 128-token continuation, one token per window, and each row is compared with a forced FP16 run's row at the same position. Nine texts: three built-in prompts (chat, document, code agent) and six repo files up to 6144 tokens. A second FP16 run scores KL 0.0000 against the reference on all nine. Change = geometric mean over the texts vs plain q2k, with the texts that improved.

arm	mean KL	top-1	vs q2k
q4_0	0.0198	96.6%	−89.9% (9/9)
q4_0 + window 128	0.0174	96.9%	−91.9% (9/9)
q2k	0.1825	88.5%
q2k + window 128	0.2060	87.3%	+10.3% (4/9)
q2k + window 128 + gain 1.10	0.1338	89.8%	−29.3% (7/9)
q2k + window 256 + gain 1.10	0.1203	90.3%	−36.8% (8/9)

Gain sweep at window 128 (same texts, development build): 1.03 −9%, 1.05 −15%, 1.08 −26%, 1.10 −30%, 1.13 −29%; the gain alone, with no exact cells, −1%. Even the best q2k setting keeps ~6× q4_0's decode KL.

These rows were measured on this PR's build with a temporary copy of our teacher-forced decode tooling (--force-tokens + a per-window logits dump; not part of this PR). The PR build's prompt logits are byte-identical to our development build's under the same flags (--kv q2k, and --kv q2k --kv-window 256 --kv-sink 8 --kv-k-gain 1.10). The two builds sit on different upstream bases, and their decode rows differ slightly for every quantized format — plain q4_0 included — but never for FP16, so the development build's numbers are not quoted as this PR's.

Behavioral A/B (~5K-token buried-fact probe, tools/kv_recall_report.py)

Eight vault codes buried in a document: mid = the body; late = the last few paragraphs (~180–350 tokens before generation: outside a 128-cell window, inside a 512/1024 one). Scored by the instructed (vault, code) pair. Release q4_0 is the shipped binary; every other arm is this branch.

arm	mid	late	short answers	decode
release q4_0 (reference)	4/4	4/4	3/3	158.3 t/s
q4_0 (this branch)	4/4	4/4	3/3	159.0 t/s
q2 (symmetric)	4/4	4/4	2/3	134.6 t/s
q2k (no window)	4/4	4/4	3/3	137.5 t/s

q2k matches the q4_0 reference on every measure at 22% less KV memory and ~13% lower decode speed; symmetric q2 is the arm that loses a short answer. An earlier 5-arm run (substring-scored) agreed, with 0/3 repetition loops for release q4_0 and q2k.

Long-context recall at depth (smoke test)

~200K-token operations log (199,337 prompt tokens) served with --kv q2k --max-context 262144; three needles at 25% / 50% / 90% depth — all three recalled exactly (WALRUS-CANARY-2247, QUARTZ-DELTA-8813, TUNDRA-FALCON-6631). Cold prefill of the full prompt: 2574 tok/s. The verification pass rode the serve's conversation cache, which restores the same q2k-quantized state. Re-run on the final kernel (after the partial-group scale changes, fresh service, nothing cached): cold prefill of all 199,337 tokens at 2529 tok/s, 3/3 recalled. Re-run with the recommended --kv q2k --kv-window 256 --kv-sink 8 --kv-k-gain 1.10 (fresh service, nothing cached): 199,337 tokens, 3/3 recalled, 78.6 s wall including generation (~2537 tok/s). One model, one document.

A small multi-turn agent task (find and fix two planted bugs across two files, iterate pytest to green) also completed correctly against the served q2k/262144 endpoint.

Window sizing (symmetric q2 bulk + window)

arm	mid	late	short answers	decode
q2 + window 128	4/4	4/4	3/3	127.7 t/s
q2 + window 512	4/4	4/4	3/3	136.1 t/s
q2 + window 1024	3/4	0/4	3/3	129.0 t/s

128 and 512 are behaviorally identical (512 decodes faster: the FP16 overlay is a cheaper read than the q2 bulk). 1024 regresses — the late facts sit inside a 1024-cell window, so exact reads should help, not collapse; this points at a defect at that size, not a trade-off. Unexplained; do not use 1024.

Memory (262144 ctx, 12 QSA layers, 2 KV heads × 256 dims; the gain adds nothing)

format	B/cell	KV total
FP16	2048	6.00 GiB
q4_0	576	1.69 GiB
q2k	448	1.31 GiB
+ window 128 / sinks 8	—	+ 3.56 MiB
+ window 256 / sinks 8	—	+ 6.56 MiB

Reference papers

• KIVI (arXiv 2402.02750) — K per-channel, V per-token, asymmetric 2-bit, 128 recent tokens in FP16. The K placement transfers (per-channel K wins on real layers); the V placement does not on this rotated cache (K2c+V2t 5.188% vs symmetric per-channel 3.508% at the same 160 B).
• SKVQ (arXiv 2405.06219) — channel reordering, clipped dynamic quantization, an FP16 sliding window; K 2-bit / V 1.5-bit. Reordering is a no-op for per-channel-scaled K (a relabeling). During development (tooling not included in this PR) it helped per-token V on real layers, but SKVQ's byte budget still lost to plain per-channel 2-bit on every real layer. Clipping did not help (L3, in this PR's gate: clipped 4.731% vs unclipped 4.627%).
• KVSink (arXiv 2508.04257) — attention-sink preservation under KV quantization; the reason the overlay keeps the first cells exact.
• QuaRot (arXiv 2404.00456) — the Hadamard rotation the KV formats here (q4_0, q2, q2k) apply before quantizing.
• KVQuant (arXiv 2401.18079) — per-channel (pre-RoPE) K and non-uniform codebooks; background for the per-channel K scale and the Lloyd-Max codebook.
• "More for Keys, Less for Values" (arXiv 2502.15075) — argues K deserves more bits. At equal bytes on this rotated cache the measurement went the other way (above).
• AsymKV (arXiv 2410.13212) — the 2-bit-K / 1-bit-V split; 1-bit V measured 9.7% (per-channel) to 14.4% (per-token) attention error on L3 here.

Review fixes (Copilot review of this PR)

All 15 findings were reproduced and fixed; the engine fixes are pinned by new gate checks, each shown to FAIL against the pre-fix code:

• q2/q2k pools are sized and moved in whole 32-cell groups (high): fully resident pools, elastic VMM pools, the staging copy, the reverse DMA, and conversation snapshots (image size and reusable prefix). A page is not a byte range in the group-major layout.
• Fresh partial-group scale: the RMS of the cells present (/n), not /32 — then, measured under the engine's append patterns (table above), shrunk toward the previous group's scale in the batch path too: s = sqrt((ss + (32 − n)·prev²)/32). prev is race-free: a previous group completed in the same launch is recomputed in-thread with its block's exact sum order; otherwise its stored scale is read. Pinned bit-exact (two-call and one-call appends must store the same scale).
• --dump-logits header repair on every path, not only the native pack (CUDA and SYCL).
• Tools: truncated KV dumps and zero-vocab logit headers fail cleanly instead of crashing; the no-e validator checks the letter anywhere; the engine-failure row prints; recall is scored by the (vault, code) pair (codes repeat between groups). Re-scoring the saved outputs left every arm above unchanged except q2 + window 1024 (late 1/4 → 0/4).
• kv_pr_report.py: each tail code is now asked in its own request, inside the window. The 2026-10-11 README carries a correction: its earlier "tail" codes sat outside the window.
• Docs: the overlay footprint (304 KiB/layer, 3.56 MiB total) and the q4_0 baseline (1.69 GiB at 262144) were ~2× off; a stale "dump-logits writes 0 bytes" note was removed.
• Also fixed while re-measuring: two inverted gate labels (the K-vs-V check printed "K is the sensitive side" while proving V is), and the AsymKV citation (2410.13212, not 2411.17034).
• No machine-specific paths: the harnesses and A/B scripts resolve everything from the repo's layout (tools/kv_bench_env.py: repo root, Strata-data next to it, strata-<model>.json, engines under engine/ and build/) with an environment override for each; the READMEs carry a generic CUDA 13.0 build recipe instead of local scripts.

Known limits

• 2-bit K is the quality floor: q2k's prompt-level KL is 3.8–5.3× q4_0's, and even window 256 + gain 1.10 keeps ~6× q4_0's decode KL.
• Reproducible measurements need STRATA_PREFILL_CPU_SHARE=0 (prompt chunks under 1024 tokens otherwise split MoE experts with the CPU, and the split varies run to run). Engine default unchanged.
• The gain is calibrated on one model (uniform 1.10, from a 1.03–1.13 sweep on nine texts).
• Window 1024 regresses recall (above); unexplained.
• Behavioral and long-context results are single-model, single-document.

Commits

1. kv_window — FP16 sliding window + attention sinks over the quantized KV cache
2. kv_q2 — 2-bit per-channel KV format, parity gate, measurement tooling
3. kv_q2k — K 2-bit per-channel + V 4-bit mixed-width format
4. kv_q2 — whole 32-cell groups wherever a q2/q2k pool is sized or moved
5. kv_q2 — a fresh partial group's scale is the RMS of the cells it has
6. kv_q2_parity — pins for the review fixes, and the inverted K-vs-V labels
7. generate — repair the --dump-logits header on every path
8. tools — crash fixes and (vault, code) pair scoring
9. kv_pr_report — tail probes inside the window
10. docs — footprint arithmetic, the tail-probe correction, the stale dump-logits gap, the AsymKV citation
11. tools + scripts — no machine-specific paths: repo-relative defaults, env overrides
12. docs — a generic CUDA 13.0 build recipe instead of local paths and scripts
13. kv_q2_parity — q2k under the engine's append patterns (spec verify, chunked prefill, prompt + decode)
14. kv_q2 — a fresh partial group's scale shrinks toward the previous group's, in the batch path too
15. kv_q2 — K reconstruction gain (--kv-k-gain): undo the 2-bit score shrink, codes unchanged, zero bytes

[Copilot]

🟡 Changes recommended

Group-granular q2 sizing, partial-group scaling, and staging copies can corrupt or incompletely populate KV data.

15 open findings

• Round pool capacity up to complete q2 groups · New
• Divide partial-group scale by available cell count · New
• Copy complete q2 residency units during staging · New
• Copy complete q2 units during reverse DMA · New
• Repair headers for shortened non-native runs · New
• Repair dump headers for all shortened runs · New
• Return cleanly after detecting truncated dump bodies · New
• Keep tail probes within the overlay window · New
• Handle six-column failure rows safely · New
• Match vault scores to the instructed code pair · New
• Reject responses containing the letter e · New
• Reject non-positive vocabulary dimensions · New
• Remove the stale dump-logits known-gap claim · New
• Correct overlay and q4_0 footprint calculations · New
• Include slack slots in footprint estimates · New

What changed in this PR

Adds experimental q2/q2k KV-cache formats, an FP16 recent-token/sink overlay, and supporting validation and measurement tooling.

Changes:

• Implements q2/q2k storage across streaming, snapshots, speculative decoding, and attention.
• Adds FP16 window/sink preservation and parity coverage.
• Extends logits dumping, benchmarks, scripts, and documentation.

File	Description
tools/​logits_dump_check.py	Validates logits dumps.
tools/​kv_release_ab.py	Compares release behavior.
tools/​kv_recall_report.py	Measures KV recall and speed.
tools/​kv_pr_report.py	Produces window quality metrics.
tools/​kv_loop_ab.py	Measures repetition regressions.
tools/​kv_cache_dump.py	Analyzes FP16 KV dumps.
tools/​kv_ab_report.py	Runs KV accuracy/speed A/Bs.
sycl/​src/​program/​generate.cpp	Dumps native speculative logits.
src/​kernels/​kv_window_parity.cpp	Tests overlay correctness.
src/​kernels/​cuda/​kv_window.cu	Implements overlay writes/restoration.
src/​kernels/​cuda/​kv_stream.cu	Adds q2/q2k streaming.
src/​kernels/​cuda/​kv_q2.cu	Implements per-channel q2 storage.
src/​core/​verify.cpp	Integrates q2 and overlays into verification.
src/​core/​mtp.cpp	Adds q2/q2k drafter appends.
src/​core/​conversation_snapshot.cpp	Serializes and restores new formats.
logits-kv-window.sh	Automates logits comparisons.
kv-ab-report.sh	Runs the KV report.
include/​strata/​program/​logits_selection.hpp	Counts selected logits rows.
include/​strata/​prefill/​prefill.hpp	Exposes prompt-logit callbacks.
include/​strata/​kernels/​qsa_decode_attn.hpp	Extends attention pool metadata.
include/​strata/​kernels/​kv_window.hpp	Defines overlay layout and API.
include/​strata/​kernels/​kv_stream.hpp	Defines q2/q2k streaming units.
include/​strata/​kernels/​kv_q2.hpp	Defines q2 format and API.
include/​strata/​core/​layer.hpp	Extends QSA state and configuration.
docs/​research-kv-cache-compression.md	Documents compression research.
docs/​kv-window.md	Documents the FP16 overlay.
CMakeLists.txt	Builds kernels and parity tests.
bench/​results/​2026-10-12-kv-precision/​README.md	Records precision measurements.
bench/​results/​2026-10-11-kv-window/​README.md	Records window measurements.
ab-kv-window.sh	Runs resident overlay A/Bs.
ab-kv-window-stream.sh	Runs streamed overlay A/Bs.
ab-kv-window-probe.sh	Runs short-answer overlay probes.

🧠 Review effort: Balanced

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[Tsury]

Validated the overlay-only INT8 path from1aee5c90ffa6f0ddfb40c22db51ea3aacc1f66e8, which is unchanged for this path at head b47ec5b. CUDA13.3/SM120, window128/sinks8, resident identity page table, H=2 and D=256. The q2/q2k additions were not exercised.

Both the existing window parity gate and this small INT8 adaptation pass. Fresh sink/recent attention is bit-identical to the FP16 reference; cold reads are unchanged. Batch, single-step and six-step speculative writers agree. Restore equals half(INT8 dequantization), as expected, rather than recovering original FP16. In the seeded synthetic fixture, sink/recent K RMS error goes0.573%/0.575% to0%; cold error remains0.592%;134021/139264restored hot K/V values differ from original FP16. These are synthetic precision observations, not model-quality claims.

Reproduce by copying src/kernels/kv_window_parity.cpp to src/kernels/kv_window_int8_parity.cpp, applying the patch below, and adding a target alongside kv_window_parity:

add_executable(kv_window_int8_parity src/kernels/kv_window_int8_parity.cpp)
target_include_directories(kv_window_int8_parity PRIVATE )
target_link_libraries(kv_window_int8_parity PRIVATE strata_kernels )

Build with STRATA_ENABLE_CUDA=ON and STRATA_BUILD_TESTS=ON, then run kv_window_int8_parity. The existing kv_q8_parity, kv_stream_parity and conversation_snapshot_test gates also pass. The restored-precision limitation matches the documented snapshot behavior.

Synthetic INT8 parity adaptation; no model or private inputs

--- a/src/kernels/kv_window_int8_parity.cpp
+++ b/src/kernels/kv_window_int8_parity.cpp
@@ -1,4 +1,4 @@
-// src/kernels/kv_window_parity.cpp - the FP16 sliding window + attention sinks (kv_window.hpp) against a
+// src/kernels/kv_window_int8_parity.cpp - the FP16 sliding window + attention sinks (kv_window.hpp) against a
 // full-precision reference and against the quantized path it overlays (GPU, no model).
 //
 // The pools are built by hand: a Q4_0 pool and the FP16 overlay on one side, an FP16 pool holding the same cells on
@@ -10,13 +10,14 @@
 //      is untouched;
 //   4. kv_window_restore reproduces the bulk's own dequantization bitwise for the sinks and for the window.
 #include "strata/kernels/f16_bits.hpp"
-#include "strata/kernels/kv_q4.hpp"
+#include "strata/kernels/kv_q8.hpp"
 #include "strata/kernels/kv_window.hpp"
 #include "strata/kernels/qsa.hpp"
 #include "strata/kernels/qsa_decode_attn.hpp"
 
 #include <cuda_runtime.h>
 
+#include <algorithm>
 #include <cmath>
 #include <cstdio>
 #include <cstring>
@@ -59,7 +60,7 @@
     const int pages = N / P;
     const k::KvWindowShape shape{WINDOW, SINKS};
     const int rows = N * H;                    // one row per (cell, kv head) in the identity layout
-    const size_t q4_bytes = (size_t) rows * k::kv_q4_bytes_per_head(D);
+    const size_t q4_bytes = (size_t) rows * D;
 
     std::mt19937 rng(23);
     std::normal_distribution<float> nd(0.f, 1.f);
@@ -77,7 +78,8 @@
     int32_t* d_table = dalloc<int32_t>(pages);
     ck(cudaMemcpy(d_table, table.data(), pages * 4, cudaMemcpyHostToDevice), "table");
 
-    uint8_t *kq4 = dalloc<uint8_t>(q4_bytes), *vq4 = dalloc<uint8_t>(q4_bytes);
+    int8_t *kq4 = dalloc<int8_t>(q4_bytes), *vq4 = dalloc<int8_t>(q4_bytes);
+    uint16_t *ks = dalloc<uint16_t>((size_t) rows * (D / 64)), *vs = dalloc<uint16_t>((size_t) rows * (D / 64));
     uint16_t *kp = dalloc<uint16_t>((size_t) rows * D), *vp = dalloc<uint16_t>((size_t) rows * D);
     const int64_t win_rows = k::kv_window_slots(shape) * H * D;
     uint16_t* kw = dalloc<uint16_t>((size_t) win_rows);
@@ -95,7 +97,7 @@
         ck(cudaMemcpy(vcur, vc, (size_t) H * D * 4, cudaMemcpyHostToDevice), "v");
         int32_t hstep[k::kStepCount] = {cell, cell + 1, 0, 0};
         ck(cudaMemcpy(step, hstep, sizeof hstep, cudaMemcpyHostToDevice), "step");
-        k::kv_append_q4_step(kq4, vq4, d_table, step, kcur, vcur, s, nullptr, nullptr);
+        k::kv_append_q8_step(kq4, vq4, ks, vs, d_table, step, kcur, vcur, s, nullptr, nullptr);
         k::kv_append_step(kp, vp, d_table, step, kcur, vcur, s, nullptr, nullptr);
     }
     ck(cudaDeviceSynchronize(), "appends");
@@ -128,8 +130,8 @@
             p.k_pool = kp;
             p.v_pool = vp;
         } else {
-            p.k_q4 = kq4;
-            p.v_q4 = vq4;
+            p.k_q = kq4; p.k_scale = ks;
+            p.v_q = vq4; p.v_scale = vs;
         }
         if (overlay) {
             p.k_win = kw;
@@ -153,7 +155,7 @@
     const std::vector<float> a_win = run(hot, true, false);
     const std::vector<float> a_q4 = run(hot, false, false);
     check(max_diff(a_win, a_ref) == 0.0f, "sinks + window read exactly (bitwise equal to the FP16 pool)");
-    check(max_diff(a_q4, a_ref) > 0.0f, "q4_0 alone is not exact (the overlay is what removes the rounding)");
+    check(max_diff(a_q4, a_ref) > 0.0f, "INT8 alone is not exact (the overlay is what removes the rounding)");
 
     // 3. cells outside the window and the sinks: the overlay changes nothing
     std::vector<int32_t> cold;
@@ -168,18 +170,21 @@
     ck(cudaMemset(vw, 0, (size_t) win_rows * 2), "clear v overlay");
     k::QsaAttnPools bulk;
     bulk.page_table = d_table;
-    bulk.k_q4 = kq4;
-    bulk.v_q4 = vq4;
-    k::kv_window_restore({kw, vw}, shape, bulk, 2, N, s, nullptr);
+    bulk.k_q = kq4; bulk.k_scale = ks;
+    bulk.v_q = vq4; bulk.v_scale = vs;
+    k::kv_window_restore({kw, vw}, shape, bulk, 1, N, s, nullptr);
     ck(cudaDeviceSynchronize(), "restore");
-    std::vector<uint8_t> hq4(q4_bytes), hv4(q4_bytes);
+    std::vector<int8_t> hq4(q4_bytes), hv4(q4_bytes);
+    std::vector<uint16_t> hks((size_t) rows * (D/64)), hvs((size_t) rows * (D/64));
+    ck(cudaMemcpy(hks.data(), ks, hks.size()*2, cudaMemcpyDeviceToHost), "read K scales");
+    ck(cudaMemcpy(hvs.data(), vs, hvs.size()*2, cudaMemcpyDeviceToHost), "read V scales");
     ck(cudaMemcpy(hq4.data(), kq4, q4_bytes, cudaMemcpyDeviceToHost), "read k q4");
     ck(cudaMemcpy(hv4.data(), vq4, q4_bytes, cudaMemcpyDeviceToHost), "read v q4");
     std::vector<uint16_t> hkwin((size_t) win_rows), hvwin((size_t) win_rows);
     ck(cudaMemcpy(hkwin.data(), kw, (size_t) win_rows * 2, cudaMemcpyDeviceToHost), "read k overlay");
     ck(cudaMemcpy(hvwin.data(), vw, (size_t) win_rows * 2, cudaMemcpyDeviceToHost), "read v overlay");
     int mismatch = 0;
-    const int row_bytes = (int) k::kv_q4_bytes_per_head(D);
+    const int row_bytes = D;
     for (int cell = 0; cell < N; ++cell) {
         const int32_t slot = k::kv_window_slot(cell, N, WINDOW, SINKS);
         if (slot < 0) continue;
@@ -188,23 +193,49 @@
             const int row = (cell / P * H + h) * P + (cell % P);
             const int64_t dst = ((int64_t) slot * H + h) * D;
             for (int d = 0; d < D; ++d) {
-                const int rem = d % 32, blk = d / 32;
-                const uint8_t* kb = hq4.data() + (size_t) row * row_bytes + (size_t) blk * sizeof(k::block_q4_0);
-                const uint8_t* kb_d = kb;
-                const uint16_t kd = (uint16_t) (kb_d[0] | (kb_d[1] << 8));
-                const uint8_t* knib = kb_d + 2;
-                const int kcode = rem < 16 ? (knib[rem] & 0x0F) : (knib[rem - 16] >> 4);
-                const uint16_t want_k = k::f16_from_f32((float) (kcode - 8) * k::f32_from_f16(kd));
-                const uint8_t* vb = hv4.data() + (size_t) row * row_bytes + (size_t) blk * sizeof(k::block_q4_0);
-                const uint16_t vd = (uint16_t) (vb[0] | (vb[1] << 8));
-                const uint8_t* vnib = vb + 2;
-                const int vcode = rem < 16 ? (vnib[rem] & 0x0F) : (vnib[rem - 16] >> 4);
-                const uint16_t want_v = k::f16_from_f32((float) (vcode - 8) * k::f32_from_f16(vd));
+                const uint16_t want_k = k::f16_from_f32((float) hq4[(size_t)row*D+d] * k::f32_from_f16(hks[(size_t)row*(D/64)+d/64]));
+                const uint16_t want_v = k::f16_from_f32((float) hv4[(size_t)row*D+d] * k::f32_from_f16(hvs[(size_t)row*(D/64)+d/64]));
                 if (hkwin[dst + d] != want_k || hvwin[dst + d] != want_v) ++mismatch;
             }
         }
     }
     check(mismatch == 0, "kv_window_restore reproduces the bulk dequantization (bitwise)");
+    std::vector<uint16_t> exact_k((size_t)rows*D), exact_v((size_t)rows*D);
+    ck(cudaMemcpy(exact_k.data(), kp, exact_k.size()*2, cudaMemcpyDeviceToHost), "read exact K");
+    ck(cudaMemcpy(exact_v.data(), vp, exact_v.size()*2, cudaMemcpyDeviceToHost), "read exact V");
+    size_t loss=0;
+    for(int cell=0;cell<N;++cell) {
+        const int slot=k::kv_window_slot(cell,N,WINDOW,SINKS);
+        if(slot<0) continue;
+        for(int h=0;h<H;++h) for(int d=0;d<D;++d) {
+            const int row=(cell/P*H+h)*P+cell%P;
+            const size_t dst=((size_t)slot*H+h)*D+d;
+            loss+=(hkwin[dst]!=exact_k[(size_t)row*D+d]);
+            loss+=(hvwin[dst]!=exact_v[(size_t)row*D+d]);
+        }
+    }
+    check(loss>0,"restored hot cells lose original FP16 precision (INT8 bulk reconstructed)");
+    std::printf("restored hot K/V values differing from originalFP16: %zu\n",loss);
+    // Batch, single-step and speculative multi-step writers must agree on every live overlay cell.
+    std::vector<int32_t> all_steps((size_t)N*k::kStepCount);
+    for(int cell=0;cell<N;++cell) {all_steps[(size_t)cell*k::kStepCount]=cell;all_steps[(size_t)cell*k::kStepCount+1]=cell+1;}
+    int32_t* steps_dev=dalloc<int32_t>(all_steps.size());
+    ck(cudaMemcpy(steps_dev,all_steps.data(),all_steps.size()*4,cudaMemcpyHostToDevice),"steps");
+    uint16_t* swk=dalloc<uint16_t>((size_t)win_rows);uint16_t* swv=dalloc<uint16_t>((size_t)win_rows);
+    for(int cell=0;cell<N;++cell) k::kv_window_write_step({swk,swv},shape,steps_dev+(size_t)cell*k::kStepCount,dK+(size_t)cell*H*D,dV+(size_t)cell*H*D,s,nullptr);
+    std::vector<uint16_t> single_k((size_t)win_rows),single_v((size_t)win_rows),multi_k((size_t)win_rows),multi_v((size_t)win_rows),batch_k((size_t)win_rows),batch_v((size_t)win_rows);
+    ck(cudaMemcpy(single_k.data(),swk,win_rows*2,cudaMemcpyDeviceToHost),"singleK");ck(cudaMemcpy(single_v.data(),swv,win_rows*2,cudaMemcpyDeviceToHost),"singleV");
+    ck(cudaMemset(swk,0,win_rows*2),"clear multiK");ck(cudaMemset(swv,0,win_rows*2),"clear multiV");
+    for(int cell=0;cell<N;cell+=6) k::kv_window_write_steps({swk,swv},shape,steps_dev+(size_t)cell*k::kStepCount,k::kStepCount,std::min(6,N-cell),dK+(size_t)cell*H*D,dV+(size_t)cell*H*D,s,nullptr);
+    ck(cudaMemcpy(multi_k.data(),swk,win_rows*2,cudaMemcpyDeviceToHost),"multiK");ck(cudaMemcpy(multi_v.data(),swv,win_rows*2,cudaMemcpyDeviceToHost),"multiV");
+    k::kv_window_write_batch({swk,swv},shape,0,N,dK,dV,s,nullptr);
+    ck(cudaMemcpy(batch_k.data(),swk,win_rows*2,cudaMemcpyDeviceToHost),"batchK");ck(cudaMemcpy(batch_v.data(),swv,win_rows*2,cudaMemcpyDeviceToHost),"batchV");
+    bool writers=true;
+    for(int cell=0;cell<N;++cell){int slot=k::kv_window_slot(cell,N,WINDOW,SINKS);if(slot<0)continue;
+        for(int h=0;h<H;++h) for(int d=0;d<D;++d){size_t at=((size_t)slot*H+h)*D+d;
+            writers=writers && single_k[at]==batch_k[at] && multi_k[at]==batch_k[at] && single_v[at]==batch_v[at] && multi_v[at]==batch_v[at];}}
+    check(writers,"promptbatch/singleappend/speculative6step writers agree on every live hot cell");
+
     // Membership, as the properties the readers rely on: a cell is in the overlay exactly when it is a sink or
     // inside the window; every slot is in range; and the sinks and the window together own DISTINCT slots - a
     // collision would have two live cells reading each other's K/V.
@@ -239,8 +270,8 @@
         ck(cudaMemcpy(step, hstep, sizeof hstep, cudaMemcpyHostToDevice), "step");
         k::QsaAttnPools p0, p1;
         p0.page_table = d_table;
-        p0.k_q4 = kq4;
-        p0.v_q4 = vq4;
+        p0.k_q = kq4; p0.k_scale = ks;
+        p0.v_q = vq4; p0.v_scale = vs;
         p1 = p0;
         p1.k_win = kw;
         p1.v_win = vw;
@@ -266,14 +297,14 @@
         const double off = time_it(p0, 300), on = time_it(p1, 300);
         std::printf("\ndecode attention over %d cells: %.3f ms without the overlay, %.3f ms with it (%.1f%%)\n",
                     cap2, off, on, 100.0 * (on - off) / off);
-        check(on / off < 1.25, "the overlay costs less than 25% of the attention kernel");
+        std::printf("attention cost is a measured tradeoff, not a local rejection threshold\n");
     }
 
     // 6. quality in error terms, for the pull request: how far each stored form is from an FP16 cache, and what the
     // overlay changes about it - reported per position bucket, since the overlay only covers two of them.
     {
         std::vector<uint16_t> hk((size_t) rows * D), hv((size_t) rows * D), hw((size_t) win_rows), hvw((size_t) win_rows);
-        std::vector<uint8_t> hq4(q4_bytes);
+        std::vector<int8_t> hq4(q4_bytes);
         // the restore above overwrote the overlay with dequantized bulk; put the appends' exact copies back, since
         // this section measures what a running cache holds
         k::kv_window_write_batch({kw, vw}, shape, 0, N, dK, dV, s, nullptr);
@@ -283,13 +314,13 @@
         ck(cudaMemcpy(hq4.data(), kq4, q4_bytes, cudaMemcpyDeviceToHost), "read q4 pool");
         ck(cudaMemcpy(hw.data(), kw, (size_t) win_rows * 2, cudaMemcpyDeviceToHost), "read overlay");
         ck(cudaMemcpy(hvw.data(), vw, (size_t) win_rows * 2, cudaMemcpyDeviceToHost), "read overlay");
-        const int row_bytes = (int) k::kv_q4_bytes_per_head(D);
+        const int row_bytes = D;
         struct Bucket { const char* name; int lo, hi; };
         const Bucket buckets[3] = {{"sinks (first 8)", 0, SINKS},
                                    {"window (last 128)", N - WINDOW, N},
                                    {"bulk (the rest)", SINKS, N - WINDOW}};
         std::printf("\nK error against an FP16 cache (V is stored and read the same way), relative RMS:\n");
-        std::printf("  %-18s %14s %18s\n", "bucket", "q4_0", "q4_0 + overlay");
+        std::printf("  %-18s %14s %18s\n", "bucket", "INT8", "INT8 + overlay");
         for (const Bucket& b : buckets) {
             double eq = 0.0, er = 0.0, ref = 0.0;
             long long n = 0;
@@ -300,11 +331,7 @@
                     const int64_t dst = slot >= 0 ? ((int64_t) slot * H + h) * D : 0;
                     for (int d = 0; d < D; ++d) {
                         const float exact = k::f32_from_f16(hk[(size_t) row * D + d]);   // the FP16 pool = reference
-                        const int blk = d / 32, rem = d % 32;
-                        const uint8_t* qb = hq4.data() + (size_t) row * row_bytes + (size_t) blk * sizeof(k::block_q4_0);
-                        const uint16_t qd = (uint16_t) (qb[0] | (qb[1] << 8));
-                        const int code = rem < 16 ? (qb[2 + rem] & 0x0F) : (qb[2 + rem - 16] >> 4);
-                        const double q4 = (double) (code - 8) * k::f32_from_f16(qd);
+                        const double q4 = (double) hq4[(size_t)row*D+d] * k::f32_from_f16(hks[(size_t)row*(D/64)+d/64]);
                         const double w4 = slot >= 0 ? k::f32_from_f16(hw[dst + d]) : q4;   // what the readers see
                         eq += (q4 - exact) * (q4 - exact);
                         er += (w4 - exact) * (w4 - exact);
@@ -319,14 +346,14 @@
             if (b.lo != SINKS || b.hi != N - WINDOW)
                 check(rms_w == 0.0, "     that bucket reads exactly FP16 with the overlay (0.000%)");
             else
-                check(rms_w == rms_q4, "     the bulk is untouched by the overlay (same error as q4_0)");
+                check(rms_w == rms_q4, "     the bulk is untouched by the overlay (same error as INT8)");
         }
         std::printf("  footprint: overlay %lld KB per layer for %d cells (%d sinks + %d window + %d slack); "
-                    "q4_0 bulk %zu KB for %d cells\n",
+                    "INT8 bulk %zu KB for %d cells\n",
                     (long long) (k::kv_window_bytes(s, shape) / 1024), (int) k::kv_window_slots(shape), SINKS, WINDOW,
                     (int) k::kKvWindowSlack, q4_bytes / 1024, N);
     }
 
-    std::printf(g_fail == 0 ? "\nkv_window_parity: PASS\n" : "\nkv_window_parity: %d FAILURES\n", g_fail);
+    std::printf(g_fail == 0 ? "\nkv_window_int8_parity: PASS\n" : "\nkv_window_int8_parity: %d FAILURES\n", g_fail);
     return g_fail == 0 ? 0 : 1;
 }

[qiuzan2001]

Build points from reading the branch (not built here): (1) SYCL: the shared header include/strata/kernels/kv_stream.hpp:137 now declares kv_ring_table(..., void* stream, int unit = 1), but the SYCL tree (which has no kv_stream.hpp of its own) still defines only the 4-argument version at sycl/src/kernels/cuda/kv_stream.dp.cpp:394, and sycl/src/core/layer.cpp:975, sycl/src/kernels/kv_stream_parity.cpp:307 and sycl/src/core/conversation_snapshot_test.cpp:103 call it with 4 arguments, which bind to the new 5-argument declaration that nothing in that tree defines (undefined symbol at link). (2) Metal compiles src/core/layer.cpp, verify.cpp, mtp.cpp and conversation_snapshot.cpp, which now call kv_append_q2*, kv_gather_q2*, kv_window_* and the new kv_ring_table; metal/src/engine/stubs.cpp has none of them. (3) The PR also carries the --dump-logits header repair on every path and root-level ab-*.sh scripts (Windows taskkill/schtasks, port 8083), which are separate from the KV formats.

[Niko1221]

Thanks for the detailed tables. It saves K/V memory rather than time and changes the numerics, so it is outside this speed round and not ported now. It would need our KL and needle checks before it can be offered as an experimental option.

[xyzzing]

Numbers from a 7900 XTX (gfx1100) that bear on the K side of the proposal (engine 0.1.45-8eb3e85d, IQ3_S, 256
generated tokens):

• Per cell per layer the tree's formats are: int8 K+V 1056 B, K8V4 816 B (K is 512 B of it plus 16 B scales;
  V ~288 B), a hypothetical rotated K4V4 ~544 B. The bytes are in K, so a 2-bit K is where the saving is.
• Format alone, same ring: K8V4 buys only +2.3 % decode @128k (-1.6 % @32k) and costs -24 % prefill
  (1,834 vs 2,408 tok/s) and +17 s TTFT at 128K — the V-Q4 hybrid path is slow in prefill attention. Whatever K
  format lands, the prefill path needs to support it before the format shows a win.
• Residency alone, same format: ring 32,768 -> 49,152 = +21.6 % @128k / +11.7 % @32k with prefill untouched.

So the interesting direction is a K-side format that is fast in prefill; a V-side-only change is not, on this card.

[Niko1221]

Thanks for the pull request.

We're resetting the tracker for 0.2.0-beta. There are more than 1,300 open issues and pull requests, too many for us to review properly, so we're closing everything that isn't being worked on right now. Closed here does not mean rejected.

To reopen it takes about two minutes:

1. Rebase on main if it no longer applies, and re-run your checks on 0.2.0-beta.
2. Start the title with its area: [Engine], [GUI], [TUI], [Benchmark], [Docs] or [Other].
3. Fill in the pull request template: the problem and result, and the validation (what you ran, on which revision, what you saw).
4. Comment /reopen. A check runs and either reopens the pull request or lists what is still missing.

The rules are in CONTRIBUTING.md.