Metadata-Version: 2.4
Name: turbo-attn
Version: 0.66.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: Apache-2.0
Project-URL: Homepage, https://github.com/arbicity/turbo-attn
Project-URL: Repository, https://github.com/arbicity/turbo-attn
Project-URL: Issues, https://github.com/arbicity/turbo-attn/issues
Project-URL: Changelog, https://github.com/arbicity/turbo-attn/blob/main/CHANGELOG.md
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: License :: OSI Approved :: Apache Software License
Classifier: Operating System :: POSIX :: Linux
Classifier: Programming Language :: Python :: 3
Classifier: Programming Language :: Python :: 3.10
Classifier: Programming Language :: Python :: 3.11
Classifier: Programming Language :: Python :: 3.12
Classifier: Topic :: Scientific/Engineering :: Artificial Intelligence
Requires-Python: >=3.10
Description-Content-Type: text/markdown
License-File: LICENSE
License-File: NOTICE
Requires-Dist: torch>=2.1
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: cute
Requires-Dist: nvidia-cutlass-dsl[cu13]>=4.6.0; extra == "cute"
Requires-Dist: quack-kernels>=0.6; extra == "cute"
Requires-Dist: apache-tvm-ffi>=0.1.11; extra == "cute"
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: 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,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.

- **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](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 (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](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.3 s vs 22.1 s with bf16 (vLLM).
  [Record](benchmarks/baselines/minicpm5-2b/serve_matrix.md)

## 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/`](calibrations/), and `python -m tkv.auto_calibrate` makes one
for any other model ([details](#calibration)).

```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=calibrations/minicpm5-2b_v59.json \
  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

### Using the codec directly

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

```python
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/`](examples/) and [`docs/public_api.md`](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`](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 |
| 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 are also kept at 8 bits and read from there ([`docs/recent-ring.md`](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`](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`](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`.

```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
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.28.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.18, [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). 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-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
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`](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](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-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/`](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.
