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.
- 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). Fidelity ladder - 3.87× smaller KV cache than bf16 at 4 bits (4.13 bits/number measured on MiniCPM5-2B), and 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.3 s vs 22.1 s with bf16 (vLLM). Record
Quickstart
Needs one NVIDIA GPU of compute capability 8.x or 12.0 (RTX 30/40/50-series,
A-series, L40). Two settings: TKV_BITS (4 = K4V4, k5v4, …) and
TKV_CALIBRATION_FILE, the model's calibration store; stores ship in
calibrations/, and python -m tkv.auto_calibrate makes one
for any other model (details).
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=calibrations/minicpm5-2b_v59.json \
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.
More
- Reports: fidelity ladder, GPQA Diamond
- Details: codec, calibration, kernels, engines and hardware, flags, precision levers, benchmarks
docs/·examples/·docs/public_api.md
Details
Using the codec directly
pip install turbo-attn
import torch
from tkv import TkvCodec
codec = TkvCodec(head_dim=128, bit_width=4, device="cuda")
keys = torch.randn(8, 128, device="cuda")
packed, norms = codec.compress_k(keys)
recon = codec.decompress_k(packed, norms)
TkvCodec is pure PyTorch and also runs on CPU. The GPU kernels need the
cute extra: pip install "turbo-attn[cute]". More in
examples/ and docs/public_api.md.
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 (TKV_BAKE=1) 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 |
| 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 are also kept at 8 bits and read from there (docs/recent-ring.md). |
on (TKV_RECENT_RING, TKV_RECENT_RING_W=128) |
| 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/runtime/_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
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.28.0, vllm-turbo) | vllm.general_plugins entry point tkv; --kv-cache-dtype tkv or tkv-bypass, backend TURBO_ATTN. |
| SGLang (v0.5.18, 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). The int8 prefill needs 8.0 or newer.
- Head dims: decode 64, 128, 256, 512; int8 prefill 128 and 256; the CuTeDSL prefill covers the rest.
- Attention shapes: GQA, sliding window, MTP speculative verify. MLA
compression is off by default and opt-in (
TKV_MLA_ENABLE=1);tkv-bypassserves 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:
python -m tkv.config --doc
tkv refuses to start when the environment sets a TKV_* name that is not in
the registry, so a typo or a removed flag fails at boot instead of being
ignored. 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. |
TKV_CALIBRATION_BASIS |
unset | Which basis to serve when the store has several. |
TKV_AUTO_CALIBRATE_MODEL |
unset | vLLM: fit a missing store at boot. |
TKV_PREFILL_INT8 |
on | int8 prefill kernel; off routes prefill to CuTeDSL. |
TKV_PREFILL_BYPASS |
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 / 128 |
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, the int8 prefill kernel costs no fidelity:
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). 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-05), tkv decoded 28–59% faster per stream than bf16 on vLLM, SGLang
and arbi-serve (132 against 83.5 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_modelinconfig.json). - Models:
config.jsonis MiniCPM5-2B; the published Qwen3.5-0.8B ladder usesconfig_qwen35_08b.json. - How it runs: CI drives it through
kv-ablation-ladder.ymland its nightly. The harness README has the gates and the manual runbook.
- What one row is: one engine boot per KV cache format, teacher-forced
against a fixed target and scored against an FP32 oracle.
- Serving matrix:
benchmarks/matrix_a/andmatrix-a.ymlmeasure TTFT, TPOT and throughput for vLLM, SGLang and arbi-serve, tkv against bf16. Baselines are inbenchmarks/baselines/. - Kernel benchmarks live in
benchmarks/, and the measurement records behind kernel decisions are indocs/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.66.0
For a detailed explanation of source distributions (sdists) and built distributions (wheels), please see the package formats documentation.
Source distribution (sdist)
| File | Size | Uploaded | |
|---|---|---|---|
| turbo_attn-0.66.0.tar.gz | 1.7 MB | Details |
Built distribution (wheel)
| File | Interpreter | ABI | Platform | Reset |
|---|---|---|---|---|
| turbo_attn-0.66.0-py3-none-any.whl | Python 3 | none | any | Details |
Total release size: 3.5 MB
Release files / turbo_attn-0.66.0.tar.gz
| Download URL | turbo_attn-0.66.0.tar.gz |
|---|---|
| Size | 1.7 MB |
| Tags | Source |
|
SHA-256 checksum How to use checksums |
993c2a00af026e7c1b48f9c4bfd58cdd0d6b99c8bc4a188b4a4ad6f3c85c2268
|
|
BLAKE2b-256 checksum How to use checksums |
9b5f65b3b7440f568b4e12685b8a89166915ac02cf414329c065d8de376fb012
|
| Upload date | |
|
Uploaded using Trusted Publishing? What is trusted publishing? |
No |
| Uploaded via |
uv/0.12.23 {"installer":{"name":"uv","version":"0.12.23","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.66.0-py3-none-any.whl
| Download URL | turbo_attn-0.66.0-py3-none-any.whl |
|---|---|
| Size | 1.8 MB |
| Tags | Python 3 |
|
SHA-256 checksum How to use checksums |
5d7ba46cd0d7b64d2f800c542ad7798b29b0e7bedb65db1d375b30a20869cbce
|
|
BLAKE2b-256 checksum How to use checksums |
7228f717b3d0ad96be94d2f7df4119d1c84312cb9b96ebabba162c9e62d38a56
|
| Upload date | |
|
Uploaded using Trusted Publishing? What is trusted publishing? |
No |
| Uploaded via |
uv/0.12.23 {"installer":{"name":"uv","version":"0.12.23","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}
|
