Metadata-Version: 2.4
Name: glq
Version: 0.8.18
Summary: Lattice and trellis-coded quantization for LLM weights (2-8 bpw)
License-Expression: GPL-3.0-only
Classifier: Development Status :: 3 - Alpha
Classifier: Intended Audience :: Science/Research
Classifier: Programming Language :: Python :: 3
Classifier: Topic :: Scientific/Engineering :: Artificial Intelligence
Requires-Python: >=3.12
Description-Content-Type: text/markdown
License-File: LICENSE
Requires-Dist: torch>=2.0
Requires-Dist: numpy
Provides-Extra: hf
Requires-Dist: transformers; extra == "hf"
Requires-Dist: accelerate; extra == "hf"
Provides-Extra: cuda
Requires-Dist: triton; extra == "cuda"
Provides-Extra: dev
Requires-Dist: pytest; extra == "dev"
Provides-Extra: quantize
Requires-Dist: transformers; extra == "quantize"
Requires-Dist: datasets; extra == "quantize"
Requires-Dist: safetensors; extra == "quantize"
Requires-Dist: accelerate; extra == "quantize"
Requires-Dist: sentencepiece; extra == "quantize"
Provides-Extra: hub
Requires-Dist: jinja2; extra == "hub"
Requires-Dist: huggingface_hub; extra == "hub"
Requires-Dist: pyyaml; extra == "hub"
Provides-Extra: chat
Requires-Dist: gradio>=6; extra == "chat"
Requires-Dist: openai; extra == "chat"
Provides-Extra: bench
Requires-Dist: vllm; extra == "bench"
Requires-Dist: datasets; extra == "bench"
Requires-Dist: huggingface_hub; extra == "bench"
Requires-Dist: jinja2; extra == "bench"
Requires-Dist: pyyaml; extra == "bench"
Requires-Dist: matplotlib; extra == "bench"
Requires-Dist: pandas; extra == "bench"
Dynamic: license-file

# GLQ — fit larger LLMs on smaller GPUs

**Lattice and trellis-coded post-training quantization** for LLM weights: **2–8 bits/weight**,
served on **vLLM · HuggingFace Transformers**, with **deterministic**
fused CUDA kernels. Validated
from 24 GB 3090-class GPUs (A10G, sm_86) to a 96 GB RTX PRO 6000 Blackwell
(sm_120).

The recommended codebook is **trellis-coded quantization** (QTIP-derived TCQ, `--codebook
trellis`, since v0.7): it reaches an effective quantization dimension of 256 with a
lookup-free decode, which is why it wins where the bits are scarcest — 2 bpw on SmolLM3-3B
gives PPL **11.94** against **13.79** for the lattice path (bf16 9.12). It takes uniform
integer bit-rates, 2–8.

The **E8 lattice** codebooks — each group of 8 weights as a 16-bit index into a
65,536-entry codebook — remain, and not only for the checkpoints published before v0.7:
they are what fractional and per-layer mixed bit-rates run on, which trellis refuses.

Both share the rest of the pipeline. A Randomized Hadamard Transform makes the weights
incoherent
so Euclidean nearest-neighbour rounding is near-optimal under the Hessian-weighted
proxy loss, and a fused CUDA kernel matmuls **directly against the compressed
indices** — on the GPU serving path the dense weight is never materialized, so
GPU memory drops with the compression ratio (CPU inference and a few
architecture fallbacks dequantize instead).

**What you get**
- **NEW in v0.7 — trellis (TCQ) codebook**: single-stream decode at **bf16 speed**
  (176 vs 180 tok/s, SmolLM3-3B 4 bpw vs bf16, RTX PRO 6000 / vLLM) in a **third of the
  memory**, and good quality at 2–3 bpw. See
  [Trellis codebook](#trellis-codebook---codebook-trellis--qtip-derived-tcq).
- **2–8 bpw**, **no group-size constraint**, optional **per-layer mixed precision**.
- **Serve anywhere** — a vLLM plugin (weight + MoE + embedding) and an
  HF Transformers integration. `pip install glq`, load, run.
- **Small footprint** — smallest of the ~4-bit quantizers we measured (vs AWQ /
  NVFP4 on a 26B); a 31B fits ≈16.5 GiB at 5 bpw where bf16 needs ≈58 GiB, with
  quality within noise of bf16 on our paired reasoning evals.
- **Deterministic kernels** — bit-identical logits across runs (reproducible
  lm-eval scoring / on-policy RL rollouts).

**Pick your path** → [run a model](#run-a-pre-quantized-model) ·
[fit a bigger model on your card](#available-pre-quantized-checkpoints) ·
[quantize your own](#quantize-your-own-model) ·
[how GLQ compares](#how-glq-compares) · serve with
[vLLM](#docker-image-nvidia-gpu) ·
[how it works](#how-it-works)

## System requirements

**Linux x86_64.** An NVIDIA GPU is the fast path; since 0.8.16 a CPU-only machine is
also supported — the installer detects the hardware and sets up the matching vLLM
backend automatically.

- **GPU (recommended)**: NVIDIA Ampere-class or newer (`sm_86+`). Developed for 24–32 GB
  cards (3090 / 4080 / 4090 / L4 / L40S); validated live on an Ada L4 (`sm_89`) and an
  RTX PRO 6000 Blackwell (`sm_120`). A recent NVIDIA driver is required; a CUDA
  *toolkit* is not — `pip install glq` ships prebuilt kernels (CPython 3.12–3.14,
  x86_64), and `nvcc` is only needed for the from-source fallback (and for FlashInfer's
  sampler JIT on Blackwell — without it `glq-chat` falls back to vLLM's built-in sampler).
- **CPU-only**: no GPU → the installer sets up vLLM's CPU backend and GLQ's own fused
  CPU kernels (AVX2 floor, AVX-512 tiers used when present); `--cpu` forces this on a
  GPU machine. Trellis checkpoints only — dense **or MoE** since 0.8.17 (e8p/shell need
  CUDA) — and expect single-digit tok/s. Measured serving under vLLM on an 8-vCPU
  Sapphire Rapids (AVX-512-FP16), one request at a time, 64 decoded tokens, vLLM's
  default core binding: `gemma-4-26B-A4B` trellis-4bpw at **3.0 tok/s**, 0.4 s to first
  token, 18.6 GiB resident (13.9 GiB weights + a 4 GiB KV pool) — against **2.7 tok/s**
  for the dense `gemma-4-E4B` on the same box, because an MoE reads only its top-k
  experts per token. `glq-chat` also gives the kernels the core vLLM holds back
  (`VLLM_CPU_NUM_OF_RESERVED_CPU=0`), measured at **3.4 tok/s**; decode is
  memory-bound, so more threads than physical cores add nothing. Usable for short
  answers and background work, not fast chat. Model recommendations size against
  system RAM, and `glq-chat` sizes the KV pool from the checkpoint **and from the RAM
  that is actually free** — on CPU the weights, the pool, the runtime and the page cache
  all come out of the same RAM, so a browser already holding 10 GiB is 10 GiB this cannot
  have, and a machine with no swap (the cloud default) thrashes rather than failing
  cleanly when they do not fit. It prints the arithmetic before starting and says how
  much it is over by if it does not fit; `VLLM_CPU_KVCACHE_SPACE` overrides it.
- **Distros**: GLQ has only been tested end-to-end on **Ubuntu**. The installer's
  pre-flight is additionally exercised (Docker + GPU) on Ubuntu 24.04/26.04, Debian 13,
  Fedora 43/44, AlmaLinux 9, Arch and openSUSE Tumbleweed
  (`tests/test_installer_distros.py`), so other distros are expected to install and run.
- **Windows**: no native support (no Windows wheels, and vLLM is Linux-only). WSL2 with
  the NVIDIA CUDA driver should work in principle — untested.
- **macOS**: not supported. There is no CUDA on Apple hardware, and Docker Desktop on
  macOS cannot pass through an NVIDIA GPU, so a container does not help.
- **Python**: 3.12–3.14 (the installer creates its own venv).

## Quickstart

### Installer command (venv, glq, vLLM, chat UI, optional pi agent)

```bash
curl -fsSL https://raw.githubusercontent.com/cnygaard/glq/main/install.sh | bash
```

Or preselect components — `bash -s --` is how arguments reach a piped script. With
the quantize deps, for making your own checkpoints:

```bash
curl -fsSL https://raw.githubusercontent.com/cnygaard/glq/main/install.sh | bash -s -- --components core,vllm,chat,quantize
```

Or with the pi coding agent (installs node via nvm; afterwards `glq-code` serves a
tool-calling vLLM and runs pi against it):

```bash
curl -fsSL https://raw.githubusercontent.com/cnygaard/glq/main/install.sh | bash -s -- --components core,vllm,chat,picode
```

CPU-only serving on a machine that has a GPU you don't want used (no flag needed on a
machine without one — the installer detects it):

```bash
curl -fsSL https://raw.githubusercontent.com/cnygaard/glq/main/install.sh | bash -s -- --cpu
```

Creates a venv at `~/.glq/venv`, then discovers the published checkpoints, sizes
them against your GPU and offers the ones that fit. When it finishes it offers to
start GLQ and open the chat; answer no and it just prints the steps. It refuses to
run as root and never calls `sudo`.

In a terminal the installer **prompts** for components and checkpoint (the prompts
read `/dev/tty`, so they appear even though the script itself arrives on stdin).
Scripted — `--yes`, `--dry-run`, or no terminal at all (CI, `ssh host 'cmd'`,
`docker build`) — it prompts for **nothing** and installs the defaults, so use
`--components` there to get anything non-default. The **components**:

| Component | What it installs | Default |
|---|---|:-:|
| `core` | the venv + glq itself | always |
| `vllm` | vLLM — the OpenAI-compatible server behind the chat UI and picode | ✓ |
| `chat` | the Gradio chat UI | ✓ |
| `picode` | the pi coding agent (installs node via nvm) — run with `glq-code`, which starts a tool-calling vLLM for it and frees the GPU when pi exits | opt-in |
| `quantize` | the deps for quantizing your own models (`glq[quantize]`) | opt-in |

Other flags: `--dry-run` prints every command without running it; `--list` shows the
checkpoints and exits; `--start` / `--no-start` decide the handoff without being
asked; `--no-modify-path` leaves your shell rc file alone (by default the venv's
`bin/` is appended to `PATH` in `~/.bashrc` / `~/.zshrc` so `vllm`, `glq-chat` and
`ninja` resolve by name in new shells).

`~/.glq/venv/bin/glq-chat` is the one command afterwards: it starts vLLM, waits for
it, serves the Gradio UI on <http://localhost:7860>, and stops the server again when
you press Ctrl-C — vLLM has no idle unload, so a server left running keeps its share
of the card. It sizes the VRAM reservation from the checkpoint — weights, runtime
overhead and a usable cache — and sizes the context window from the card's KV
headroom, in tiers from 8192 up to 65536, clamped to the model's own maximum (a
24 GB card serving a 26B stays at 8192; a 96 GB card reaches 65536).
`--gpu-memory-utilization` / `--max-model-len` pin either; `--no-serve` attaches to
a server you started yourself.

The first start takes minutes — weights download, model load, CUDA-graph capture —
so it reports progress in vLLM's own words while it waits and writes the full server
log to `~/.glq/vllm.log`. `--verbose` streams that log instead of summarising it.

It also publishes a public `https://….gradio.live` link by default, so the chat can be
opened from a phone or another machine with no port forwarding. That link is
**unauthenticated for as long as the chat runs — anyone holding it can use your GPU**;
`--no-share` keeps everything on localhost.


The rest of this section document the manual path.

### Run a pre-quantized model

```bash
pip install 'glq[hf]'   # glq + transformers + accelerate; requires PyTorch ≥ 2.0
```

```python
import glq.hf_integration  # registers GLQ with transformers
from transformers import AutoModelForCausalLM, AutoTokenizer

model = AutoModelForCausalLM.from_pretrained(
    "xv0y5ncu/SmolLM2-360M-Instruct-GLQ-block-diagonal-4bpw",
    device_map="auto",
)
tok = AutoTokenizer.from_pretrained("xv0y5ncu/SmolLM2-360M-Instruct-GLQ-block-diagonal-4bpw")
print(tok.decode(model.generate(
    **tok("The capital of France is", return_tensors="pt").to(model.device),
    max_new_tokens=20,
)[0], skip_special_tokens=True))
```

`import glq.hf_integration` registers `quant_method="glq"` with HF
Transformers; `from_pretrained` then swaps `nn.Linear` for `E8RHTLinear`
and uses the fused CUDA C kernel on inference. CPU falls back to a
naive dequantize-then-matmul.

Or serve the fastest GLQ checkpoint on vLLM (the trellis-3INST decode —
single-stream speed at bf16 parity, 1.9 GiB of weights):

```bash
pip install glq vllm      # glq ≥ 0.7.0 (trellis kernel storage layout)
vllm serve xv0y5ncu/SmolLM3-3B-trellis-3inst-4bpw-kernel --quantization glq
```

> **Blackwell (sm_120 — RTX 5090, RTX PRO 6000):** vLLM's FlashInfer sampler ships no
> prebuilt kernel for this architecture and compiles one at startup. Without the NVIDIA
> CUDA Toolkit that build fails and takes the engine down *before the first token* — GLQ's
> own kernels are fine and load normally. Either install the toolkit, or run with
> `VLLM_USE_FLASHINFER_SAMPLER=0`. `glq-chat` detects this and falls back on its own;
> `vllm serve` and the `LLM(...)` API do not.

### Available pre-quantized checkpoints

The table mirrors the live [**start-here** collection](https://huggingface.co/collections/xv0y5ncu/start-here-recommended-glq-checkpoints)
— the same curated list `glq-setup --list` and the installer picker read — in its order,
plus one downloads-earned extra at the end. Everything is on the
[**`xv0y5ncu`** HF org](https://huggingface.co/xv0y5ncu).

| Repo | Base model | bpw | License | Footprint¹ | Best for |
|---|---|---|--:|--:|---|
| [`SmolLM3-3B-trellis-3inst-4bpw-kernel`](https://huggingface.co/xv0y5ncu/SmolLM3-3B-trellis-3inst-4bpw-kernel) | SmolLM3-3B | 4.0 trellis | Apache 2.0 | 1.9 GiB | **flagship / fastest GLQ decode** — single-stream at bf16 parity |
| [`gemma-4-26B-A4B-it-GLQ-trellis-3inst-4bpw`](https://huggingface.co/xv0y5ncu/gemma-4-26B-A4B-it-GLQ-trellis-3inst-4bpw) | Gemma-4-26B-A4B (MoE) | 4.0 trellis | Apache 2.0 | 14.4 GiB | best quality-per-GB (MoE) |
| [`gemma-4-26B-A4B-it-GLQ-trellis-3inst-3bpw`](https://huggingface.co/xv0y5ncu/gemma-4-26B-A4B-it-GLQ-trellis-3inst-3bpw) | Gemma-4-26B-A4B (MoE) | 3.0 trellis | Apache 2.0 | 11.9 GiB | the 26B for 12–16 GB cards — AIME-2026 avg@8 82.1% vs 86.25% at 4 bpw |
| [`Qwen3.8-27B-GLQ-trellis-3inst-4bpw`](https://huggingface.co/xv0y5ncu/Qwen3.8-27B-GLQ-trellis-3inst-4bpw) | Qwen3.8-27B (hybrid GDN) | 4.0 trellis | Apache 2.0 | 16.7 GiB | the biggest model on a 24 GB card — AIME-2026 avg@8 90.4% (bf16 87.9%, n=30) |
| [`Gemma-4-31B-it-GLQ-5.0bpw-mix3-8`](https://huggingface.co/xv0y5ncu/Gemma-4-31B-it-GLQ-5.0bpw-mix3-8) | Gemma-4-31B | 5.0 mix | Apache 2.0 | 16.5 GiB | a 31B on **one** 24–32 GB card (vs 57.9 GiB bf16) |
| [`Gemma-4-12B-it-GLQ-5.0bpw`](https://huggingface.co/xv0y5ncu/Gemma-4-12B-it-GLQ-5.0bpw) | Gemma-4-12B | 5.0 mix | Apache 2.0 | 6.9 GiB | 12B, mixed 3–8 bpw allocation |
| [`gemma-4-E4B-it-GLQ-trellis-3inst-4bpw`](https://huggingface.co/xv0y5ncu/gemma-4-E4B-it-GLQ-trellis-3inst-4bpw) | Gemma-4-E4B (8B, multimodal) | 4.0 trellis | Apache 2.0 | 6.58 GiB | a capable model on an 8–12 GB card |
| [`SmolLM3-3B-GLQ-block-diagonal-3.5bpw`](https://huggingface.co/xv0y5ncu/SmolLM3-3B-GLQ-block-diagonal-3.5bpw) | SmolLM3-3B | 3.5 mix | Apache 2.0 | 1.8 GiB | fractional-bpw example; small + fast |
| [`SmolLM2-360M-Instruct-GLQ-trellis-3inst-6bpw`](https://huggingface.co/xv0y5ncu/SmolLM2-360M-Instruct-GLQ-trellis-3inst-6bpw) | SmolLM2-360M | 6.0 trellis | Apache 2.0 | 0.31 GiB | near-lossless tiny — wikitext-2 PPL within +0.2% of bf16 |
| [`SmolLM2-135M-Instruct-GLQ-block-diagonal-4bpw`](https://huggingface.co/xv0y5ncu/SmolLM2-135M-Instruct-GLQ-block-diagonal-4bpw) | SmolLM2-135M | 4.0 | Apache 2.0 | 0.1 GiB | smallest checkpoint; CI smoke tests |
| [`Devstral-Small-2-24B-Instruct-GLQ-4bpw`](https://huggingface.co/xv0y5ncu/Devstral-Small-2-24B-Instruct-GLQ-4bpw) | Devstral-Small 24B | 4.0² | Apache 2.0 | ~20.5 GiB | coding / agentic (top-15 by downloads; not in start-here) |

**44 public checkpoints total** — the [HF org](https://huggingface.co/xv0y5ncu) also has
trellis builds of the whole Gemma-4 family (E2B/E4B/12B/26B/31B at 4 bpw), SmolLM3-3B and
SmolLM2-360M trellis ladders at 3–6 bpw, and the Gemma-4 12B/31B/E4B family across
**3–8 bpw** (incl. `e8p` variants). Per-model **quality** (MMLU-Pro / AIME, paired vs bf16) and **throughput** are in
each model card and in [How GLQ compares](#how-glq-compares) and
[Quality & footprint](#quality--footprint) below.

<sub>¹ **Footprint** = resident weight memory after load (vLLM's `Model loading took … GiB`)
— the figure that decides whether a model fits a 24/32 GB card. For current (block-diagonal)
checkpoints it tracks the bpw budget — that is what lets a 31B fit one GPU.</sub>

<sub>² Quantized before block-diagonal FHT became the quantizer default: power-of-2 FHT padding
is stored as real bits, so the checkpoint holds more bits per weight than its nominal rate and
the footprint is correspondingly larger than 4 bpw implies. The `-block-diagonal-` repos above
are true-to-label re-quants.</sub>

### Quantize your own model

```bash
pip install 'glq[quantize]'    # adds transformers, datasets, etc.

glq-quantize \
    --model HuggingFaceTB/SmolLM2-360M \
    --output ./smollm2-glq-4bpw \
    --bpw 4 \
    --nsamples 128 \
    --device cuda
```

This produces a **trellis** checkpoint: `--codebook` defaults to `trellis` (3INST
variant), the recommended format since v0.7. Other bit-widths: `--bpw 2` through
`--bpw 8`, **uniform integers only** — trellis has no mixed-precision encoding. For
fractional rates like `2.5`, and for per-layer mixed precision, add `--codebook e8_shell`
(or `e8p`); the refusal names them if you forget. `glq-quantize --help` lists every flag.

**`--streaming`** loads one layer at a time from safetensors. It is **required** for the
gemma-4 and Qwen3.5/Qwen3.8 families (their per-layer embeddings / multimodal wrappers
only load correctly this way), and useful for anything bigger than system RAM. Do *not*
pass it for Llama/SmolLM-style models — the streaming loader does not support their
profile — they load whole anyway.

A realistic large-model run (this is how the published Qwen3.8-27B checkpoints were made):

```bash
GLQ_TRELLIS_VARIANT=3inst glq-quantize \
    --model Qwen/Qwen3.8-27B \
    --output ./qwen38-27b-trellis-3inst-4bpw \
    --codebook trellis --bpw 4 --nsamples 128 --streaming
```

For **mixed-precision** allocation, run a two-pass flow: a profile
pass writes a per-layer `bpw_allocation.json`, then a quantize pass
applies it. See [`examples/quantize_mixed_precision.md`](examples/quantize_mixed_precision.md).

**Trellis (TCQ) quantization** — the fastest-decoding GLQ format and the
recommended pick from 2 bpw up:

```bash
GLQ_TRELLIS_VARIANT=3inst glq-quantize \
    --model HuggingFaceTB/SmolLM3-3B \
    --output ./smollm3-trellis-3inst-4bpw \
    --codebook trellis --bpw 4 --nsamples 128
```

Always set `GLQ_TRELLIS_VARIANT=3inst` for new quantizations: the
lookup-free 3INST decode is what the fused fast path is built for, and
its quality measured equal-or-slightly-better than the legacy `hyb`
lookup-table variant in our paired tests. (`hyb` remains the env default
only for back-compat with existing hyb checkpoints.)

Trellis constraints differ from the shell/e8p paths: **integer bpw only
(2–8)** — mixed precision and fractional rates are rejected rather
than silently rounded — and the fused kernel needs layer dims with
`out % 32 == 0`, `in % 64 == 0` (standard transformer shapes qualify).
Above 4 bpw the layer becomes a two-stage stacked RVQ (a K=4 code plus a
K=bpw−4 residual), which costs roughly 2× the decode of a single stage;
5–8 bpw checkpoints need **glq ≥ 0.8.0** and `GLQ_TRELLIS_VARIANT=3inst`.
Models with per-layer embeddings (Gemma-4 E2B/E4B) are handled
automatically — the PLE table quantizes via the shell codebook (requires
glq ≥ 0.7.2). `--streaming` is **required** for the Gemma-4 and
Qwen3.5/Qwen3.8 families, not just recommended.

**Mixture-of-Experts** (Gemma-4 26B-A4B and similar) needs **glq ≥ 0.8.1**,
which added the fused grouped trellis MoE decode. It serves under a full
CUDA graph — on an RTX PRO 6000 Blackwell the 26B-A4B at 4 bpw decodes
**91.6 tok/s at B=1** and 1341 tok/s at B=32, against 8.5 / 248 for the
per-expert fallback loop that preceded it. The fallback is still there for
prefill and for layer shapes the fused path declines, so a checkpoint that
cannot take the fast path still serves, just eagerly.

For how it actually works — the RHT bracket, the LDLQ + Viterbi encode, the
packed-tile storage layout and the fused decode — see
[docs/trellis-3inst.md](docs/trellis-3inst.md).
See [Trellis codebook](#trellis-codebook---codebook-trellis--qtip-derived-tcq)
for details.

## Docker image (NVIDIA GPU)

A prebuilt CUDA image ships everything needed to run GLQ models —
`glq`, PyTorch, vLLM, transformers, and lm-eval on CUDA 12.8:

```
ghcr.io/cnygaard/glq-env:latest     # CUDA 12.8 bundle: glq + vLLM + transformers + lm-eval
```

**Prerequisite — GPU access in Docker.** You need an NVIDIA GPU plus
the [NVIDIA Container Toolkit](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/install-guide.html)
installed on the host; that's what makes the `--gpus all` flag pass
the GPU into the container. Verify it works:

```bash
docker run --rm --gpus all ghcr.io/cnygaard/glq-env:latest nvidia-smi
```

If that prints your GPU table, you're set. (No toolkit → `--gpus`
errors with "could not select device driver".)

**Produce output.** Mount a host directory for the model cache (the
image's `HF_HOME` is `/cache/hf`, so models persist across runs
instead of re-downloading), then generate:

```bash
docker run --rm --gpus all \
    -v "$HOME/.cache/huggingface:/cache/hf" \
    ghcr.io/cnygaard/glq-env:latest \
    python -c '
import glq.hf_integration, torch                      # registers GLQ with HF
from transformers import AutoModelForCausalLM, AutoTokenizer
mid = "xv0y5ncu/SmolLM3-3B-GLQ-3.5bpw"
tok = AutoTokenizer.from_pretrained(mid)
model = AutoModelForCausalLM.from_pretrained(
    mid, device_map="cuda", torch_dtype=torch.float16)
ids = tok("The capital of France is", return_tensors="pt").to("cuda")
print(tok.decode(model.generate(**ids, max_new_tokens=20)[0],
                 skip_special_tokens=True))
'
```

Expected output:

```
The capital of France is Paris. It is located in the north of the country.
```

The first run downloads the model into the mounted cache; later runs
reuse it. Swap `mid` for any GLQ checkpoint (see
[Available pre-quantized checkpoints](#available-pre-quantized-checkpoints)).

**Flag reference:**

| Flag | Why |
|---|---|
| `--gpus all` | binds **all** host GPUs into the container (needs the NVIDIA Container Toolkit). Use `--gpus '"device=0"'` to pick one. |
| `-v "$HOME/.cache/huggingface:/cache/hf"` | persists downloaded weights on the host (`HF_HOME=/cache/hf` inside) so they survive `--rm`. |
| `--rm` | remove the container when it exits (drop it to keep the container around). |

**Serving (vLLM) & an interactive shell.** The image bundles vLLM, so
you can serve an OpenAI-compatible endpoint — publish the port and
mount the cache:

```bash
# Plain chat — the model ships its own chat template, so nothing extra needed:
docker run --rm --gpus all -p 8000:8000 \
    -v "$HOME/.cache/huggingface:/cache/hf" \
    ghcr.io/cnygaard/glq-env:latest \
    vllm serve xv0y5ncu/Gemma-4-E4B-it-GLQ-4bpw --max-model-len 64000
```

**Tool-calling + thinking.** Gemma-4's tool template is *not* in the model (its
bundled `chat_template.jinja` is plain chat) and *not* in the vLLM pip wheel, so
fetch it from vLLM's `examples/` first, then mount it:

```bash
curl -fsSL https://raw.githubusercontent.com/vllm-project/vllm/v0.20.2/examples/tool_chat_template_gemma4.jinja \
    -o tool_chat_template_gemma4.jinja

docker run --rm --gpus all -p 8000:8000 \
    -v "$HOME/.cache/huggingface:/cache/hf" \
    -v "$PWD/tool_chat_template_gemma4.jinja:/work/tool.jinja:ro" \
    ghcr.io/cnygaard/glq-env:latest \
    vllm serve xv0y5ncu/Gemma-4-E4B-it-GLQ-4bpw \
        --max-model-len 64000 \
        --enable-auto-tool-choice \
        --tool-call-parser gemma4 \
        --reasoning-parser gemma4 \
        --chat-template /work/tool.jinja \
        --default-chat-template-kwargs '{"enable_thinking": true}'
```

Pass `--chat-template` the **in-container** mount path (`/work/tool.jinja`), not
the host path. `--default-chat-template-kwargs '{"enable_thinking": true}'`
defaults Gemma-4 reasoning on. The `gemma4` parsers and all of these flags are
accepted by the image's bundled vLLM 0.20.2 and the model loads; note that
startup runs a multi-minute `torch.compile` + CUDA-graph capture before the
endpoint is ready. See the
[vLLM Gemma-4 recipe](https://docs.vllm.ai/projects/recipes/en/latest/Google/Gemma4.html)
for the full tool-calling / reasoning reference.

> **Image vs. pip:** the image is a convenience bundle and may lag the PyPI
> package — `pip install -U glq` inside the container (or your own venv) always
> gets the newest release. The pip package is the source of truth; the image just
> saves you assembling a matching CUDA + vLLM + transformers stack.

The image's default command is a shell (`docker run --rm -it --gpus all
ghcr.io/cnygaard/glq-env:latest`) if you'd rather poke around interactively.

## How GLQ compares

GLQ sits next to the other post-training quantizers — AWQ, GPTQ, NVIDIA's NVFP4
(TensorRT Model Optimizer), and Unsloth's dynamic mixed-precision. They optimize for
different things; here is the honest layout.

| | **GLQ** | AWQ | GPTQ | NVFP4 (ModelOpt) | Unsloth dynamic |
|---|---|---|---|---|---|
| Bits / weight | **2–8** + mixed | 4 (grouped) | 3–4 (grouped) | ~4 (W4A4) | mixed 2–8 (selective) |
| Group-size constraint | **none** | g64/g128 | g64/g128 | — | — |
| Core method | E8 lattice + RHT + LDLQ | activation-aware scale | block error-feedback | FP4 + per-tensor scale | per-layer bit allocation |
| Footprint at ~4-bit | **smallest** (no per-group scales/zeros) | + group scales/zeros | + group scales/zeros | + FP8 scales | varies |
| Speed on Blackwell | W4A16; **single-stream at bf16 parity** (trellis 3INST, 3B measured) | W4A16 (Marlin) | W4A16 (Marlin) | **fastest at batch** (native FP4) | n/a (GGUF) |
| Serving stack | **vLLM · HF** | vLLM · HF · TRT | vLLM · HF · TRT | vLLM · TRT-LLM | **llama.cpp / Ollama** (GGUF) |
| Bit-exact deterministic kernels | **yes** | — | — | — | — |
| Fine-tuning (QLoRA) | — | — | — | — | **yes** |

<sub>A "—" means the feature is absent or not advertised by that method — we have not
tested the other methods for those properties.</sub>

**Where each wins** — pick by your constraint, not by a single "best":

- **GLQ** — the smallest footprint of the ~4-bit methods we measured, the widest bit-range (2–8) with
  per-layer mixed precision, deterministic kernels, and built-in KV compression. The pick
  when you are **memory-bound** — fit a bigger model or longer context on a 24–32 GB card —
  and serve on vLLM / HF.
- **NVFP4 (ModelOpt)** — **fastest on Blackwell** (native FP4 tensor cores). The pick when
  you have a Blackwell GPU with memory to spare and want raw decode speed; it trades a larger
  footprint than GLQ for that speed.
- **AWQ / GPTQ** — mature, ubiquitous 4-bit weight-only with fast Marlin kernels. The safe
  default on any GPU when a ~4-bit footprint is enough and you don't need <4 bpw.
- **Unsloth dynamic** — selective mixed-precision for the **GGUF / llama.cpp** stack, plus a
  strong QLoRA fine-tuning story. A different target (CPU/edge/Ollama, or training) than GLQ's
  GPU serving — GLQ also does per-layer mixed precision, but for the vLLM/HF stack.

Where GLQ speed stands after v0.7.1 (measured, SmolLM3-3B on an RTX PRO 6000, vLLM 0.25):
**single-stream (B=1) trellis-3INST decode is at bf16 parity** — 176 vs 180 tok/s — while
using a third of the weight memory. At batch, bf16 still leads (2,423 vs 4,887 tok/s at
concurrency 32) — the batched GEMM is the remaining gap, though v0.8.14 narrowed it: a
parallel-column reduce in the batched kernel is worth −26–29% kernel time at B=16–64
(RTX PRO 6000, 27B shapes; −12–21% on an L4) and +11.8% serving throughput at
concurrency 32 (Qwen3.8-27B-3bpw, same box/workload), bit-identical outputs. On a
bandwidth-starved L4 the compressed batched kernel now beats cuBLAS-on-dense through
B=32. The shell/e8p codebooks decode
slower than trellis; their draw is the 2–8 bpw range and mixed precision. A
weight-quantization method's headline win remains **footprint** — the freed VRAM as KV /
longer-context headroom, and fitting models bf16 can't — but at 4 bpw trellis the
single-user speed cost of that footprint is now ~zero on the hardware we measured.

### Matched 4-bit head-to-head (measured)

Numbers behind the matrix: GLQ vs AWQ vs NVFP4 on **one** base model (`gemma-4-26B-A4B-it`, an
MoE), **all at ~4-bit** so footprint is directly comparable, on a single RTX PRO 6000 Blackwell
(vLLM 0.23, seed 0, thinking mode). bf16 is the uncompressed ceiling.

| Method | bits | Weights, exact on disk¹ | MMLU-Pro n=60² | AIME-2026 n=30² |
|---|--:|--:|--:|--:|
| bf16 (ceiling) | 16 | ~50 GB | 91.7% | 90.0% |
| **GLQ 4 bpw** | 4.0 | **14.97 GB** | **93.3%** | **90.0%** |
| AWQ 4-bit | ~4.25 | 17.19 GB | 86.7% | 83.3% |
| NVFP4 (W4A4) | ~4 | 18.78 GB | 86.7% | **90.0%** |

**GLQ has the smallest footprint, the top MMLU-Pro, *and* matches the bf16 ceiling on hard
reasoning** — AIME-2026 GLQ 90.0% = bf16 90.0% = NVFP4 90.0% > AWQ 83.3% (GLQ reproduces bf16's
27/30 exactly, at <1/3 the footprint); MMLU-Pro is saturated (all 86.7–93.3%, inside the ±8% n=60
band → a tie). The trade is **decode speed**: GLQ ran the same 30-problem AIME in **58 min** vs
~24 min for AWQ/NVFP4 and ~32 min for bf16 — its W4A16 kernel is L2-codebook-gather-bound on
Blackwell, where NVFP4's native FP4 tensor cores, AWQ's Marlin, and even bf16 all decode faster.
So GLQ's win is **footprint + quality, not speed** — the pick when you are memory-bound. Full
tables, truncations, chain lengths, and caveats:
[`benchmarks/_quant_compare_gemma4_26b.md`](benchmarks/_quant_compare_gemma4_26b.md).

<sub>¹ Exact safetensors weight bytes summed from tensor headers (tower-independent — all three ship
the identical 1.15 GB vision tower); loaded footprints track this ordering (GLQ ~14.0 < AWQ 15.6 <
NVFP4 17.1 GiB). ² Thinking mode, single-sample pass@1, seed-fixed subsets; n=60 MMLU-Pro 95% CI ≈
±8%, n=30 AIME ≈ ±15% — a fidelity comparison on one model / one GPU, not a leaderboard. NVFP4 is
W4A4 (4-bit activations) and defaults to fp8 KV, a speed edge the W4A16 methods don't take;
calibration differs per vendor. The paired GLQ-vs-bf16 view (incl. AIME-2024) is in
[Quality & footprint](#quality--footprint) below.</sub>

### Trellis 3INST vs bf16 vs NVFP4 (measured)

The same exercise for the **trellis-3INST** format on a dense model: SmolLM3-3B, one
RTX PRO 6000 Blackwell, vLLM 0.25.0, glq 0.7.1. Speed = `vllm bench sweep serve`
(random 128-in/256-out, ignore-eos, seed 42, mean of 3 runs); quality = wikitext-2 PPL +
AIME-2026 in thinking mode (32k budget, 8 samples/problem, avg@8).

| | **GLQ trellis-3INST 4 bpw** | bf16 | NVFP4 (W4A4)¹ |
|---|--:|--:|--:|
| Weights on disk | **1.9 GB** | 5.8 GB | 2.5 GB |
| Decode, 1 stream (tok/s) | **176** | 180 | 301 |
| Decode, 32 streams (tok/s) | 2,423 | 4,887 | 7,600 |
| PPL (wikitext-2) | 9.23 | 9.12 | n/a² |
| AIME-2026 avg@8 (thinking) | **41.7%** | 47.5% | 32.5% |
| 240-generation batch job, wall-clock³ | 52 min | 36 min | 32 min |

**The read**: at 4 bpw trellis, GLQ decodes **single-stream at bf16 parity** in a third of
the memory, and keeps most of the reasoning quality (−5.8 pts AIME vs bf16) where this
NVFP4 checkpoint's 4-bit activations cost −15 pts. NVFP4 is decisively faster — but its
FP4 tensor cores exist **only on Blackwell**; on the 24 GB 3090-class (sm_86) / 4090-class
(sm_89) cards GLQ targets, NVFP4 cannot use FP4 hardware, while GLQ's fp16-mma kernels have
no Blackwell dependency (core GLQ kernels validated on sm_86 / sm_89 / sm_120; the trellis
kernels to date on sm_120).

<sub>¹ [`Firworks/SmolLM3-3B-nvfp4`](https://huggingface.co/Firworks/SmolLM3-3B-nvfp4), a
community w4a4 checkpoint — results are scoped to it, not to NVFP4 at large. ² The
checkpoint does not load in HF transformers; quality measured via vLLM (AIME row).
³ 30 AIME problems × 8 samples at up to 32k tokens each, launched as one batch — fixed
load/compile overhead and the low-concurrency straggler tail compress the steady-state
speed ratios (3.1× at c=32 becomes 1.6× on the real job).</sub>

## Quality & footprint

### SmolLM3-3B at matched 4.5 bpw vs GPTQ

Blackwell RTX PRO 6000, 128 calibration samples,
`lm-evaluation-harness` limit=200/task (GSM8K n=500, MMLU 50/subtask).
GLQ 4.5 bpw uses two-pass mixed allocation (91 layers @ 4 bpw + 161 @
5 bpw, avg 4.64 bpw).

| Task                     | bf16   | **GLQ 4.5 bpw** | GPTQ W4 g128 |
|--------------------------|--------|-----------------|--------------|
| ARC-challenge (acc_n)    | 0.490  | **0.475**       | 0.420        |
| ARC-easy (acc_n)         | 0.745  | **0.735**       | 0.695        |
| HellaSwag (acc_n)        | 0.660  | 0.660           | **0.675**    |
| MMLU (acc)               | 0.617  | **0.603**       | 0.589        |
| TruthfulQA mc2           | 0.529  | **0.545**       | 0.515        |
| WinoGrande               | 0.655  | 0.660           | **0.670**    |
| WikiText-2 ppl ↓         | 10.67  | **10.90**       | 11.33        |
| GSM8K flex (n=500)       | 0.722  | **0.738**       | 0.688        |
| IFEval prompt-strict     | 0.310  | 0.310           | 0.285        |
| IFEval prompt-loose      | 0.325  | **0.330**       | 0.295        |
| IFEval inst-strict       | 0.478  | 0.472           | 0.453        |
| IFEval inst-loose        | 0.494  | 0.491           | 0.469        |

GLQ beats GPTQ on 10/12 metrics. WikiText-2 ppl gap to bf16: +2.2 %
(GLQ) vs +6.2 % (GPTQ). GSM8K flex matches bf16; GPTQ drops 0.034.

### Small models: SmolLM2-360M-Instruct at 4 bpw

GPTQ requires a group-size dividing the hidden dim; SmolLM2-360M's
hidden=960 is not divisible by 128, forcing `group_size=64` (~4.5 eff
bpw) and losing quality. GLQ has no group-size constraint.

| Method        | bpw stored | 5-task avg | % of bf16 |
|---------------|------|------------|-----------|
| bf16          | 16.0 | 0.557      | 100 %     |
| **GLQ 4-bit block-diagonal** (current) | 4.0 | — | **97.9 %**¹ |
| GLQ 4-bit legacy (padded) | ~6.4 eff¹ | 0.555 | 99.6 % |
| GPTQ W4 (g64) | ~4.5 | 0.486      | 87.2 %    |

5-task = ARC-e, HellaSwag, PIQA, WinoGrande, LAMBADA; 128 calibration
samples; L40S. GPTQ's LAMBADA collapses to 0.346; GLQ preserves 0.508.

<sub>¹ The original "4 bpw" checkpoint predates block-diagonal FHT and stored its
power-of-2 padding as real bits — effectively ~6.4 bits/weight, so its 99.6 % was
earned with extra storage. The true-4-bpw re-quant
([`SmolLM2-360M-Instruct-GLQ-block-diagonal-4bpw`](https://huggingface.co/xv0y5ncu/SmolLM2-360M-Instruct-GLQ-block-diagonal-4bpw))
scores 97.9 % of bf16 on the same 5 tasks (paired GLQ/bf16 runs on the vLLM stack)
— still well clear of GPTQ. The group-size constraint on GPTQ applies regardless.</sub>

### Gemma-4 family — GLQ vs bf16, paired (thinking mode)

Each GLQ checkpoint was run **head-to-head against its bf16 base on the same
questions**, in thinking mode (these are reasoning models), on a single RTX PRO 6000
Blackwell with vLLM 0.23. These are **small-n fidelity** comparisons (95% CI ≈ ±8%
MMLU-Pro, ±15% AIME), single-sample pass@1 — not leaderboard scores.

| Model (GLQ vs bf16) | bpw | Footprint GLQ / bf16 | MMLU-Pro n=60 | AIME-2024 n=30 |
|---|--:|--:|--:|--:|
| Gemma-4-31B-it | 5.0 mix | **16.5** / 57.9 GiB | 90.0% vs 86.7% | 90.0% vs 86.7% |
| Gemma-4-26B-A4B-it (MoE) | 4.0 | **~15** / ~50 GiB | 93.3% vs 91.7% | 93.3% vs 93.3% |
| Gemma-4-12B-it | 5.0 mix | **6.9** / 24 GiB | 81.7% vs 78.3% | 83.3% vs 93.3%† |

On these runs GLQ is **within noise of bf16** — ahead on MMLU-Pro for all three,
ahead/tied on AIME-2024 for the 31B and 26B-A4B. †At n=30 the 12B AIME-2024 gap
(bf16 +3 items) is not statistically significant; likewise the 31B's AIME-2026
(GLQ 83.3% vs bf16 90.0%, bf16 +2 items). Footprint is the consistent win — a 31B in
16.5 GiB fits one 24–32 GB card where bf16 (~58 GiB) needs three. See each model card
for the full paired tables, thinking budgets, and caveats.

For decode speed vs bf16 on the current stack, see
[How GLQ compares](#how-glq-compares) — single-stream trellis-3INST at bf16
parity, bf16 ahead at batch (measured numbers there).

## How it works

Every codebook shares the same bracket: rotate (RHT), quantize with error
feedback (LDLQ), then serve from the compressed form with fused kernels.
What differs is the codebook in the middle — **trellis is the default**
since v0.8.8; the E8 lattice family covers fractional and mixed precision.

1. **Randomized Hadamard Transform.** Random sign flips followed by a
   Fast Walsh–Hadamard Transform rotate both weights and Hessian.
   After RHT the Hessian is approximately diagonal, so plain Euclidean
   nearest-neighbour against the codebook is near-optimal under the
   Hessian-weighted proxy loss. Large layers use a block-diagonal
   transform; the inverse runs inside the inference kernels.

2. **LDLQ error feedback.** Block-LDL decomposition of the Hessian
   drives a sequential sweep — GPTQ-style, but over multi-weight blocks
   instead of scalar columns. Each block's quantization error
   propagates forward to correct downstream blocks.

3. **Codebooks.**
   - **Trellis (default; QTIP-style TCQ).** No stored codebook at all: a
     16-bit state slides along the weight stream emitting two weights per
     step, and each weight costs its K bits (2–4 native). The shipped
     `3inst` variant decodes a state arithmetically — an integer
     multiply-add, a mask/xor, and one fp16 add — so decode needs no
     table lookup and no shared memory. Encoding is a Viterbi
     (add-compare-select) sweep over the trellis, fused in Triton.
     5–8 bpw stack two trellis stages as residual VQ (each stage K ≤ 4,
     residual scale fitted per layer).
   - **E8 lattice (`e8_shell`, `e8p`).** 65,536 vectors from the first
     seven shells of the E8 lattice; each 8-weight group is one 16-bit
     index (2 bpw primary), refined by 8-bit or 16-bit residual stages.
     This is the path for fractional (e.g. 3.5 bpw) and per-layer mixed
     precision; `e8p` adds a block-diagonal tensor-core decode.

4. **Fused inference kernels.** Custom CUDA and Triton kernels keep the
   weights compressed end-to-end — the dense matrix is never
   materialized, so VRAM scales with the compression ratio. Trellis
   decode runs entirely in registers and feeds tensor-core `mma`
   directly (B=1 GEMV, a batched kernel through B=64 with a
   parallel-column reduce, dense fallback for long prefills); E8 decode
   gathers from the L2-cached 1 MiB codebook. Decode is bit-exact
   against the Python reference and deterministic run-to-run — no
   atomics, fixed reduction order — which is what lets the test suite
   gate kernels with `torch.equal` rather than tolerances.

## Advanced

### CUDA-graph decode wrapper

The B=1 autoregressive decode path is Python-dispatch-bound in eager
mode. `CUDAGraphWrapper` captures the fixed-shape decode and replays
it; benchmarks below are on SmolLM3-3B 3.5bpw, L40S.

| Mode       | GLQ 3.5 bpw | bf16   |
|------------|-------------|--------|
| Eager      | 25 tok/s    | 40     |
| CUDA graph | 37 tok/s    | 40     |

```python
from glq.cuda_graph import CUDAGraphWrapper
wrapper = CUDAGraphWrapper(model)
logits = wrapper(input_ids)   # first call captures; replays after
```

The wrapper falls back to eager for variable shapes (prefill, batch>1,
extra kwargs). For 24B models the matmul is compute-bound at B=1, so
graphs don't help (Devstral-24B GLQ 4 bpw: 6.6 tok/s eager vs 6.4
graphed).

### Tuning vLLM CUDA-graph capture sizes (v0.3.4+)

vLLM 0.20 captures both **FULL** model-forward graphs (single replay
per fixed shape) and **PIECEWISE** subgraphs split at attention. The
default capture set is derived from `max_num_seqs * 2`, so a
single-sequence harness only gets FULL captures for `[1, 2]`. For
batched serving, raise the list explicitly:

```python
from vllm import LLM
llm = LLM(model="xv0y5ncu/Gemma-4-E4B-it-GLQ-4bpw",
          compilation_config={
              "cudagraph_capture_sizes": [1, 2, 4, 8, 16],
          })
```

Measured impact on Gemma-4-E4B-it-GLQ-4bpw, RTX PRO 6000 Blackwell,
256-token decode:

| Mode | B=1 tok/s | B=4 tok/s (total) |
|---|---:|---:|
| Eager | 14.4 | 35.0 |
| Piecewise + default capture `[1, 2]` | 39.4 | 132.7 |
| Piecewise + capture `[1, 2, 4, 8, 16]` | **40.0** | **157.3 (+18.5 %)** |

At B=1 the FULL graph was already captured (no change). At B=4 the
extended list keeps the FULL graph active where the default
degenerated to PIECEWISE-only, recovering ~6 tok/s per sequence.

Cost: ~10-20 MB VRAM per captured shape on 3B / E4B models (vLLM
prints the total at "Graph capturing finished in N s, took X GiB").
On 24-31B models budget ~100-200 MB per shape. Capture time is
~1 s per shape, one-time at LLM init.

### Bit widths

| bpw | Primary | Residual stages |
|-----|---------|-----------------|
| 2   | 16 b    | —                |
| 3   | 16 b    | + 8 b            |
| 4   | 16 b    | + 16 b           |
| 5   | 16 b    | + 16 b + 8 b     |
| 6   | 16 b    | + 16 b + 16 b    |
| 7   | 16 b    | + 16 b + 16 b + 8 b |
| 8   | 16 b    | + 16 b + 16 b + 16 b |

One global scale per layer; no group-size parameter. Non-power-of-2
hidden sizes use block-diagonal FHT (v0.2.9+) — e.g. 2688 is
decomposed as `2048 + 512 + 128` so on-disk storage matches the
nominal rate exactly.

The table above applies to the shell and e8p codebooks. The
[trellis codebook](#trellis-codebook---codebook-trellis--qtip-derived-tcq)
takes integer rates only: 2–4 bpw as a single K-bit trellis code, 5–8 bpw
as two stacked codes (K=4 plus a K=bpw−4 residual).

### Trellis codebook (`--codebook trellis`) — QTIP-derived TCQ

`--codebook trellis` replaces the per-8-weight lattice lookup with **trellis-coded
quantization** (TCQ) over 256-weight sequences, following
[QTIP](https://github.com/Cornell-RelaxML/qtip) (Tseng et al., 2024): a Viterbi search
encodes each row against a tail-biting trellis, so neighbouring weights share state and
the effective codebook is exponentially larger than a flat lookup at the same rate. In our
SmolLM3-3B tests it is GLQ's best format across **2–8 bpw** — at 2 bpw it clearly beats the
e8p codebook (PPL 11.74 vs 13.21) — and at 4 bpw its decode is GLQ's fastest
(single-stream at bf16 parity, see
[the measured table](#trellis-3inst-vs-bf16-vs-nvfp4-measured)).

**Variants.** `GLQ_TRELLIS_VARIANT=3inst` decodes each 16-bit trellis state
*arithmetically* (a hash + two fp16 halves — "3 instructions", no lookup table), which is
what the fused fast path is built around; the default `hyb` uses a 512-entry lookup
table and exists for back-compat with earlier hyb checkpoints. Quality measured
equal-or-slightly-better for 3INST in our paired tests — **use 3INST for new quants**.

**What the v0.7.1 kernels do** (the 129 → 179 tok/s single-stream jump): the input
Hadamard transform and fp16 cast run *inside* every decode block instead of as separate
1-block launches; the five lowest FHT butterfly stages run as warp shuffles; and the
per-shard output transforms of fused QKV / gate-up layers batch into single launches.
All bit-exact — wikitext-2 PPL is unchanged to the fourth decimal across the entire
optimization series. Runtime opt-outs, should you ever need the unfused paths:
`GLQ_TRELLIS_FUSE_INPUT=0`, `GLQ_TRELLIS_BATCH_OUT_RHT=0`.

**Rates 5–8 bpw (stacked RVQ, v0.8.0).** A trellis code's window collapses as K grows —
at 16 bits of state a native K=6 keeps far less history than K=4 — so above 4 bpw GLQ
stacks **two** codes instead of widening one: a K=4 primary plus a K=(bpw−4) residual
fitted per layer. Measured on SmolLM3-3B, that is worth it: 6 bpw reaches wikitext-2 PPL
**9.1310** against bf16's 9.1220, closing 92% of the 4 bpw → bf16 gap, at exactly 6.00
bits/weight on disk. The cost is decode: two stages means ~2× the state decodes, which is
irreducible — measured 1.9× on sm_120 and 2.3–2.7× on sm_89 at B=1. These serve on vLLM
(`--quantization glq`, 3INST only; HYB has no two-stage kernel), and a checkpoint whose
stored stages disagree with its declared bpw is rejected at load rather than served a
stage short.

**Storage layout.** Trellis checkpoints store indices in the decoder-native "kernel"
layout (`trellis_layout: "kernel"` in the config) — loading them requires
**glq ≥ 0.7.0**, or **≥ 0.8.0** for the 5–8 bpw two-stage layout; older versions abort
with a layout error rather than decode garbage. Mixed precision and fractional rates are
rejected at quantize time. Quantization cost is the Viterbi encode: ~35 min for a 3B on
one GPU (CUDA-graph-cached).

### E8P codebook (`--codebook e8p`) — derivative of QuIP#

`--codebook e8p` is **derivative work that ports the E8P codebook and its
tensor-core decode kernels from
[QuIP#](https://github.com/Cornell-RelaxML/quip-sharp)** (Tseng et al., 2024).
It swaps the default E8-shell codebook for QuIP#'s 2-bit **E8P** (padded-D̂8)
grid decoded on tensor cores (`mma.sync` GEMV), plus residual vector
quantization for the higher rates. The port is self-contained (no `quiptools`
dependency); full credit and the citation are in
[Acknowledgments](#acknowledgments).

| bpw | RVQ recipe (stages)         | added per stage |
|-----|-----------------------------|-----------------|
| 2   | `[E8P]`                     | 16-bit E8P (primary) |
| 3   | `[E8P, E81B]`               | + 8-bit E81B residual |
| 4   | `[E8P, E8P]`                | + 16-bit E8P residual |
| 5   | `[E8P, E8P, E81B]`          | + 8-bit E81B |
| 6   | `[E8P, E8P, E8P]`           | + 16-bit E8P |
| 7   | `[E8P, E8P, E8P, E81B]`     | + 8-bit E81B |
| 8   | `[E8P, E8P, E8P, E8P]`      | + 16-bit E8P |

Each E8P stage is a 16-bit tensor-core decode (`mma.sync` GEMV, +2 bpw); odd
bit-widths end in a single 8-bit **E81B** residual (WMMA lookup-matmul, +1 bpw).

```bash
glq-quantize --model HuggingFaceTB/SmolLM2-360M --output ./out \
    --codebook e8p --bpw 4 --nsamples 128
```

The entire linear — input RHT → E8P tensor-core decode (all RVQ stages) →
×Wscale → output RHT — runs as a **single fused CUDA op**
(`glq_fused_linear_e8p_cuda`), so vLLM captures B=1 decode in a **FULL CUDA
graph** exactly like the default path (HF inference is supported too). On
SmolLM3-3B 4 bpw (RTX PRO 6000 Blackwell, vLLM 0.23) that fused op runs B=1
decode at **~85 tok/s**; collapsing the per-linear dispatch into one opaque op
is what makes cudagraph a win here — the equivalent unfused multi-op path is
~7× slower under capture. **The full 2–8 bpw range serves on both HF and vLLM
as of v0.6.7** (the 5–8 bpw N-stage decode landed there); the fused op is
bit-exact against the unfused reference across all of 2–8 bpw (decode reproduces
the quantize-side weight to ~66–68 dB SQNR per bit-width).

### Devstral-24B tokenizer

`transformers` 5.x auto-routes Mistral/Devstral models through
`mistral_common`, which rejects the standard `tokenizer.json`. Use
`PreTrainedTokenizerFast` explicitly:

```python
from huggingface_hub import snapshot_download
from transformers import AutoModelForCausalLM, PreTrainedTokenizerFast

path = snapshot_download("xv0y5ncu/Devstral-Small-2-24B-Instruct-GLQ-4bpw")
tok = PreTrainedTokenizerFast(tokenizer_file=f"{path}/tokenizer.json")
tok.pad_token, tok.eos_token, tok.bos_token = "<pad>", "</s>", "<s>"
model = AutoModelForCausalLM.from_pretrained(
    "xv0y5ncu/Devstral-Small-2-24B-Instruct-GLQ-4bpw",
    device_map="cuda", dtype="float16",
)
```

`examples/inference_hf.py` includes a `load_tokenizer()` helper that
handles this automatically.

### transformers compatibility

For models ≤ 1B parameters use `transformers >= 5.0`. Transformers
4.57.x has a weight-loading bug that produces garbage output for small
GLQ models. Larger models (3B+) work with both 4.x and 5.x.

## Inference kernels

`glq/inference_kernel.py` + `glq/csrc/glq_cuda.cu` provide CUDA C and
Triton kernels that compute `Y = X @ dequant(W)^T` without
materializing the weight matrix. Each kernel iterates over `N/8`
codebook blocks per output row, gathers 8-D vectors from the L2-cached
codebook, and accumulates the matmul directly against indices.

| Path | When | Notes |
|------|------|-------|
| **CUDA C Tensor Core** | B ≥ 2 (prefill) | inline PTX `mma.sync` against codebook-loaded registers; 3-5× faster than Triton |
| **CUDA C split-K matvec** | B = 1 (decode) | 4 rows/warp + `__shfl_xor_sync` reduction; 2.7× faster than Triton |
| **CUDA C shared-mem FHT** | RHT step | double-buffered butterfly; low 5 stages as warp shuffles (v0.7.1) |
| **Trellis 3INST fused linear** | trellis checkpoints, B = 1 | lookup-free decode with the input RHT + cast computed in-block — one kernel per linear (v0.7.1) |
| **Shard-batched output RHT** | trellis, fused QKV / gate-up | per-shard output transforms in one `grid.y` launch (v0.7.1) |
| **Triton fallback** | no `ninja`, or `n_pad > 32 768` | always available |

**Bit-exact determinism.** Every kernel reduces partial sums in a fixed
order (scratch-buffer split-K or in-block reduction) instead of
`atomicAdd` across k-splits, so running the
same prompt at B=1 decode or B=8 prefill produces identical logits
across runs — required for reproducible lm-eval scoring and
on-policy RL rollouts.

Direct kernel access:

```python
from glq.inference_kernel import glq_dequant_matmul
y = glq_dequant_matmul(x, Qidxs, codebook, Wscale,
                       Qidxs2=Qidxs2, codebook2=codebook2,
                       inv_resid_scale=inv_rs)  # 3/4 bpw two-stage
```

## Architecture

```
glq/
  codebook.py          # E8ShellCodebook: enumeration, encode/decode
  hadamard.py          # Fast Walsh-Hadamard Transform
  rht.py               # Randomized Hadamard Transform
  ldlq.py              # Block-LDL quantization with error feedback
  quantize_model.py    # Full model pipeline + CLI
  quantized_linear.py  # E8RHTLinear: drop-in nn.Linear replacement
  inference_kernel.py  # Triton kernels + CUDA dispatch
  csrc/glq_cuda.cu     # CUDA C kernels (split-K matvec, TC, FHT)
  hf_integration.py    # HuggingFace Transformers integration
  kv_cache.py          # INT8 quantized KV cache
  cuda_graph.py        # B=1 decode wrapper
glq_vllm/              # vLLM integration: weight + KV cache (v0.3.0+)
```

## Acknowledgments

This project builds on **[QuIP#](https://github.com/Cornell-RelaxML/quip-sharp)**
(Tseng et al., 2024). The default E8-shell pipeline is *inspired by* its
Hadamard-incoherence + lattice-codebook formulation, and the optional
`--codebook e8p` path is **derivative work that ports QuIP#'s E8P (padded-D̂8)
codebook and its tensor-core decode / residual kernels** — the `grid_packed_abs`
codebook plus the `decode_matvec_e8p`, `decompress_packed_e8p`, and E81B
lookup-matmul kernels in [`glq/csrc/glq_e8p.cu`](glq/csrc/glq_e8p.cu). All credit
for the E8P codebook and those kernels belongs to the QuIP# authors; this
repository is an independent port, not an official QuIP# release. See the
[QuIP# repository](https://github.com/Cornell-RelaxML/quip-sharp) for the paper
and citation.

The `--codebook trellis` path is likewise **derivative of
[QTIP](https://github.com/Cornell-RelaxML/qtip)** (Tseng et al., 2024): the tail-biting
trellis formulation, the hybrid lookup codebook, the "3INST" lookup-free decode idea, and
the bit-unpack structure of the kernels in
[`glq/csrc/glq_trellis.cu`](glq/csrc/glq_trellis.cu) all originate there — GLQ grafts
them onto its own RHT + LDLQ pipeline and fused-linear kernel architecture. All credit
for TCQ and the 3INST decode belongs to the QTIP authors; see the
[QTIP repository](https://github.com/Cornell-RelaxML/qtip) for the paper and citation.

Other foundations:

- E8 lattice: Korkin & Zolotarev (1872); Gosset (1900); Conway & Sloane, *Sphere Packings, Lattices and Groups*; [Viazovska (2016)](https://arxiv.org/abs/1603.04246) — sphere-packing optimality in 8 dimensions.
- Block-feedback quantization: [GPTQ](https://arxiv.org/abs/2210.17323) (Frantar et al., 2022).
- INT8 KV cache: [KIVI](https://arxiv.org/abs/2402.02750) (Liu et al., 2024).

## License

GNU General Public License v3.0 (GPL-3.0). See [LICENSE](LICENSE).
