Metadata-Version: 2.4
Name: turbo-attn
Version: 0.67.0
Summary: Optimized CUDAgraph-enabled kernels and attention backend for vLLM, SGLang and more based on TurboQuant near-lossless KV cache compression.
Author-email: "Dmitri Evseev (Arbi City)" <dmitri.evseev@arbi.city>
Maintainer-email: "Dmitri Evseev (Arbi City)" <dmitri.evseev@arbi.city>
License-Expression: Apache-2.0
Project-URL: Homepage, https://github.com/arbicity/turbo-attn
Project-URL: Repository, https://github.com/arbicity/turbo-attn
Project-URL: Documentation, https://github.com/arbicity/turbo-attn/tree/main/docs
Project-URL: Issues, https://github.com/arbicity/turbo-attn/issues
Project-URL: Changelog, https://github.com/arbicity/turbo-attn/releases
Keywords: kv-cache,quantization,vllm,sglang,flash-attention,llm-inference,transformers,cuda
Classifier: Development Status :: 4 - Beta
Classifier: Intended Audience :: Developers
Classifier: Intended Audience :: Science/Research
Classifier: Environment :: GPU :: NVIDIA CUDA
Classifier: Operating System :: POSIX :: Linux
Classifier: Programming Language :: Python :: 3
Classifier: Programming Language :: Python :: 3 :: Only
Classifier: Programming Language :: Python :: 3.10
Classifier: Programming Language :: Python :: 3.11
Classifier: Programming Language :: Python :: 3.12
Classifier: Typing :: Typed
Classifier: Topic :: Scientific/Engineering :: Artificial Intelligence
Classifier: Topic :: Software Development :: Libraries :: Python Modules
Requires-Python: >=3.10
Description-Content-Type: text/markdown
License-File: LICENSE
License-File: NOTICE
Requires-Dist: torch>=2.1
Requires-Dist: nvidia-cutlass-dsl[cu13]>=4.6.0
Requires-Dist: quack-kernels>=0.6
Requires-Dist: apache-tvm-ffi>=0.1.11
Provides-Extra: codebook
Requires-Dist: scipy>=1.10; extra == "codebook"
Provides-Extra: store
Requires-Dist: cryptography>=41; extra == "store"
Provides-Extra: triton
Requires-Dist: triton>=3.0; extra == "triton"
Provides-Extra: vllm
Requires-Dist: vllm>=0.19; extra == "vllm"
Provides-Extra: flash-attn
Requires-Dist: flash-attn>=2.5; extra == "flash-attn"
Provides-Extra: fa4
Requires-Dist: flash-attn-4==4.0.0b24; extra == "fa4"
Provides-Extra: eval
Requires-Dist: lm-eval>=0.4.5; extra == "eval"
Requires-Dist: ray; extra == "eval"
Requires-Dist: datasets; extra == "eval"
Requires-Dist: langdetect; extra == "eval"
Requires-Dist: immutabledict; extra == "eval"
Requires-Dist: nltk; extra == "eval"
Requires-Dist: sacrebleu; extra == "eval"
Requires-Dist: absl-py; extra == "eval"
Provides-Extra: calibration-data
Requires-Dist: datasets>=4.0; extra == "calibration-data"
Provides-Extra: calibrate-hf
Requires-Dist: transformers>=4.56; extra == "calibrate-hf"
Provides-Extra: dev
Requires-Dist: import-linter>=2.0; extra == "dev"
Requires-Dist: pyright>=1.1.380; extra == "dev"
Requires-Dist: pytest>=7.0; extra == "dev"
Requires-Dist: pytest-xdist>=3.0; extra == "dev"
Requires-Dist: pytest-timeout>=2.1; extra == "dev"
Requires-Dist: scipy>=1.10; extra == "dev"
Requires-Dist: triton>=3.0; extra == "dev"
Provides-Extra: all
Requires-Dist: turbo-attn[codebook,eval,flash-attn,store,triton,vllm]; extra == "all"
Dynamic: license-file

# Turbo Attention (turbo-attn)

[![CI](https://github.com/arbicity/turbo-attn/actions/workflows/ci.yml/badge.svg)](https://github.com/arbicity/turbo-attn/actions/workflows/ci.yml)
[![PyPI](https://img.shields.io/pypi/v/turbo-attn.svg)](https://pypi.org/project/turbo-attn/)
[![License](https://img.shields.io/badge/license-Apache--2.0-blue.svg)](LICENSE)

turbo-attn (`tkv`) compresses the KV cache to 2–8 bits so long context and
high concurrency fit on NVIDIA RTX GPUs, serving attention straight from the
compressed cache with its own kernels.

## Install and try it (no GPU needed)

```bash
pip install turbo-attn
```

```python
import torch
from tkv import TkvCodec

codec = TkvCodec(head_dim=128, bit_width=4)  # pure PyTorch: runs on CPU
keys = torch.randn(8, 128, generator=torch.manual_seed(0))  # 8 key vectors
packed, norms = codec.compress_k(keys)  # 4-bit codes plus one norm per vector
print(f"{codec.compression_ratio:.2f}x smaller than bf16, relative error {(codec.decompress_k(packed, norms) - keys).norm() / keys.norm():.2f}")
```

```text
3.88x smaller than bf16, relative error 0.10
```

Check what a machine can do with `tkv doctor` (it works without a GPU and exits
`1` when something that stops tkv from serving is wrong), and see every setting
with `tkv config --doc`:

```bash
tkv doctor
tkv config --check      # unknown or malformed TKV_* variables, with suggestions
tkv inspect-store "$(python -c 'import tkv.calibrations as c; print(c.path("qwen3.5-0.8b"))')"   # what a calibration store holds
tkv warmup --head-dim 128 --num-kv-heads 2 --num-heads 16   # warm the decode autotune on this GPU
```

From Python, `tkv.configure(bits=4)` and `with tkv.override(bits="k4v5"):` set flags
without touching `os.environ`, and refuse a bad value with the flag's name.

The wheel is pure Python, and `pip install turbo-attn` brings the CuTeDSL kernel
stack (`nvidia-cutlass-dsl`, `quack-kernels`, `apache-tvm-ffi`) that every GPU path
needs (there is no separate `cute` extra). Extras are optional; quote them in a shell:

| Extra | Adds | For |
|---|---|---|
| `store` | `cryptography` | verifying prebuilt kernels from the kernel store |
| `calibrate-hf` | `transformers` | calibrating your own model without an engine |
| `codebook` | `scipy` | fitting codebooks |
| `triton` | `triton` | the Triton compress path |
| `vllm`, `flash-attn` | upstream vLLM, flash-attn | running next to upstream vLLM, the flash-attn interop example |
| `eval` | lm-eval and friends | the accuracy harnesses |
| `all` | `vllm`, `flash-attn`, `triton`, `codebook`, `eval`, `store` | everything except `calibrate-hf` and the benchmark-only `fa4` |

Pin with `turbo-attn~=0.65` (see
[`docs/compatibility.md`](docs/compatibility.md)). More runnable code is in
[`examples/`](examples/); the supported API is in
[`docs/public_api.md`](docs/public_api.md), with [every public name](docs/api_reference.md),
[every `TKV_*` flag](docs/flags.md) and the
[engine API version history](docs/public_api.md#api-version-history).

## Why

- **Most faithful 4-bit KV cache in vLLM:** 98.92% of next tokens match FP32,
  vs 97.63% (`int4_per_token_head`) and 97.20% (`turboquant_4bit_nc`). The
  figures come from the fidelity ladder report on an external host (linked
  below, not stored in this repository), so they are not reproducible from the
  repository alone; the harness that produced them is
  [`benchmarks/eval/paper_kv_ablation/`](benchmarks/eval/paper_kv_ablation/).
  "Match" is 1 minus the total-variation distance under the served sampler,
  first position excluded. The ladder scores against an HF fp32 oracle; the
  release gate ([`tests/test_fidelity_gate_gpu.py`](tests/test_fidelity_gate_gpu.py))
  scores against the same engine's bf16 KV instead, so the two figures are not
  interchangeable.
  [Fidelity ladder](https://arbi-tkv-reports.s3.uk.io.cloud.ovh.net/kv-fidelity-ladder/latest.html)
- **3.87× smaller KV cache than bf16** at 4 bits. That ratio is slot-size
  arithmetic (4.13 bits/number on MiniCPM5-2B); the pools the engines actually
  allocated hold 3.7–3.8× the tokens
  ([record](benchmarks/baselines/minicpm5-2b/serve_matrix.md)). FP8-level
  fidelity at 5 bits.
- **No loss on hard reasoning:** GPQA Diamond on MiniCPM5-2B, bf16 68.9% vs
  tkv K4V4 67.5% / K5V5 68.2%, every difference within noise.
  [GPQA report](https://arbi-tkv-reports.s3.uk.io.cloud.ovh.net/benchmarks/gpqa-diamond/minicpm5-2b/index.html)
- **Long prompts read faster:** 126K-token prompt on one RTX 4090, first token
  in 13.7 s vs 22.3 s with bf16 (vLLM). This speed comes from the int8
  attention route (`TKV_PREFILL_INT8`, on by default), which quantizes Q and P
  to int8 and has a measured fidelity cost: see
  [Precision levers](#precision-levers) and
  [`docs/records/default-on-off-priced.md`](docs/records/default-on-off-priced.md).
  [Record](benchmarks/baselines/minicpm5-2b/serve_matrix.md)

## Serve with a GPU engine

Needs one NVIDIA GPU of compute capability 8.x or 12.0 (RTX 30/40/50-series,
A-series, L40). One setting: `TKV_BITS` (`4` = K4V4, `k5v4`, …). A served model is
always calibrated: from `TKV_CALIBRATION_FILE` (the model's calibration store), else
the store that ships for it, else one fitted from the live model at boot; a boot that
can do none of these is refused, and `TKV_CALIBRATION=none` is the explicit way to
serve uncalibrated. The wheel ships the
stores for the models we have calibrated: `tkv.calibrations.names()` lists them
and `tkv.calibrations.path("qwen3.5-0.8b")` is the file to give
`TKV_CALIBRATION_FILE` (`gpt-oss-20b` is one of them). `python -m tkv.auto_calibrate` makes one for any other
model ([details](#calibration)). The Docker commands below need a clone of the
repository for the compose files.

```bash
mkdir tkv && cd tkv
git clone https://github.com/arbicity/turbo-attn
git clone https://github.com/arbicity/vllm-turbo   vllm-fork     # for vLLM
git clone https://github.com/arbicity/sglang-turbo sglang-fork   # for SGLang
mkdir -p ~/models
hf download Qwen/Qwen3.5-0.8B --local-dir ~/models/Qwen3.5-0.8B
cp turbo-attn/calibrations/qwen3.5-0.8b_v59.json ~/models/
cd turbo-attn/docker
```

**vLLM** ([vllm-turbo](https://github.com/arbicity/vllm-turbo)):

```bash
TKV_MODELS_ROOT=~/models TKV_CALIBRATION_FILE=/models/qwen3.5-0.8b_v59.json TKV_CUDA_ARCH=8.9 \
  docker compose -f compose.vllm.yaml up -d        # OpenAI API on :8000
```

or, where vllm-turbo and turbo-attn are installed:

```bash
VLLM_PLUGINS=tkv TKV_BITS=4 TKV_CALIBRATION_FILE="$(python -c 'import tkv.calibrations as c; print(c.path("minicpm5-2b"))')" \
  vllm serve openbmb/MiniCPM5-2B --kv-cache-dtype tkv --attention-backend TURBO_ATTN
```

**SGLang** ([sglang-turbo](https://github.com/arbicity/sglang-turbo)):

```bash
TKV_MODELS_ROOT=~/models TKV_CALIBRATION_FILE=/models/qwen3.5-0.8b_v59.json TKV_CUDA_ARCH=8.9 \
  docker compose -f compose.sglang.yaml up -d      # OpenAI API on :30000; flag: --kv-cache-dtype tkv
```

**arbi-serve** ([PyPI](https://pypi.org/project/arbi-serve/)):

```bash
uvx --python 3.12 arbi-serve run Qwen/Qwen3.5-0.8B
```

`TKV_CUDA_ARCH`: `8.6` RTX 30, `8.9` RTX 40/L40, `12.0` RTX 50. All knobs:
[`compose.vllm.yaml`](docker/compose.vllm.yaml),
[`compose.sglang.yaml`](docker/compose.sglang.yaml).

## More

- [Reports](https://arbi-tkv-reports.s3.uk.io.cloud.ovh.net/index.html): fidelity ladder, GPQA Diamond
- [Details](#details): codec, calibration, kernels, engines and hardware, flags, precision levers, benchmarks
- [`docs/`](docs/) · [`examples/`](examples/) · [`docs/public_api.md`](docs/public_api.md)

---

## Details

### First boot

A cold machine compiles the kernels its model needs (nvcc, cached in
`TORCH_EXTENSIONS_DIR`) and times the decode kernel variants once per KV shape
(cached in `TKV_CACHE_DIR`, default `~/.cache/tkv`). The compose files bake
kernels into the image (the compose files pass `TKV_BAKE=1`; the
Dockerfile's own default is `TKV_BAKE=0`, so a plain `docker build` does not
bake) and keep both caches on named volumes,
so later boots skip both steps. The boot log prints a line per compiled kernel
and per tuned shape. Set `TKV_NO_JIT=1` to refuse any compile at serve time.

### The codec

| Piece | What it does | Default |
|---|---|---|
| Widths | K and V each 2–8 bits, independently (`TKV_BITS=4`, `k5v4`). | `TKV_BITS` unset = 4 |
| Hadamard rotation | A Walsh–Hadamard transform spreads outlier channels across the vector. Queries are rotated the same way, so scores need no un-rotation. | always |
| OSCAR rotation | A rotation fitted per layer and KV head on top of Hadamard ([`tkv/oscar.py`](tkv/oscar.py), fitted by [`tkv/calibration/oscar_rotation.py`](tkv/calibration/oscar_rotation.py)). | served when the store's basis is OSCAR; both shipped stores are. Labelled experimental and opt-in in `tkv/oscar.py` (after OSCAR, arXiv 2605.17757); no Hadamard-versus-OSCAR A/B is committed |
| Per-channel scales | Calibrated scales fill the codebook's range per channel ([`tkv/codec/_scales.py`](tkv/codec/_scales.py)). | on (`TKV_DISABLE_PRESCALE=1` is an ablation) |
| Norms | Each vector's length is stored once in bf16 and only its direction is quantized ([`tkv/codec/_norm.py`](tkv/codec/_norm.py)). Key-side norm correction keeps the decoded key's length. | on (`TKV_K_NC`) |
| vq2 | At 2–4 bits, channel pairs are quantized together against a 2-D codebook of 4^bits entries. | when the store has a vq2 entry |
| Lloyd codebooks | Scalar Lloyd–Max levels fitted per layer and width. Serve 5–8 bits, and 2–4 when there is no vq2 entry. | from the store |
| Fresh diagonal | While a prompt is read, its newest chunk attends to its own fresh values; only older context is read from the codes ([`tkv/runtime/_fresh_diagonal.py`](tkv/runtime/_fresh_diagonal.py)). | on (`TKV_FRESH_DIAGONAL`) |
| Recent ring | Each sequence's newest 128 positions (64 when that would take over 2% of KV memory) are also kept at 8 bits and read from there, for every sequence the engine can run, on layers coarser than 8 bits ([`docs/recent-ring.md`](docs/recent-ring.md)). | on (`TKV_RECENT_RING`; `TKV_RECENT_RING_W` pins the window) |
| SmartMix | A fractional `TKV_BITS` serves the per-layer `(k_bits, v_bits)` allocation the store recorded for that budget, raced against the best uniform width ([`tkv/runtime/smartmix.py`](tkv/runtime/smartmix.py)). The kernels read the mixed-width cache in one forward pass. | only for a fractional `TKV_BITS` |

### Calibration

A calibration store is one JSON file per model, schema version **59**
(`SCHEMA_VERSION` in [`tkv/bundle/schema.py`](tkv/bundle/schema.py)).
The loader accepts exactly that version. Its top-level keys are
`schema_version`, `model`, `bases` (the rotations, keyed `hadamard` or
`oscar:<hash>`), `calibrations` (one entry per `<basis>/lloyd` and
`<basis>/vq2`, with per-layer codebooks and scales), and optional
`allocations` (the SmartMix proxy table and race verdicts). When a store holds
several bases with calibrations, choose one with `TKV_CALIBRATION_BASIS`.

```bash
pip install "turbo-attn[calibrate-hf]"
python -m tkv.auto_calibrate --model /path/to/model --output mymodel_v59.json
# add --smartmix --fit-target ... --race-target ... [--budgets ...] to also race fractional budgets
python -m tkv.auto_calibrate --help     # every option
```

[`calibrations/README.md`](calibrations/README.md) documents the shipped stores (also in the wheel, see `tkv.calibrations`)
and the full recipe (OSCAR rotation, then `auto_calibrate --rotation`, then vq2).
[`docs/calibration_data.md`](docs/calibration_data.md) covers the calibration
corpus. With vLLM, `TKV_AUTO_CALIBRATE_MODEL` fits a missing store at boot
(codebooks only, no SmartMix).

### Kernels

All of these are turbo-attn's own kernels. They read the compressed cache
directly and run under CUDA graph capture.

| Kernel | Role | Source |
|---|---|---|
| `turbo_attn_int8` | Default prefill for the compressed KV cache at head_dim 128 and 256: int8 Q·Kᵀ with exact int32 accumulation and byte-LUT K/V decode (`TKV_PREFILL_INT8=1`). | [`tkv/kernels/cuda/prefill/turbo_attn_int8.py`](tkv/kernels/cuda/prefill/turbo_attn_int8.py) |
| `TurboAttnCute` | CuTeDSL tensor-core prefill. Serves the bf16 first-chunk bypass, `tkv-bypass`, and the shapes int8 declines (head_dim 64 and 512). | [`tkv/kernels/cuda/prefill/turbo_attn_cute.py`](tkv/kernels/cuda/prefill/turbo_attn_cute.py), [`docs/prefill-path-selection.md`](docs/prefill-path-selection.md) |
| Split-KV prefill | Splits a prefill launch that under-fills the GPU (`TKV_PREFILL_INT8_SPLIT`, `TKV_PREFILL_CUTE_SPLIT`, both `auto`). Eager launches only. | [`docs/records/prefill-split-kv.md`](docs/records/prefill-split-kv.md) |
| `turbo_attn_simt` decode | One split-K paged decode kernel. Its tensor-core (`mma.sync`) body is the default; the autotuner times it against the SIMT body per shape. Policies cover full attention, sliding window, MLA, latent MQA and MTP verify. | [`tkv/kernels/_turbo_attn_simt.cu`](tkv/kernels/_turbo_attn_simt.cu), [`docs/decode-kernel-design.md`](docs/decode-kernel-design.md) |
| Fused compress-store | Encodes K/V (Lloyd or vq2) and writes the paged cache in one launch. | [`tkv/kernels/cuda_compress_store.py`](tkv/kernels/cuda_compress_store.py) |

Measured kernel records:

- int8 prefill route against the CuTeDSL route on an RTX 4090, 27B-class geometry
  (24 query / 4 KV heads, head_dim 256, K4V4): **1.08–2.71×** faster across
  the prefill cells measured on 2026-09-21; 1.58–2.71× at 512 query tokens or
  more ([`docs/records/ncu-profile-2026-09-21.md`](docs/records/ncu-profile-2026-09-21.md)).
  The kernel has changed since and has not been re-measured against CuTeDSL.
- The int8 prefill kernel reaches 94% of its own structural throughput
  ceiling, 357 of 381 int8 TOPS
  ([`docs/records/where-the-105-tops-went.md`](docs/records/where-the-105-tops-went.md)).
- Prefill throughput on the served chunk compared with dense PyTorch flash SDPA:
  [`docs/prefill-path-selection.md`](docs/prefill-path-selection.md).

The kernels are policy-parametric on the KV format through the Loader
protocol, so a new format is a sibling module rather than a mainloop change
([`docs/writing_a_loader.md`](docs/writing_a_loader.md)). The codec also works
with third-party attention by decompressing to bf16
([`docs/codec_third_party_interop.md`](docs/codec_third_party_interop.md)).

### Engines and hardware

| Engine | How tkv attaches |
|---|---|
| vLLM (v0.31.0, [vllm-turbo](https://github.com/arbicity/vllm-turbo)) | `vllm.general_plugins` entry point `tkv`; `--kv-cache-dtype tkv` or `tkv-bypass`, backend `TURBO_ATTN`. |
| SGLang (v0.5.21, [sglang-turbo](https://github.com/arbicity/sglang-turbo)) | `tkv.integrations.sglang.register()` in the server process (the image's entrypoint does it); `--kv-cache-dtype tkv`, backend `turbo-attn`. |
| arbi-serve | Uses tkv's public runtime API directly; `--kv-cache-dtype tkv`. |

Why forks: neither upstream vLLM nor upstream SGLang lets a plugin register a
new KV cache dtype and attention backend at runtime yet. Both forks are thin
overlays on tagged releases ([`docker/PATCHES.md`](docker/PATCHES.md)).

- **GPUs:** compute capability 8.x (Ampere, Ada) and 12.0 (consumer
  Blackwell). Hopper (9.x) and datacenter Blackwell (10.x) are refused by
  prefill. See [`docs/support_matrix.md`](docs/support_matrix.md).
- **Head dims:** decode 64, 128, 256, 512; int8 prefill 128 and 256; the
  CuTeDSL prefill covers the rest. A head that is not a power of two (80, 96,
  112, 192) is stored and attended at the next power of two on every path,
  paying +60% KV bytes at 80 and +33% at 96 and 192
  ([cost table](docs/support_matrix.md#padded-head-dims-cost)).
- **Attention shapes:** GQA, sliding window, MTP speculative verify. MLA
  compression is off by default and opt-in (`TKV_MLA_ENABLE=1`);
  `tkv-bypass` serves MLA models on the model's own latent.

### Configuration

All runtime configuration is `TKV_*` environment variables, declared in one
registry ([`tkv/config.py`](tkv/config.py) and the `tkv/_config_*.py` modules
it assembles). Print every flag with its type, default and description:

```bash
tkv config --doc        # every flag: type, default, stability, description
tkv config --json       # the same, machine-readable, with the values set here
tkv config --check      # exit 1 on an unknown TKV_* name or a malformed value
```

An engine refuses to start when the environment sets a `TKV_*` name that is not
in the registry, and a malformed value such as `TKV_BITS=banana` fails as soon
as the flags are read; both errors name what is accepted, and an unknown name
comes with a did-you-mean suggestion. In your own Python process nothing checks
the environment until you ask: run `tkv config --check`, or call
`tkv.config.refuse_unknown_tkv_envs()`. The flags that matter most:

| Flag | Default | Meaning |
|---|---|---|
| `TKV_BITS` | unset = `4` | Int: uniform K=V. `k<K>v<V>`: uniform, K and V separate. Float: SmartMix budget (needs a store verdict). |
| `TKV_CALIBRATION_FILE` | unset | The model's calibration store. Unset: the shipped store for the model, else one fitted in process. |
| `TKV_CALIBRATION` | unset | `none`: serve the uncalibrated codec expressly (warns at every boot). |
| `TKV_CALIBRATION_BASIS` | unset | Which basis to serve when the store has several. |
| `TKV_AUTO_CALIBRATE_MODEL` | unset | The model to fit from when not the served one. |
| `TKV_AUTO_CALIBRATE` | `1` | `0` disables the in-process fit; a boot that then has no store is refused. |
| `TKV_PREFILL_INT8` | on | int8 prefill kernel; off routes prefill to CuTeDSL. |
| `TKV_FIRST_CHUNK_RAW` | on | First prompt chunk attends in raw bf16. |
| `TKV_FRESH_DIAGONAL` | on | The current prefill chunk reads its own fresh values. |
| `TKV_RECENT_RING` / `TKV_RECENT_RING_W` | on / auto (128, or 64 past 2% of KV) | 8-bit ring of each sequence's newest positions. |
| `TKV_MMA_PRECISION` | `fast` | Tensor-core decode arithmetic, `fast` or `exact`. |
| `TKV_DECODE_MMA` | auto | Tensor-core decode body; `1` pins it, `0` (deprecated) forces SIMT. |
| `TKV_FUSE_QROT` / `TKV_O_PROJ_FOLD` | on / on | Fuse the query rotation (off by itself on OSCAR layers); fold the output projection into the attention output. |
| `TKV_NO_JIT` | off | Refuse to compile kernels at serve time. |
| `TKV_CACHE_DIR` | `~/.cache/tkv` | Autotune tables and tkv's user cache. |
| `TKV_KERNEL_STORE_URL` | public store | Prebuilt kernel store; `off` disables it. |
| `TKV_MLA_ENABLE` | off | Compress MLA layers. |

### Precision levers

Each default below trades a little arithmetic precision for speed. The KV
fidelity ladder re-measures every one of them on each full run, as a paired
ablation of the served K4V4 row (one `TKV_*` flag changed, same prompts):

| Lever | Default | Ablation row |
|---|---|---|
| int8 prefill kernel | `TKV_PREFILL_INT8=1` | `K4V4~int8off` |
| fresh prefill chunk quantized to s8 inside int8 prefill | `TKV_PREFILL_INT8_FRESH_S8=1` | `K4V4~freshbf16` |
| tensor-core decode arithmetic | `TKV_MMA_PRECISION=fast` | `K4V4~mmaexact` |
| decode SMEM tile in fp16 | `TKV_DECODE_SMEM_HALF=1` | `K4V4~smemfp32` |

The rows are defined in
[`benchmarks/eval/paper_kv_ablation/config.json`](benchmarks/eval/paper_kv_ablation/config.json)
(`ablations`). On MiniCPM5-2B end to end, no int8 prefill cost resolves:
1 − TVD against FP32 is **0.97839 with it on and off** (32 prompts, 60,518
decode positions; paired 32-prompt run, 2026-10-05), which says the cost is
below what that sample can see, not that it is zero. Elsewhere it is
measured: at the attention-kernel level the int8 route has 2.6 to 3.0 times
the attention-arithmetic error of the bf16 route
([`docs/records/default-on-off-priced.md`](docs/records/default-on-off-priced.md)),
and the records on larger, weight-quantized checkpoints
([`cross-engine-tip-rescore.md`](docs/records/cross-engine-tip-rescore.md),
[`pdown-quad-gate.md`](docs/records/pdown-quad-gate.md)) and the Qwen3.5-0.8B
record
([`p-scale-levers-what-is-reachable.md`](docs/records/p-scale-levers-what-is-reachable.md))
score the int8-off arm better. `TKV_PREFILL_INT8=0` buys that fidelity back at
the price of the prefill speed above. On GPQA Diamond, K4V4
with int8 prefill off scored 70.0% against 67.5% with it on, a difference whose
95% range includes zero ([GPQA report](https://arbi-tkv-reports.s3.uk.io.cloud.ovh.net/benchmarks/gpqa-diamond/minicpm5-2b/index.html)).

### What is not claimed

Decode speed is measured only on MiniCPM5-2B on one RTX 4090: in
[`benchmarks/baselines/minicpm5-2b/serve_matrix.md`](benchmarks/baselines/minicpm5-2b/serve_matrix.md)
(2026-10-06, mixed request lengths), tkv decoded 11–64% faster per stream than bf16
on vLLM, SGLang and arbi-serve from c1 to c256 (134 against 81.6 tok/s in vLLM at
126K context), and other
models and GPUs are not measured. No committed record
compares turbo-attn's speed with other KV cache quantization backends.

### Benchmarks and how to reproduce them

- **KV fidelity ladder:**
  [`benchmarks/eval/paper_kv_ablation/`](benchmarks/eval/paper_kv_ablation/).
  - What one row is: one engine boot per KV cache format, teacher-forced
    against a fixed target and scored against an FP32 oracle.
    - The primary metric is fidelity = 1 − TVD under the served sampler.
    - Bits and bytes per token are read off the live engine.
  - What it covers: every tkv width plus every vLLM-native KV quantization,
    the lever ablations above, and task benchmarks (GPQA Diamond, LongBench v2,
    LiveCodeBench) served with the model card's sampling (`sampling_by_model`
    in `config.json`).
  - Models: `config.json` is MiniCPM5-2B; the published Qwen3.5-0.8B ladder
    uses `config_qwen35_08b.json`.
  - How it runs: CI drives it through
    [`kv-ablation-ladder.yml`](.github/workflows/kv-ablation-ladder.yml) and
    its nightly. The harness README has the gates and the manual runbook.
- **Serving matrix:** [`benchmarks/matrix_a/`](benchmarks/matrix_a/) and
  [`matrix-a.yml`](.github/workflows/matrix-a.yml) measure TTFT, TPOT and
  throughput for vLLM, SGLang and arbi-serve, tkv against bf16. Baselines are
  in [`benchmarks/baselines/`](benchmarks/baselines/).
- **Kernel benchmarks** live in [`benchmarks/`](benchmarks/README.md), and
  the measurement records behind kernel decisions are in
  [`docs/records/`](docs/records/README.md). Methodology notes:
  [`docs/BENCHMARKING.md`](docs/BENCHMARKING.md).

Published reports: [index](https://arbi-tkv-reports.s3.uk.io.cloud.ovh.net/index.html).

### Kernel distribution

Kernels resolve fastest-first: kernels baked into the image, then the local
cache, then the kernel store, then a local nvcc compile. The kernel store is a
public, content-addressed object store whose manifests are signed with
ed25519. An object is used only if a trusted signed manifest lists its sha256,
and verification needs `pip install "turbo-attn[store]"`. Any failure is a
miss that falls back to compiling. The same store serves arbi-serve. The store
is new and is still being populated, so today most kernels are compiled
locally or baked into the image. See
[`docs/kernel-store.md`](docs/kernel-store.md). Set
`TKV_KERNEL_STORE_URL=off` to never contact it.

### Repository map

- `tkv/codec/`, `tkv/core.py`: the codec.
- `tkv/kernels/`: prefill, decode, compress-store and the kernel store client.
- `tkv/runtime/`: calibration loading, SmartMix, cache layout, autotune.
- `tkv/integrations/`: vLLM and SGLang.
- `tkv/calibration/`, `tkv/auto_calibrate.py`: calibration pipeline.
- `docker/`: engine images and compose files ([`docker/README.md`](docker/README.md)).
- More: [`ARCHITECTURE.md`](ARCHITECTURE.md), [`docs/README.md`](docs/README.md),
  [`CONTRIBUTING.md`](CONTRIBUTING.md).

### License and citation

Apache License 2.0 ([`LICENSE`](LICENSE), [`NOTICE`](NOTICE)). Vendored
third-party code and technique provenance are listed in
[`ATTRIBUTION.md`](ATTRIBUTION.md). To cite turbo-attn, use
[`CITATION.cff`](CITATION.cff), and please also cite TurboQuant (Zandieh et
al., *TurboQuant: Near-optimal KV Cache Quantization for LLM Inference*,
ICLR 2026). Release notes are on the
[GitHub Releases](https://github.com/arbicity/turbo-attn/releases) page.
