Benchmarks

One machine (RTX 4090, RunPod), measured across the top local LLMs (Qwen, Llama, Mistral, DeepSeek, Gemma-2), every number measured and reproducible. No invented figures: if a metric is absent from the sources, it is absent here too. New here? Start with the plain-language overview of what Trapetum does.

November 2026 update: corrected 4-bit perplexity, faster decode, exact speculative decoding, Llama-2-70B and Llama-3.1-70B-Instruct on one RTX 4090

Measured from 5 to 8 November 2026 on rented RunPod GPUs, one GPU per run, each run on a fresh pod, so hosts differ between runs. An NVIDIA GeForce RTX 4090 (24 GB, sm_89) ran every step; an NVIDIA A40 (sm_86) also ran the first and the final fusion validations and one Python-only sparsity diagnostic. Batch 1, greedy decoding, CUDA-graph decode (TRAPETUM_GRAPH=1) unless a figure is marked eager. Three timing windows appear: the mean per 128-position bin of one long greedy decode (TRAPETUM_SWEEP), the 16-token continuation after a 6-token prompt (positions 6 to 22, the window of the CBK4 table under model-level results), and a 512-token decode benchmark (the Llama-2-70B 2x8 and Llama-3.1-70B-Instruct runs also report a 256-token decode, and the speculative runs an end-to-end throughput with prefill). Each per-step gain is a same-host A/B; figures from different steps come from different hosts, and cross-host comparisons are marked as such. Apart from the exporter A/B, every Llama-2-7B 4-bit speed below was measured on the new exporter's file. The code is on the repository's main branch (commits 2e8c170 to 6b8cb42) and is not in an app release yet: the installed app and its downloadable models keep the previous runtime and file format. The November results are summarized on this page; their raw logs are not in the repository.

Correction. This page listed the scalar 4-bit runtime file (CBK3) of Llama-2-7B at 3.66 GB and PPL 6.34, and said the 2-bit file trades +0.33 PPL for its size. The 6.34 came from a research script (mid-grid k-means start, LM head kept in fp16), not from that file. Read back from the file on 5 November 2026, the 4-bit file written by the runtime's exporter until now scores 6.79 in the repository's 30-window harness (6.35 on the full wikitext-2 test set). The +0.33 PPL statement compared the 2-bit checkpoint with the research script's 6.34 and is withdrawn: against the real file, the 2-bit file was slightly better in quality. The exporter now writes a better 4-bit file, 6.10 (5.72 on the full test set), and against that file the 2-bit file is 0.57 PPL behind. Earlier versions also gave this file as 3.5 GB and 3.9x smaller than fp16, a MiB figure written as GB: it is 3.66 GB (legacy exporter) or 3.75 GB (new exporter), 3.7x or 3.6x smaller than the 13.48 GB of fp16 weights.

Llama-2-7B files read back from disk: wikitext-2 perplexity

Llama-2-7B Format File PPL, 30 windows PPL, full test set
fp16 reference fp16 - 5.83 5.47
Scalar 4-bit, legacy exporter (current app models) corrected CBK3, 4-bit LM head 3.66 GB 6.79 6.35
Scalar 4-bit, new exporter (the new default) new CBK4, 8-bit LM head 3.75 GB 6.10 5.72
AQLM PV 2x8, 2-bit CBK4, 8-bit LM head 2.05 GB 6.67* 6.26*
Research script, scalar 4-bit, LM head in fp16 (the 6.34 previously shown for the file) no file - 6.34 -

Harnesses: the repository's harness averages the first 30 windows of 2,048 tokens of the wikitext-2 test set (fp16 5.83); the full harness covers all 166 full windows of 2,048 tokens of the 341,469-token test set (fp16 5.47). File sizes: 3,655,827,760, 3,752,083,768 and 2,050,610,744 bytes. * Measured on the PV checkpoint (6.666); the exported 2-bit file reads back at 6.665 in the 30-window harness and was not run on the full test set.

The new exporter runs k-means from a mid-grid start, moves an empty centroid to 0, runs 10 iterations, weights every centroid update by E[x2] of its input channel over 128 wikitext-2 train windows, and stores an 8-bit LM head as CBK4. The file is 2.6% larger, and the gap to fp16 in the 30-window harness falls from 0.96 PPL to 0.27, 72% smaller. In the legacy file the 4-bit LM head alone cost about 0.30 PPL (6.49 with an fp16 head, 6.79 as written); the 8-bit head costs 0.0004 (6.0978 with an fp16 head, 6.0982 as written). Worst relative error of the prompt logits against the Hugging Face reference: 2.0e-3 for the new file, 5.1e-3 for the legacy one. Unused centroid slots across the 224 projections: 1.19% in the legacy file, 0.004% in the new one.

Speed cost of the new file with the runtime published on 5 November (630168e), same RTX 4090 host, two rounds: 1.1 to 1.6% slower (CUDA graph at positions 6 to 22: 150.6 against 153.0 tok/s; eager: 138.1 against 140.3; CUDA graph over a 256-token window: 107.3 against 108.5), peak GPU memory 4,246 against 4,172 MiB, and greedy 16/16 against the reference in all 8 runs.

2-bit against 4-bit: against the new 4-bit file, the 2-bit file is 0.57 PPL behind (6.665 against 6.098), for 45% less disk (2.05 against 3.75 GB) and 36% less peak GPU memory (2,736 against 4,246 MiB with CUDA graphs). Against the legacy 4-bit file it was already slightly better in quality, 0.13 PPL from the unrounded values (6.665 against 6.791), while 44% smaller (2.05 against 3.66 GB).

Scope: the new recipe's perplexity was measured on Llama-2-7B only; the AQLM, DeepSeek and Gemma exporters are unchanged; runtimes older than this change reject the new files. The 8-bit LM head is chosen automatically only on GPUs that can run its kernel (69,632 bytes of shared memory per block), and the runtime refuses such a layer at load time on Turing and Pascal GPUs (verified by reading the code; no Turing GPU was available to test). The repository's current llama_quant.py no longer reproduces the research script's 6.34: since commit c9d8b72 (2026-06-28) it starts its centroids from a min-to-max grid and gives 6.41.

Decode speed versus context: split-KV attention

Before this work, decode attention walked the KV cache one position at a time in one block per head, about 20.6 microseconds per cached position per token on Llama-2-7B (RTX 4090, CUDA graph). At 1,024 positions, 73% (4-bit) and 78% (2-bit) of the decode time depended on the context, against a threshold of 30% fixed before the run. Split-KV decode attention cuts the KV cache into 64-position chunks; one block per (head, chunk) runs an online softmax, and the chunks merge in a fixed order with no atomics, so results are bitwise reproducible. No decode step allocates during CUDA-graph capture, and TRAPETUM_ATTN=seq keeps the old kernel.

Llama-2-7B, decode tok/s First 128 positions Positions 1,024 to 1,152 Positions 3,968 to 4,096
4-bit, sequential attention (before, host A) 131.1 35.0 11.2
4-bit, sequential attention (host B, same-host control) 129.7 34.3 not run
4-bit, split-KV (host B) new 157.5 146.2 115.8
2-bit, sequential attention (before, host A) 167.8 37.1 11.4
2-bit, split-KV (host B) new 214.5 194.4 143.9

RTX 4090, batch 1, greedy, CUDA graph, mean per 128-position bin of one long greedy decode; 4-bit is the new exporter's file, 2-bit the AQLM PV 2x8 file. Hosts A and B are two different RTX 4090 pods. Same host: 4.3x at 1,024 positions for 4-bit (146.2 against 34.3 tok/s) and 18% less time per token over the first 128 positions (6.35 against 7.71 ms). Across hosts: 4-bit 10x at 4,096 positions (115.8 against 11.2), 2-bit 5x at 1,024 and 13x at 4,096. Validation: 88 of 88 cases match an f64 CPU reference (worst relative error 2.2e-4), each one deterministic, across MHA and GQA, head_dim 64, 128 and 256, and lengths 1 to 4,096 around the chunk boundaries; greedy 16/16 against the reference in every generate run. The attention kernel alone (32 heads, head_dim 128) takes 10.0 microseconds per call at 1,024 positions, against 569.5 for the old kernel.

Short-context decode, step by step (Llama-2-7B, 4-bit)

Fused decode kernels for scalar 4-bit projections: q, k and v in one GEMV; gate and up in one GEMV with SiLU(gate) times up applied in the reduce; the residual add inside the o and down reduces; RoPE and the KV-cache append in one launch. About 28 launches per layer become 13 (counted in the code; no profiler run), and TRAPETUM_FUSE=0 restores the separate kernels. The fused path produces prompt logits byte-identical to the separate kernels (checked on an RTX 4090 and on an A40), greedy 16/16 in every run, and a GQA model with head_dim 64 matches a CPU reference (relative error 7.9e-4, top-1 5/5). Then the 4-bit codebook GEMV moved to tensor cores (mma.sync m16n8k16, f16 inputs, f32 accumulation, weights read straight from the 4-bit index bytes with no repacking or extra memory, partial sums added in a fixed order): one kernel shared by decode and speculative verify.

Step (same-host A/B) Timing window Before After Change
Split-KV attention first 128 positions 7.71 ms (129.7 tok/s) 6.35 ms (157.5 tok/s) 18% less time
Fused decode kernels first 128 positions 6.36 ms (157 tok/s) 5.28 ms (189 tok/s) 17% less time
Tensor-core 4-bit GEMV (sm_80 and newer) new 512-token generation 5.42 ms (184.3 tok/s) 5.14 ms (194.5 tok/s) +5.5% tok/s
Activation sparsity 25% (opt-in, lossy, +0.12 PPL) first 128 positions 5.44 ms (184 tok/s) 4.69 ms (213 tok/s) 1.16x faster

RTX 4090, CUDA graph, the new exporter's 4-bit file; each row is a same-host A/B on its own host (one run per kernel for split-KV, two rounds for fusion and sparsity, three alternated runs per build for the tensor-core row). On an A40 the fusion A/B gave 10.49 to 9.31 ms per token. A first fused version that read the projections from a table passed by value was about 1.8 ms per token slower on an A40; passing them as scalar arguments made it fast. Over the 16-token continuation (positions 6 to 22) on the tensor-core host: 184.6 to 205.0 tok/s with CUDA graphs and 175.8 to 195.4 tok/s eager. Greedy stays 16/16 against the Hugging Face reference, with a worst prompt-logits error of 1.94e-3 (1.98e-3 before), and TRAPETUM_TC=0 and the sparse path give prompt logits byte-identical to the previous commit. Sparsity is not part of the default path; the sparse path keeps the FMA kernels.

Cumulative, across hosts: at short context the step went from 7.63 to 7.71 ms per token (130 to 131 tok/s) before split-KV to 5.28 to 5.44 ms (184 to 189 tok/s) after split-KV and fusion, +40% to +46% depending on the host pair. At 1,024 positions it went from 34.3 to 35.0 tok/s to 170.9 tok/s, about 5x. At 4,096 positions only split-KV has been measured (115.8 tok/s against 11.2, 10x); no measurement exists at that length after fusion or the tensor-core GEMV. For scale, one decode path with bit-identical prompt logits measured 184.6, 185.5 and 187.8 tok/s (positions 6 to 22, CUDA graph) on three RTX 4090 hosts.

Opt-in activation sparsity (lossy)

TEAL-style: TRAPETUM_SPARSITY=<file> loads per-layer thresholds; inputs of the q/k/v, o, gate/up and down projections below their threshold become zero, and the 4-bit GEMVs skip the index rows of zero inputs. It is off by default, lossy and CUDA only; without a file the kernels and their bits are unchanged, and Metal and the batched speculative path refuse it. Position 0, the attention sink, always runs dense: sparsifying it as well cost +0.66 PPL at 25% sparsity, against +0.12 with it kept dense. The chosen setting, 25% on every site, decodes 1.16x faster over the first 128 positions (5.44 to 4.69 ms per token) and 1.15x at 1,024 positions (5.85 to 5.08 ms), for +0.12 PPL in the 30-window harness (6.10 to 6.22) and +0.11 on the full test set (5.72 to 5.83). That perplexity is computed in Python with the 4-bit file's weights; the runtime's own NLL on the first 2,048 wikitext-2 test tokens reproduces the simulated loss (+0.0174 against +0.0164 nats per token), and the two agree within 2.4e-5 nats when dense. Per-site mixes with the o and down inputs at 50 to 60% reach 1.19 to 1.26x but cost +0.17 to +0.30 PPL. Sparse decoding changes the output: at 25% the 16-token greedy continuation matches the dense reference for 3 of 16 tokens.

Batch-invariant speculative verify

A draft of M = 1, 2, 4, 8 or 16 tokens is verified in one forward whose logits are, row by row, the same bits as M decode steps, so a speculative decoder built on it emits exactly what plain greedy decoding emits. 15 of 15 checks pass (M = 1, 2, 4, 8 and 16 at bases 6, 56 and 1,016, attention over up to 17 chunks, CUDA graphs replayed at other bases, state scrambled before every checked call, next decode step compared), also with TRAPETUM_FUSE=0 and with TRAPETUM_TC=0. Verify cost relative to one decode step (Llama-2-7B, CUDA graph, short context, mean of two rounds): 1.05x at M=2, 1.06x at M=4, 1.12x at M=8 and 1.21x at M=16. The cost grows with context: 1.46x at M=16 near 1,100 positions. An FMA-only verify, kept as the fallback, reached the same bits but cost 3.27x a decode step at M=16 (3.10x after a second attempt), in two earlier runs on other hosts.

PARD speculative decoding: Llama-3.1-8B-Instruct, exact greedy output, 2.51x

PARD (arXiv 2504.18583) now runs in the Rust runtime in target-independent mode, with the PARD-2 drafter (arXiv 2605.08632). The drafter proposes K tokens in one forward over the new real tokens plus K-1 mask tokens; the target checks the last token and the K drafts in one batch-invariant verify; the longest matching prefix plus the target's next token are emitted. Both caches are addressed by position, so a rejection needs no copy. Prompts are prefilled through the same verify graphs (chunks of 16, then 8, 4, 2 and 1 rows), whose rows and KV-cache rows carry the decode step's bits, so a prefilled prompt continues exactly like a decoded one; plain and speculative runs share this prefill.

Decision run. Target Llama-3.1-8B-Instruct in 4-bit (new exporter defaults); drafter amd/PARD2-Llama-3.1-8B (1.2 billion parameters, 16 layers) in 4-bit with a 4-bit head; K=7, so the target verifies 8 rows; at most 256 new tokens with EOS stop; RTX 4090; CUDA graphs in both arms; 260 prompts in the chat template with a system prompt, up to 437 tokens long; two repetitions with the arm order alternated; both arms on the same file in the same process. 520 of 520 speculative outputs are identical, token for token up to EOS or 256 tokens, to plain greedy decoding of the same 4-bit target.
Prompt set Prompts Plain tok/s PARD tok/s Speedup Tokens per verify round
GSM8K80166.4451.02.71x4.55
HumanEval80163.6471.82.88x5.13
MATH-50020165.4449.12.72x4.62
MT-Bench, first turns80167.0340.02.04x3.21
All four sets, pooled (520 runs, 108,102 tokens) new260165.5415.62.51x4.19
Pooled, drafter with an 8-bit head (one repetition)260165.4381.92.31x4.22

Speed is end to end with prefill included: total new tokens over total wall time from prefill start to last token. Tokens per round include the target's bonus token. GSM8K, HumanEval and MATH-500 prompts come from PARD's repository; the MT-Bench prompts are FastChat's first turns. In a preliminary acceptance check with PARD's own code, MT-Bench acceptance came out 22.5% below the PARD-2 paper's vLLM figure, while GSM8K, MATH-500 and HumanEval were within 5%. With the 8-bit drafter head, 260 of 260 outputs are identical; per set 2.485x (GSM8K), 2.651x (HumanEval), 2.501x (MATH-500) and 1.874x (MT-Bench).

The target file: Llama-3.1-8B-Instruct through the new exporter (weighted k-means, 8-bit LM head; 5,041 MiB, about 5.29 GB) decodes on the tensor-core GEMV with prompt logits within 1.93e-3 relative error of the Hugging Face reference and top-1 agreement 6/6. Its 16-token greedy continuation follows the reference for 4 tokens and then splits on a near-tie (top-two logit margin 0.004 at the fifth token). Plain decode: 173.5 tok/s (5.76 ms per token), CUDA graph, 512-token generation (175.2 tok/s in an earlier run on another host). Batch-invariant verify on this GQA model: 15 of 15 bit checks pass (M = 1 to 16 at bases 6, 56 and 1,016); verifying 8 tokens costs 1.19x a decode step and 16 tokens 1.34x to 1.36x at short context, rising to 1.58x near 1,050 to 1,170 positions. The 4-bit drafter's 16-row forward costs 0.27 of a target decode step with a 4-bit head and 0.43 with an 8-bit head.

For context, the July two-model table under speculative decoding measured plain decode of Llama-3.1-8B at 117.9 tok/s, with a Llama-3.2-1B drafter at parity (0.85x to 0.90x). That table used another harness and the previous runtime, so the comparison is not a controlled A/B.

Llama-2-70B in 2 bits on one RTX 4090

The 2-bit path (CBK4 additive codebooks) also serves ISTA-DASLab/Llama-2-70b-AQLM-2Bit-2x8-hf, an AQLM 2x8 checkpoint (not PV-tuned), on a single RTX 4090 (24 GB). The exporter copies the AQLM codes, codebooks and scales bit for bit and quantizes the fp16 LM head to 8 bits, a 17.96 GB file; the runtime keeps the embedding off the GPU (524 MB in fp16 in the file, about 1.05 GB as f32 in host RAM, one row uploaded per token). Llama-2-70B uses the Llama-2-7B tokenizer, so the perplexities compare directly. Measured on 7 November 2026, one run on one host, against a rule fixed before the run; all five criteria passed (the model fits on the card; prompt logits within 5e-3 relative error of the reference, top-1 on every prompt position and greedy 16/16 in eager and CUDA-graph mode; the runtime's own NLL within 0.005 nats of the reference; full-test perplexity at most 4.70; at least 10 tok/s).

Llama-2 File PPL, 30 windows PPL, full test set Decode tok/s, positions 6 to 22
7B, fp1613.48 GB5.835.47-
7B, 4-bit, new exporter3.75 GB6.105.72205.0
70B, 2-bit AQLM 2x8 new17.96 GB4.684.4938.4

Wikitext-2, sequence length 2048, computed in Python on the served weights (AQLM codes and the 8-bit head). Decode is batch 1 with CUDA graphs on an RTX 4090; the two models ran on different hosts.

Agreement with the reference: against the Hugging Face AQLM model carrying the same 8-bit head, prompt logits are within 1.77e-3 to 2.57e-3 relative error (top-1 6/6) and the 16-token greedy continuation matches, in eager mode and in three CUDA-graph runs; the smallest top-two logit margin along it is 0.17. On the first 2,048 wikitext-2 test tokens the runtime's own mean NLL is 1.228539 nats per token against 1.228504 for the reference. Memory: with a 1,024-position KV cache, decoding peaks at 17,378 MiB of the card's 24,564 MiB (17,698 MiB with the 1,216-position cache of the decode to 1,152 positions, 18,004 MiB with the 2,112-position cache of the 2,048-token NLL check). Speed: 38.3 and 38.6 tok/s at positions 6 to 22 (two rounds), 38.1 over a 256-token decode, 37.5 near 1,024 to 1,152 positions; about 17.4 GB of weights read per token, so about 670 GB/s, two thirds of the card's 1,008 GB/s. On the same host the 7B 2-bit PV file decodes at 213.4 and 213.9 tok/s.

The checkpoint is 2x8 and not PV-tuned (no PV-tuned 2x8 Llama-2-70B is published): at context 4,096 the model cards give 4.21 for it, 3.83 for the 1x16 checkpoint and 3.78 for the PV-tuned 1x16 one, which the runtime now also serves (next section). The 4.06 of the 2-bit 70B in the multi-method comparison below is that 1x16 checkpoint (not PV-tuned), run in Python on an H100, not this file.

Llama-2-70B PV-tuned in 1x16, with exact speculative decoding

The PV-tuned Llama-2-70B is published only in AQLM's 1x16 layout, ISTA-DASLab/Llama-2-70b-AQLM-PV-2Bit-1x16-hf (AQLM, Egiazarian et al., ICML 2024; PV-Tuning, Malinovskii et al., NeurIPS 2024): each weight matrix has one codebook of 65,536 entries of 8 fp16 values and one 16-bit code per group of 8 weights, still 2 bits per weight. The exporter copies these codes, codebooks and scales bit for bit too (8-bit LM head as above), an 18.54 GB file. A 1 MiB codebook does not fit in shared memory, so, like AQLM's own 1x16 CUDA kernel, the decode kernel gives each output row to one warp and gathers every codebook entry from the L2 cache; it computes the products in f32 rather than fp16 and uses no atomics, so its results are the same bits from run to run, eagerly and in CUDA graphs. Measured on 7 November 2026 on one RTX 4090, in two runs, each against a rule fixed before it.

Llama-2, 2-bit AQLM File PPL, 30 windows PPL, full test set Decode tok/s
13B PV, 1x164.00 GB5.745.4042.6
70B, 2x8 (above)17.96 GB4.684.4938.4
70B PV, 1x16 new18.54 GB4.174.028.5
70B PV, 1x16, speculative (K = 7) new+ 2.05 GB drafter4.174.0222.5

Wikitext-2, sequence length 2048, computed in Python on the served weights; for scale, the new 7B 4-bit file (3.75 GB) scores 6.10 and 5.72. Decode is batch 1 with CUDA graphs, positions 6 to 22, except the speculative row: end-to-end throughput, prefill included, over 12 prompts of 8 to 65 tokens with up to 128 generated tokens each (plain decoding in the same harness: 8.1 tok/s). The 2x8 row comes from another host.

Agreement with the reference: prompt logits are within 1.62e-3 relative error of the Hugging Face AQLM model on its dequantize-then-matmul path (top-1 6/6), and the 16-token greedy continuation matches in eager and CUDA-graph mode. Against its default path the gap is 5.91e-3: aqlm 1.1.7 runs a prompt of 6 tokens or fewer through its GEMV kernel, which adds the 8 products of each group in fp16, and that path is itself 4.96e-3 away from the dequantized one. On the first 2,048 wikitext-2 test tokens the runtime's mean NLL is 1.132485 nats per token against 1.132619 for the reference (13B: 1.362944 against 1.362946). Peak memory: 17,930 MiB with a 1,024-position KV cache (18,560 MiB with the 2,112-position cache of the NLL check).

Plain decode is slow: 8.5 tok/s at positions 6 to 22, and 8.4 up to 1,152 positions. Every group of 8 weights costs a random 16-byte codebook read that L2 serves as a 32-byte sector, and counting the whole decode step the 7B PV 1x16 and the 70B move about 2.0 and 2.3 TB/s of such sectors. On the same host the 7B PV 1x16 decodes at 78 tok/s against 215 for its 2x8 version.

Exact speculative decoding: the 2-bit ISTA-DASLab/Llama-2-7b-AQLM-PV-2Bit-2x8-hf (2.05 GB, same tokenizer) drafts K tokens with CUDA-graph decode steps, and the 70B verifies the K + 1 rows in one pass with a kernel that gives every row the decode kernel's bits and reads each code and codebook entry once for all rows (at short context 8 rows cost 1.11 decode steps and 16 rows 1.81; near 1,050 to 1,070 positions, 1.14 and 1.87). The verify rows are identical to decode steps at three bases, for M = 1, 2, 4, 8 and 16, eager and CUDA graph. Over 12 prompts (8 windows of 64 wikitext-2 test tokens and 4 short prompts, BOS included), 128 greedy tokens each at most, the outputs equal plain greedy decoding in 36 of 36 runs (K = 3, 7, 15). The throughput, end to end from the start of the prefill to the last token (the drafter's prefill, one decode step per prompt token, included), is 22.5 tok/s at K = 7 against 8.1 for plain decoding in the same harness (2.79x, 4.49 tokens per verify pass), 18.3 at K = 3 and 17.1 at K = 15. Peak memory with both models and the verify buffers: 21,295 MiB.

The first run met its criteria of fit, top-1, greedy, NLL agreement and quality (full-test perplexity at most 4.20) but not its speed criterion (at least 10 tok/s), and its logits criterion (5e-3) failed against the default reference only. The second run met all four of its criteria: identical outputs, verify rows identical to decode steps, at least 12 tok/s at K = 7, and logits within 5e-3 of the dequantized reference.

Llama-3.1-70B-Instruct on one RTX 4090, with exact speculative decoding

The same 1x16 path serves the conversational ISTA-DASLab/Meta-Llama-3.1-70B-Instruct-AQLM-PV-2Bit-1x16 (PV-tuned, 1x16, vocabulary of 128,256 tokens), with ISTA-DASLab/Meta-Llama-3.1-8B-Instruct-AQLM-PV-2Bit-2x8-hf (2x8, 3.46 GB file) as the drafter; the 70B file is 21.00 GB with the 8-bit LM head. The 70B checkpoint's current configuration declares no RoPE scaling (8,192 positions), unlike Meta's Llama 3.1 and the checkpoint's own first revision on Hugging Face (llama3 scaling, 131,072 positions); it is served as published: the runtime, like Hugging Face, uses plain RoPE with theta 500,000 (the 8B drafter declares llama3 scaling, which the exporter bakes into its file). On the 24 GB card the export keeps the 2.1 GB input embedding on the CPU, computes the perplexity loss over chunks of positions, frees the dense LM head before loading its dequantized copy and transposes the head on the CPU for its quantization; its highest memory sample (every 300 ms, with PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True) was 23,091 MiB of 24,564. Measured on 8 November 2026 on one RTX 4090, one run, against a rule fixed before it; all seven criteria passed (fit; logits within 5e-3 of the dequantized-path reference with greedy 16/16; runtime NLL within 0.005 nats of the reference; full-test perplexity at least 15% below the drafter's; outputs identical to plain greedy decoding; verify rows identical to decode steps; at least 12 tok/s at K = 7).

Llama 3.1 Instruct, 2-bit AQLM File PPL, 30 windows PPL, full test set Decode tok/s
8B PV, 2x8 (drafter)3.46 GB11.2810.61183.2
70B PV, 1x16 new21.00 GB6.626.288.7
70B PV, 1x16, speculative (K = 7) new+ 3.46 GB drafter6.626.2821.3

Llama 3 tokenizer (141 windows of 2,048 tokens), so these perplexities do not compare with the Llama 2 ones, and no unquantized model was measured. Decode is batch 1 with CUDA graphs, positions 6 to 22, except the speculative row: end-to-end throughput, prefill included, over 16 chat prompts of 53 to 178 tokens with up to 128 generated tokens each (plain decoding in the same harness: 7.9 tok/s).

Agreement with the reference: prompt logits are within 1.40e-3 relative error of the Hugging Face AQLM model on its dequantize-then-matmul path (1.61e-3 against its default path; the two paths are 2.14e-3 apart), top-1 6/6, and the 16-token greedy continuation matches in eager mode and in two CUDA-graph runs. On the first 2,048 wikitext-2 test tokens the runtime's mean NLL is 1.593123 nats per token against 1.593003 for the reference. Peak memory: 18,730 MiB with a 1,024-position KV cache (19,360 MiB with the 2,112-position cache of the NLL check), 21,774 MiB with the drafter and the verify buffers.

Plain decode: 8.7 tok/s at positions 6 to 22 (two rounds) and over a 256-token decode. Over 16 chat prompts (the first 4 of PARD's gsm8k, humaneval and math_500 sets and of the MT-Bench first turns, with the Llama 3.1 chat template and the system prompt "You are a helpful assistant."), the speculative outputs equal plain greedy decoding in 32 of 32 runs (K = 3 and 7), and the end-to-end throughput is 21.3 tok/s at K = 7 against 7.9 for plain decoding in the same harness (2.7x, 5.59 tokens per verify pass), and 17.0 at K = 3. At K = 7: 22.75 tok/s on gsm8k, 20.87 on humaneval, 23.88 on math_500 and 18.67 on MT-Bench (4.10 tokens per pass). At short context a verify pass of 8 rows costs 1.13 decode steps and one of 16 rows 1.85 (1.16 and 1.92 near 1,050 to 1,070 positions), and the verify rows are identical to decode steps at three bases, for M = 1, 2, 4, 8 and 16, eager and CUDA graph. The drafter alone decodes at 183.2 tok/s.

Limits of this update

Software: the pytorch/pytorch:2.4.0-cuda12.1-cudnn9-devel image (Python 3.11.9, PyTorch 2.4.0, CUDA 12.1) on every pod, Rust 1.99.1 on every pod that built the runtime, and NVIDIA driver 595.91.07 on the RTX 4090 hosts whose logs record it.

End-to-end Pareto: speed, memory, accuracy, energy (June 2026)

Batch-1 decode, all decode metrics measured over a 512-token generation (iso-context). Wikitext-2 PPL over 30 000 tokens, context 2048 (generation-independent). Energy via pynvml power.draw. The Rust runtime row uses the same 4-bit projection weights as the Python row, with no Python at runtime (its LM head is quantized too, see the PPL notes). You can test these numbers yourself in the live compare demo, 4-bit versus fp16 on two identical RTX 4090s. Newer decode kernels are in the November 2026 update.

Iso-context note. Every row was measured over the same 512-token generation so that energy and throughput numbers are directly comparable across methods. J/token is context-dependent (in this June run it drops at longer context for all methods because attention over the KV cache is memory-bound and lower-power; the November split-KV attention changes long-context decode, and energy was not re-measured). Comparing rows at different context lengths would be misleading, so all decode metrics are measured at the same 512-token window.
Method Bits VRAM (GB) Wikitext PPL Decode tok/s J/token gCO2/1k FR gCO2/1k US
fp16 (cuBLAS baseline) 16.0 13.48 5.28 45.2 5.45 0.076 0.61
Trapetum 4-bit, Python op ours 4.05 3.81 5.92 28.4 4.00 0.056 0.44
aqlm-2bit (official kernel) 2.0 2.15 7.01 24.2 3.86 0.054 0.43
Trapetum 4-bit, Rust runtime ours 4.05 4.73* 5.92† 81 2.58 0.036 0.29

* 4.73 GB measured peak (3.66 GB weights + KV cache + activations). The short-context Rust throughput reached 135 tok/s in June 2026; the November 2026 kernels are in the update.

† The Rust row's 5.92 is the Python row's figure: bench/pareto.py's quantization of the 224 projections, evaluated with the LM head kept in fp16. The runtime's .cbk file quantizes the LM head too, and its perplexity was not measured in this run; perplexities read back from the files are in the November 2026 update.

PPL protocol: RTX 4090 energy run, perplexity measured over the 512-token generation context (fp16 5.28, codebook-4bit 5.92, AQLM 2-bit 7.01). The AQLM 2-bit row uses the 2x8 checkpoint (ISTA-DASLab Llama-2-7b-AQLM-2Bit-2x8-hf).

gCO2/1k-tok is a derived projection, not a measurement: gCO2 = J/token x grid intensity at France-like (~50 gCO2/kWh) or US-like (~400 gCO2/kWh) intensity. RunPod's actual grid mix is unknown. J/token is what is measured.

Pareto chart: decode speed vs perplexity, bubble area proportional to VRAM, Rust runtime highlighted in orange
RTX 4090, Llama-2-7B, batch-1, 512-token iso-context. The Rust runtime (orange) is the fastest and lowest-energy path at 4-bit. fp16 is most accurate but 2.9x heavier in VRAM (13.5 GB fp16 vs 4.7 GB Rust runtime; the 4-bit weights file is 3.66 GB, 3.7x smaller than the 13.48 GB of fp16 weights).

Reading the table honestly

Speculative decoding: lossless throughput on top

The same bandwidth law has a second payoff. At batch 1 the arithmetic units sit idle while the weights stream from memory, so the one weight read that makes a token can just as well verify several. A small drafter proposes K tokens; the target checks all K+1 in a single forward. Because decode is bandwidth-bound, that verify is nearly free. The output is token-for-token identical to plain greedy decode, lossless, not an approximation.

Measured on RTX 4090. Acceptance rate α (fraction of positions where the drafter's greedy token matches the 4-bit Llama-2-7B target) between the target and two 4-bit drafters, with the speedup projected by the standard estimate S(K) = (1-αK+1)/(1-α) ÷ (1 + K·tdraft/ttarget). The batched verify is validated lossless on CUDA (K = 1, 2, 3) and on Apple Metal. Since November 2026 the CUDA runtime also has a batch-invariant verify for drafts of 1, 2, 4, 8 or 16 rows, used by the PARD benchmark (see the update); the batched verify above is unchanged.
Drafter Acceptance α K = 1 K = 2 K = 3 Best
TinyLlama-1.1B 0.734 1.23x 1.25x 1.20x K=2 → 1.25x
llama-160m best 0.648 1.38x 1.50x 1.49x K=2 → 1.50x

The smaller 160M drafter wins: lower acceptance (0.648 vs 0.734) but so much cheaper per draft that the net speedup is higher, 1.50x vs 1.25x. Draft cost, not raw acceptance, is the binding constraint, and K=2 (verify three tokens per target forward) is optimal for both. This lossless throughput is orthogonal to and compounds with the quantization memory and energy savings above.

Wall-clock, two models, measured (July 2026)

The real harness (drafter decoding incrementally in its own KV cache), after fixing the verify kernel (see below). 128 tokens, lossless at every K:

Pair (target + drafter)VocabPlain tok/sK=1K=2K=3Verdict
Llama-2-7B + llama-160m32k124.61.19x1.32x1.24xWIN, 164 tok/s
R1-Distill-Qwen-7B + Qwen2.5-1.5B152k100.60.76x0.76x0.76xparity
Llama-3.1-8B + Llama-3.2-1B128k117.90.85x0.90x0.80xparity

The 32k-vocab pair wins as projected (1.32x measured vs the 1.50x ceiling). In this July harness, with the previous runtime, the large-vocab pairs sat at parity because every drafter forward read back the full logits (500+ KB) for a host argmax. In November 2026, PARD speculative decoding on Llama-3.1-8B-Instruct (128k vocabulary) measured 2.51x end to end with output identical to plain greedy; that is a different harness, drafter and runtime, not a controlled A/B against this table (see the update). The Qwen2.5 pair was not re-measured.

Reading it honestly

Model coverage: top local LLMs, 4-bit, on one 4090

Popular local models compressed to 4-bit and run in the pure-Rust runtime. Pareto columns are HF-loop decode (fp16 vs 4-bit); the fused runtime is faster. Four beyond-Llama ports now run end-to-end: DeepSeek-V2-Lite (MLA + Mixture-of-Experts), Gemma-2 (GeGLU + logit softcapping + 4-norm layers), Phi-4 14B (fused qkv/gate_up projections, torch-free export), and DeepSeek-R1 671B, the largest open-source model, with its experts resident in host RAM and decoded on the CPU (the inversion). Browse the models that fit your GPU to see which run on the card you have.

ModelArchfp16 VRAM4-bit VRAMWikitext PPL (fp16 → 4-bit)J/token net (fp16 → 4-bit)Runtime
Qwen2.5-7B #1 localGQA15.23 GB5.49 GB7.24 → 8.224.66 → 2.45~101 tok/s (verified)
DeepSeek-R1-Distill-Qwen-7BGQA15.23 GB5.49 GB25.78 → 30.734.63 → 2.42runs
Llama-2-7BMHA13.48 GB3.81 GBn/a3.87 → 2.13194.5 tok/s, HF 16/16
Llama-3.1-8BGQA16.06 GB5.64 GB7.35 → 8.454.96 → 2.38173.5 tok/s
Mistral-7BGQA14.50 GB4.07 GB5.77 → 6.194.68 → 2.22runs
DeepSeek-V2-Lite (16B) MLA+MoEMLA + 64-expert MoEn/a~7 GBn/an/acoherent, ~10 tok/s
Gemma-2-9B GeGLU+softcapGeGLU + softcaps + 4-norm~18 GB~6 GBn/an/acoherent, 74 tok/s
Phi-4 (14B) torch-freePhi3, fused qkv/gate_upn/a~7.9 GBn/an/acoherent, 79.9 tok/s
DeepSeek-R1 671B inversionMLA q_lora + 256-expert MoE1.34 TB (bf16)350 GB, one noden/a+1.5% (mixed 8-bit)2.46 tok/s, full model, one box

Runtime column for Llama-2-7B and Llama-3.1-8B: November 2026 decode kernels (not yet in an app release), the new exporter's files, CUDA graph, 512-token generation on an RTX 4090. The Llama-3.1-8B figure is for the Instruct model, whose greedy continuation follows the Hugging Face reference for 4 of 16 tokens before splitting on a near-tie (top-two logit margin 0.004 at the fifth token). Earlier runtime figures, not comparable (different window or harness): Llama-2-7B 135 tok/s at short context in June 2026; Llama-3.1-8B 118 tok/s in July 2026. For these two rows, the VRAM, PPL and J/token columns are July 2026 HF-loop measurements of fp16 and of the previous 4-bit files, and were not re-measured.

Consistent across models: ~2.8x less VRAM, ~47% less energy per token at 4-bit. Gemma-2 ran correctly on its first execution (softcapping, GeGLU, the 4-norm residual and embedding scaling all baked right); DeepSeek-V2-Lite's MLA + MoE required a multi-bug debug pass. Llama-3.1-8B and Mistral-7B are measured above; Phi-4 (14B) was exported torch-free (safetensors read directly + GPU k-means, no PyTorch or transformers) and decodes coherently at 79.9 tok/s; its fp16 Pareto columns are pending. The 671B row is now a serving result, not a reach probe: the full DeepSeek-R1 decodes at 2.46 tok/s on a single commodity node (64 vCPU, one L40S, 497 GB RAM), a x10.2 gain over the 0.24 tok/s disk-offload baseline. The lever is inversion: routed experts stay resident in host RAM and are decoded in place on the CPU (one vpshufb per 32 weights on AVX2), while the GPU holds only MLA attention and the dense layers (20.9 GB VRAM). At batch one only the ~10 GB of routed-expert bytes per token matter, and they no longer travel to the GPU. The gain came from removing eight measured walls (0.24, 0.44, 0.96, 1.31, 1.35, 1.67, 1.84, 2.46 tok/s); routed decode reaches ~67 GB/s aggregate, near the memory read floor, so throughput scales with host memory bandwidth. Quality is honest, four measured wikitext-2 perplexities: full 8-bit is essentially lossless at +0.03% vs fp16 (5.70), 4-bit costs +7.1% (6.10), and a mixed-precision variant (shared experts + output head at 8-bit) captures most of the recovery at +1.5% (5.78) for +0.7% size. Watch the recorded 671B run replayed at its true speed, or read the raw per-token logs in the repo.

Does it run on your GPU?

Trapetum is CUDA, so it runs on any modern NVIDIA GPU. You recompile the kernel for your architecture with one -arch flag, or ship a multi-arch build that covers them all in a single binary. The 4-bit memory saving is universal; only the speedup depends on the card.

Your GPU Arch Memory win Decode speed
RTX 30 / 40 series (3060, 3090, 4070, 4090, ...) sm_86 / sm_89 ~3.5x less 2.20x, bandwidth-bound consumer card
A40 / L40 (A100 runs the same arch) sm_80 / sm_86 / sm_89 ~3.5x less largest gain (most bandwidth-limited), 2.34x
H100 / H200 sm_90 ~3.5x less parity (fp16 already near roofline)
Turing and older (GTX 16, RTX 20, V100) sm_70 / sm_75 ~3.5x less untested, rebuild needed
Apple Silicon (M1 to M4) Metal unified supported (see the Apple M4 section below)
AMD, CPU - - AMD (ROCm) in development; CPU not supported

The point: the memory win (a 7B in 3.66 GB, so it fits an 8 GB card) holds on every NVIDIA GPU. The speed win is largest exactly where most people run, on bandwidth-limited consumer cards. Validated on Ampere, Ada and Hopper (sm_80 to sm_90); pre-Ampere is untested.

Apple Silicon (Metal): the same win on a MacBook

The fused 4-bit decode kernel is ported to Metal and runs on the Apple GPU, no NVIDIA needed. Measured on an Apple M4 (10-core GPU), batch-1 decode, same methodology as the CUDA numbers. The bandwidth law holds: the fused decode reads a quarter of the bytes, so it beats a dense fp16 GEMV.

Fused 4-bit GEMV vs dense fp16 (M4, 200 iters/shape)

Shape (IC x OC)4-bitfp16 denseSpeedup
4096 x 40960.279 ms0.431 ms1.55x
4096 x 110080.959 ms1.588 ms1.66x
11008 x 40960.809 ms1.771 ms2.19x

End-to-end decode on the M4

Model (4-bit)Size on diskDecode throughput
Llama-3.2 1B1.2 GB~28.5 tok/s
Llama-2 7B3.66 GB~6.6 tok/s

Why the 1B is 1.2 GB and not smaller: Llama-3.2 uses a 128k-token vocabulary (4x Llama-2's 32k), and the embedding table is kept in fp16 rather than quantized, so it dominates a small model's size. That is expected, not a rounding error.

The 4-bit path reproduces the CPU reference numerically (fused GEMV to 3e-7, attention, RMSNorm, SwiGLU all validated). Same Rust runtime, same .cbk files as the CUDA build; only the GPU backend differs. Perf still has headroom (the prefill is not yet batched), so these are conservative.

Kernel microbenchmarks

Synthetic matrices (IC = OC = 4096, fp16 activations/output, batch M = 1), fixed seed mt19937(0). Timings use CUDA events over 50-500 iterations after warmup. All variants verified numerically against cuBLAS on the same random data.

Decode GEMV (A40, sm_86, CUDA 11.8)

Decode is memory-bound: reading fewer weight bytes directly buys latency. The lever is bits per index.

Kernel Scheme Latency (ms) vs cuBLAS Max rel. err
cuBLAS fp16 dense (baseline) dense fp16 0.0612 1.00x -
gemv_codebook.cu uint8, K = 256 0.0562 1.09x 2.9e-5
gemv_codebook.cu uint8, K = 64 0.0389 1.57x 2.4e-5
gemv_codebook_4bit.cu fastest 4-bit, K = 16 0.0261 2.34x 2.6e-5

Cross-GPU decode (4-bit, K = 16, 4096x4096, batch 1)

The bandwidth law: the more bandwidth-limited the GPU, the larger the speedup from reading 1/4 the weight bytes. On a bandwidth-rich H100 the fp16 GEMV is already near roofline, so the kernel only ties.

GPU Bandwidth class Speedup vs cuBLAS fp16 Rel. err
RTX 4090 (sm_89) ~1.0 TB/s 2.20x 3e-4
A40 (sm_86) ~0.7 TB/s 2.34x 3e-4
H100 PCIe (sm_90) ~3.3 TB/s 0.99x (parity) 3e-4

Per-layer PyTorch op speedup vs torch dense fp16 (RTX 4090)

Speedup of the custom codebook_gemv op versus torch.mv (the batch-1 path a real fp16 Linear runs), by layer shape. Tiny matrices are overhead-bound and lose; the large MLP layers that dominate LLM parameter count win by 3.7x.

Layer (IC x OC) Speedup vs torch dense fp16
3072 x 768 (small) 0.58x (loses, overhead-bound)
4096 x 4096 2.10x
11008 x 4096 (MLP down) 3.70x
4096 x 11008 (MLP up) 3.77x

Apply the kernel selectively to large layers. Reproduce with python kernels/shapes_test.py.

Dequantization bandwidth (A40)

Kernel Effective bandwidth
naive (redundant shared-memory staging) 31.8 GB/s
dequant_l2 (L2-cached gather) 213 GB/s (6.7x)

Additive vector-quantization GEMV (AQLM-style, 2-bit)

A fused decode kernel for additive (AQLM-style) codebooks: builds a per-group LUT in shared memory, then reads only the codes. At 2 bits (M=2, K=256, D=8) that is 4.2 MB of codes versus 33.6 MB of fp16 weight. Measured vs live cuBLAS fp16 GEMV at 4096x4096, rel. err 2.3e-4.

GPU 2-bit (M=2) vs cuBLAS 4-bit (M=4) vs cuBLAS
RTX 4090 (~1.0 TB/s) 1.71x -
A40 (~0.7 TB/s) 4.30x 2.39x

The key optimization for v3: vectorized 32-bit code reads (four 8-bit codes per thread), which lifted the A40 from 2.35x to 4.30x. The kernel decodes real AQLM weights (Llama-2-7b-AQLM-2Bit-2x8-hf format maps one-to-one onto this kernel), and the runtime now serves 2-bit AQLM checkpoints end to end (see the 2-bit table under model-level results).

Model-level results: Llama-2-7B end-to-end (RTX 4090)

All 224 projection layers quantized to Trapetum 4-bit (K = 16, per-column k-means). Evaluated end to end on an RTX 4090.

Integration path Decode tok/s VRAM (GB) vs fp16
fp16 baseline (cuBLAS) 61.6 13.58 1.00x
naive per-layer custom-op swap 44 (est.) 4.73 0.73x
cast-free eager (no per-op float32 cast) ~52 (est.) 4.73 ~0.85x
CUDA-graphed decode path (clean integration) ships 123.4 4.73 2.0x

The progression 0.73x (naive) to 0.85x (cast-free) to 2.0x (CUDA-graphed) shows that the kernel win is real: it takes proper CUDA-graph integration to see it end to end. Memory drops 2.9x (13.58 GB to 4.73 GB). Note: 123.4 tok/s here is the CUDA-graphed decode path (vs fp16 61.6); it is distinct from the pure-Rust runtime, which reached 135 tok/s at short context and 81 tok/s at a 512-token context in June 2026 (the November 2026 kernels, not yet in an app release, reach 194.5 tok/s with CUDA graphs over a 512-token generation, see the update).

Llama-2-13B on A40 (bandwidth law at scale)

Model/method Decode tok/s VRAM (GB) Speedup Memory ratio
Llama-2-13B fp16 (A40) 20.0 26.17 1.00x -
Llama-2-13B Trapetum 4-bit (A40) 49.0 8.50 2.45x 3.08x less

Accuracy: 4-bit accuracy attempts on Llama-2-7B

Scheme (4-bit, Llama-2 7B) Wikitext-2 PPL
fp16 baseline 5.83
Trapetum scalar, research script (unweighted k-means, LM head in fp16) 6.34
+ activation-aware calibration 6.17
+ per-channel scale search 6.18
+ full AWQ pipeline (scale + clip) 6.21
+ incoherence processing (QuIP# lever) 6.29
beam-search + least-squares additive VQ (this work) 6.13
Runtime file, legacy exporter (4-bit LM head), read back from disk 6.79
Runtime file, new exporter (E[x2]-weighted k-means, 8-bit LM head) new 6.10

PPL protocol: custom codebook eval harness (ablations), the first 30 windows of 2,048 tokens of the wikitext-2 test set; fp16 baseline 5.83 on this harness. Numbers are not directly comparable to the HF-forward seqlen-2048 harness (fp16 5.47) or the RTX 4090 energy run (fp16 5.28), which use different protocols. The research-script rows keep the LM head in fp16; the two runtime-file rows (November 2026) are read back from the .cbk files, LM head included. Correction: the research script's 6.34 was previously presented as the runtime file's perplexity.

The calibration and VQ attempts listed above did not close the accuracy gap. One later recipe narrowed it: centroid updates weighted by E[x2] of each input channel, which the runtime's exporter now uses (6.10 here, 5.72 on the full test set). That is the best 4-bit result in this table, ahead of the trained additive VQ (6.13 PPL), which also decodes at 2.39x (A40). The gap to fp16 of the runtime's 4-bit file falls from 0.96 to 0.27 PPL, and the scalar codebook still trails AWQ on the full test set (5.72 against 5.60 in the H100 table below, fp16 5.47 in both). The value of this scheme is still mainly memory and kernel speed.

2-bit AQLM served end to end (CBK4), Llama-2-7B on RTX 4090, 5 November 2026

The runtime now serves real 2-bit AQLM checkpoints. A CBK4 file keeps the AQLM codes, codebooks and scales byte for byte, the dense Llama blocks dispatch each projection to the scalar or the additive kernel, and the LM head is stored as an 8-bit scalar codebook. Shown: the PV-tuned ISTA-DASLab/Llama-2-7b-AQLM-PV-2Bit-2x8-hf, which reproduces the Hugging Face greedy continuation 16/16.

Llama-2 7B in the runtime File (GB) Peak VRAM (MiB) PPL (custom harness) Decode tok/s, eager Decode tok/s, CUDA graph CUDA graph, 256-token window
Trapetum scalar 4-bit (CBK3), legacy exporter corrected 3.66 4,166 6.79 143 157 111
AQLM PV 2x8, 2-bit (CBK4) new 2.05 2,730 6.67 182 205 132

RTX 4090, batch 1, decode at positions 6 to 22 unless noted, with the runtime published on 5 November 2026 (before the kernels of the November 2026 update). Correction: this table showed PPL 6.34 for the CBK3 file, a research-script figure; read back from the file, it scores 6.79 (6.35 on the full test set). The statement that the 2-bit file decodes faster "for +0.33 PPL" is withdrawn: against this CBK3 file the 2-bit file was slightly better in quality (6.665 against 6.791 unrounded, a 0.13 gap) while 44% smaller and 1.2x to 1.3x faster with these kernels. The runtime's new exporter writes a 3.75 GB 4-bit file that scores 6.10; against it the 2-bit file is 0.57 PPL behind, for 45% less disk and 36% less peak GPU memory (2,736 against 4,246 MiB with CUDA graphs). The 4-bit path has since gained fused projection kernels and a tensor-core GEMV, which the 2-bit path has not (it gets only the RoPE and KV-append fusion), and the two have not been compared on one host since, so no 2-bit speed advantage is claimed for the new kernels. The 6.67 was measured on the PV checkpoint (the exported file reads back at 6.665); the same checkpoint scores 6.26 on the full wikitext-2 test set (fp16 5.47); the non-PV 2x8 checkpoint we first loaded scores 7.63 (7.23). CUDA-graph decode (TRAPETUM_GRAPH=1) replays the whole step with one launch and saved about 0.6 ms per token in both formats in this measurement. Limits: four PV checkpoints exist in the 2x8 format (since 7 November 2026 the runtime also serves the 1x16 format, see the update), the 2-bit path is CUDA only, and its 2x8 decode sums partials with atomics, so logits vary in the last digits from run to run.

Multi-method comparison (H100, from the paper)

A fair single-harness benchmark of fp16, AWQ, and AQLM on Llama-2 7B and 70B. Full wikitext-2 PPL, HuggingFace forward path, greedy, batch 1, seed 0. The fp16 7B PPL of 5.47 and AWQ 70B PPL of 3.41 match published values, which validates the harness. VRAM in this table is H100 peak VRAM (the fp16 7B figure of 15.3 GB is the H100 peak, higher than the ~13.5 GB fp16 weight footprint on the 4090 pages).

Model Method Decode (tok/s) VRAM (GB) Wikitext-2 PPL
Llama-2 7B fp16 43.7 15.3 5.47
Llama-2 7B AWQ 4-bit 26.8 5.7 5.60
Llama-2 7B AQLM 2-bit 25.0 4.2 6.34
Llama-2 70B fp16 ~138 GB, does not fit on one 80 GB GPU
Llama-2 70B AWQ 4-bit 9.1 38.4 3.41
Llama-2 70B AQLM 2-bit 9.2 20.7 4.06

PPL protocol: H100 HuggingFace forward path, seqlen 2048, median of 3 (matches published Llama-2). The AQLM 2-bit rows use the 1x16 checkpoint (ISTA-DASLab .1x16-hf: 7B PPL 6.34, 70B PPL 4.06), a different checkpoint from the 2x8 AQLM 2-bit (PPL 7.01) in the RTX 4090 Pareto table above.

At 7B on H100, every quantized method decodes slower than fp16: bandwidth is high enough that fp16 is already fast and the dequant overhead dominates. At 70B, the story flips: AQLM 2-bit fits in 20.7 GB on a single 24 GB consumer GPU, and at PPL 4.06 it is markedly more accurate than fp16 7B (PPL 5.47) at comparable memory.

Methodology and hardware

Hardware

RTX 4090 (RunPod), CUDA 11.8/12.x (as reported per run), sm_89 for 40-series, sm_86 for A40, sm_90 for H100 PCIe. All kernel microbenchmarks at matrix shape IC = OC = 4096.

Randomness and reproducibility

All kernel randomness is generated by std::mt19937(0) (fixed seed 0). Timings use cudaEvent over 50-500 iterations after warmup. Every kernel is checked for numerical correctness against cuBLAS on the same random data; the maximum relative error is reported.

Energy measurement

pynvml power.draw sampled over the full 512-token generation window. J/token = (mean power W) x (time per token s). The Rust runtime row: 209 W mean, 2.58 J/token. gCO2/1k-tok is derived from J/token x grid intensity (not measured).

Wikitext-2 PPL

Full 30 000 token evaluation, context window 2048, stride 512. Generation-independent (does not depend on decode length). HuggingFace forward path, greedy, batch 1, seed 0.

The file perplexities of the November 2026 update use two other protocols: the repository's harness, which averages the first 30 windows of 2,048 tokens of the wikitext-2 test set (fp16 Llama-2-7B 5.83), and the full-test-set harness, which covers all 166 full windows of 2,048 tokens of the 341,469-token test set (fp16 5.47), both with the Llama 2 tokenizer; the Llama-3.1 rows use the Llama 3 tokenizer (141 full windows of a 289,077-token test set, no fp16 reference).

Reproducing the numbers

All code is public, and so are the raw JSON results and figures of the earlier runs; the November 2026 results are summarized in the update above. The full table regenerates itself on your GPU:

# one kernel (adjust -arch: sm_80/A100, sm_86/A40, sm_89/RTX40, sm_90/H100)
nvcc -O3 -arch=sm_90 kernels/gemv_codebook_4bit.cu -o gemv4 && ./gemv4

# full kernel reference table (writes results.json + results.md)
python kernels/benchmark.py --arch sm_90

# end-to-end Pareto (Python rows)
python bench/pareto.py --gen 512

# Rust runtime row (requires the .cbk quantized model)
generate <model.cbk> ... 512

# per-layer PyTorch op shapes
python kernels/shapes_test.py

# additive VQ kernel (2-bit, K=16, A40)
nvcc -O3 -arch=sm_86 -DGT=16 kernels/avq_gemv3.cu -lcublas -o avq3 && ./avq3

Honest limits

Convinced by the numbers?

Every figure above, except those of the November 2026 update, reproduces with one command. The next step takes one line.

Run it on your machine Star on GitHub Read the paper

All figures were measured on the specific hardware, drivers and configuration described above and are provided for information only. They are not a performance guarantee; your results will vary with hardware, models and settings. Pre-release software, provided as is without warranty. Terms.