Turbo Attention (turbo-attn)

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)

pip install turbo-attn

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}")

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:

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). More runnable code is in examples/; the supported API is in docs/public_api.md, with every public name, every TKV_* flag and the engine 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/. "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) scores against the same engine's bf16 KV instead, so the two figures are not interchangeable. Fidelity ladder
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). 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
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 and docs/records/default-on-off-priced.md. Record

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). The Docker commands below need a clone of the repository for the compose files.

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):

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:

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):

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):

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, compose.sglang.yaml.

Reports: fidelity ladder, GPQA Diamond
Details: codec, calibration, kernels, engines and hardware, flags, precision levers, benchmarks
docs/ · examples/ · 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`, fitted by `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`).	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`). 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`).	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`).	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`). 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). 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.

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 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 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`
`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`, `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`
`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`, `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`

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). 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).
Prefill throughput on the served chunk compared with dense PyTorch flash SDPA: 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). The codec also works with third-party attention by decompressing to bf16 (docs/codec_third_party_interop.md).

Engines and hardware

Engine	How tkv attaches
vLLM (v0.31.0, vllm-turbo)	`vllm.general_plugins` entry point `tkv`; `--kv-cache-dtype tkv` or `tkv-bypass`, backend `TURBO_ATTN`.
SGLang (v0.5.21, 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).

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.
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).
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 and the tkv/_config_*.py modules it assembles). Print every flag with its type, default and description:

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 (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), and the records on larger, weight-quantized checkpoints (cross-engine-tip-rescore.md, pdown-quad-gate.md) and the Qwen3.5-0.8B record (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).

What is not claimed

Decode speed is measured only on MiniCPM5-2B on one RTX 4090: in 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/.
- 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 and its nightly. The harness README has the gates and the manual runbook.
Serving matrix: benchmarks/matrix_a/ and matrix-a.yml measure TTFT, TPOT and throughput for vLLM, SGLang and arbi-serve, tkv against bf16. Baselines are in benchmarks/baselines/.
Kernel benchmarks live in benchmarks/, and the measurement records behind kernel decisions are in docs/records/. Methodology notes: docs/BENCHMARKING.md.

Published reports: index.

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. 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).
More: ARCHITECTURE.md, docs/README.md, CONTRIBUTING.md.

License and citation

Apache License 2.0 (LICENSE, NOTICE). Vendored third-party code and technique provenance are listed in ATTRIBUTION.md. To cite turbo-attn, use 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 page.

Metadata

Release files for turbo-attn 0.67.0

For a detailed explanation of source distributions (sdists) and built distributions (wheels), please see the package formats documentation.

Source distribution (sdist)

Source distribution for turbo-attn 0.67.0
File	Size	Uploaded
turbo_attn-0.67.0.tar.gz	21.5 MB	Oct 9, 2026	Details

Built distribution (wheel)

Table of built distributions (wheels) for turbo-attn 0.67.0
File	Interpreter	ABI	Platform	Reset
turbo_attn-0.67.0-py3-none-any.whl 21.6 MB Oct 9, 2026	Python 3	none	any	Details

Total release size: 43.0 MB

Release files / turbo_attn-0.67.0.tar.gz

Download URL	turbo_attn-0.67.0.tar.gz
Size	21.5 MB
Tags	Source
SHA-256 checksum How to use checksums	`dfdc8a459d89ec04127ba0f0942260b45995eb56f66474a5ef9001ba27e9d58a`
BLAKE2b-256 checksum How to use checksums	`531e2d60a0d1b1366d25e890ecb6a4ee68152ccfc9c21022edb399f50cc3fbec`
Upload date	Oct 9, 2026
Uploaded using Trusted Publishing? What is trusted publishing?	No
Uploaded via	`uv/0.12.24 {"installer":{"name":"uv","version":"0.12.24","subcommand":["publish"]},"python":null,"implementation":{"name":null,"version":null},"distro":{"name":"Ubuntu","version":"24.04","id":"noble","libc":null},"system":{"name":null,"release":null},"cpu":null,"openssl_version":null,"setuptools_version":null,"rustc_version":null,"ci":true}`

Release files / turbo_attn-0.67.0-py3-none-any.whl

Download URL	turbo_attn-0.67.0-py3-none-any.whl
Size	21.6 MB
Tags	Python 3
SHA-256 checksum How to use checksums	`15f19cd73021c3c6df354fb1ade54c5fcf2e8771b9816a7d22827cef9f41a729`
BLAKE2b-256 checksum How to use checksums	`7c7d5034a88f1bf2f66dc513f02a0ba712f6e788c043330da2d03e95b382311a`
Upload date	Oct 9, 2026
Uploaded using Trusted Publishing? What is trusted publishing?	No
Uploaded via	`uv/0.12.24 {"installer":{"name":"uv","version":"0.12.24","subcommand":["publish"]},"python":null,"implementation":{"name":null,"version":null},"distro":{"name":"Ubuntu","version":"24.04","id":"noble","libc":null},"system":{"name":null,"release":null},"cpu":null,"openssl_version":null,"setuptools_version":null,"rustc_version":null,"ci":true}`