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.
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.
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 |
|---|---|---|---|---|---|
| GSM8K | 80 | 166.4 | 451.0 | 2.71x | 4.55 |
| HumanEval | 80 | 163.6 | 471.8 | 2.88x | 5.13 |
| MATH-500 | 20 | 165.4 | 449.1 | 2.72x | 4.62 |
| MT-Bench, first turns | 80 | 167.0 | 340.0 | 2.04x | 3.21 |
| All four sets, pooled (520 runs, 108,102 tokens) new | 260 | 165.5 | 415.6 | 2.51x | 4.19 |
| Pooled, drafter with an 8-bit head (one repetition) | 260 | 165.4 | 381.9 | 2.31x | 4.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, fp16 | 13.48 GB | 5.83 | 5.47 | - |
| 7B, 4-bit, new exporter | 3.75 GB | 6.10 | 5.72 | 205.0 |
| 70B, 2-bit AQLM 2x8 new | 17.96 GB | 4.68 | 4.49 | 38.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, 1x16 | 4.00 GB | 5.74 | 5.40 | 42.6 |
| 70B, 2x8 (above) | 17.96 GB | 4.68 | 4.49 | 38.4 |
| 70B PV, 1x16 new | 18.54 GB | 4.17 | 4.02 | 8.5 |
| 70B PV, 1x16, speculative (K = 7) new | + 2.05 GB drafter | 4.17 | 4.02 | 22.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=) 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 GB | 11.28 | 10.61 | 183.2 |
| 70B PV, 1x16 new | 21.00 GB | 6.62 | 6.28 | 8.7 |
| 70B PV, 1x16, speculative (K = 7) new | + 3.46 GB drafter | 6.62 | 6.28 | 21.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
- Not in a release yet. The code is on the repository's main branch (commits 2e8c170 to 6b8cb42), not in an app release. The desktop app's model registry is not updated: its downloadable models keep the old format and the old recipe, because installed apps (including the Windows beta 0.1.0-beta.1) cannot read the new files. A runtime release has to come before any new registry entries.
- CUDA only. Split-KV, the fused kernels, sparsity and the tensor-core GEMV are CUDA only for now. Apple Metal is unchanged: it keeps the sequential attention, runs the same separate operations as before and refuses sparsity. The new files have not been measured on Metal.
- Tensor cores. The tensor-core GEMV needs sm_80 or newer and was measured on an RTX 4090 only. Older GPUs,
TRAPETUM_TC=0, the non-default determinism modes and the sparse path keep the previous FMA kernels and their bits, and the 8-bit LM head stays on FMA kernels. - Sparsity is lossy and opt-in. +0.12 PPL at 25%, and the greedy output changes (3 of 16 tokens match the dense reference).
- Verify cost grows with context. At M=16: 1.21x a decode step at short context and 1.46x near 1,100 positions on Llama-2-7B; 1.34x to 1.36x and 1.58x near 1,050 to 1,170 positions on Llama-3.1-8B-Instruct; 1.81x at short context and 1.87x near 1,050 to 1,070 positions on Llama-2-70B 1x16; 1.85x at short context and 1.92x near 1,050 to 1,070 positions on Llama-3.1-70B-Instruct 1x16.
- PARD scope. Measured on these four prompt sets only, with a 4-bit target and a 4-bit drafter, in target-independent mode, at contexts under about 700 positions (prompts up to 437 tokens plus at most 256 new tokens). Exact means identical to plain greedy decoding of the same 4-bit target, not to the bf16 model.
- Weights not redistributed. The Llama-3.1 weights (8B-Instruct and 70B-Instruct; Llama 3.1 Community License), the Llama-2 checkpoints served here (70B, 13B and the 7B drafter; Llama 2 Community License) and the PARD-2 drafter are not redistributed.
- 2-bit path. It gets only the RoPE and KV-append fusion, since its additive-codebook projections are not fused (about 1.5% on an A40, 7.17 to 7.06 ms over the first 128 positions), and it has no tensor-core path. Its 2x8 decode sums partials with atomics, so its logits vary slightly from run to run, and the fused-kernel and tensor-core bit-identity results apply to the 4-bit path (the 1x16 decode and verify have their own bit checks, above). 2-bit and 4-bit have not been compared on one host since these changes.
- Llama-2-70B 2x8 scope. One run on one host, with a 2x8 checkpoint that is not PV-tuned. Its perplexities are computed in Python on the served weights, and the runtime's own NLL was compared with them on the first window only. The 2x8 path has no exact speculative decoding (the 1x16 path has one).
- Llama-2-70B 1x16 scope. One run per result on one host each, greedy decoding only. The speculative speed depends on how many drafts are accepted and on the prompts: at K = 7, 22.16 tok/s on the wikitext windows (4.68 tokens per verify pass, 65-token prompts whose prefill is timed) and 23.15 on the short prompts (4.13). The perplexities are computed in Python on the served weights; the runtime's NLL was compared on the first window only. Plain 1x16 decode failed its pre-registered speed criterion (at least 10 tok/s); the second run's speed criterion was for speculative decoding.
- Llama-3.1-70B scope. One run on one host, greedy decoding only, 16 prompts. The perplexities are computed in Python on the served weights; the runtime's NLL was compared on the first window only. The export leaves at most 1,473 MiB of the card free (highest 300 ms sample). The plain RoPE comes from the checkpoint's current configuration; this run does not test Meta's llama3 scaling with these weights.
- Energy not re-measured. The J/token figures on this page come from June 2026 (Pareto table, previous runtime) and July 2026 (model coverage, HF loop) runs.
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.
| 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.
Reading the table honestly
- Energy is the headline. The Rust runtime runs the same 4-bit projection weights (its LM head is quantized too) at 2.58 J/token, 2.1x less than fp16 (5.45) and ~1.5x less than the Python-wrapped kernel (4.00), iso-context.
- Speed gap, Python vs Rust. The 4-bit weights via a Python custom op decode at 28.4 tok/s, slower than fp16 (45.2). The bottleneck is Python per-op dispatch (224 calls/token), not the kernel. The same projection weights in Rust reached 81 tok/s in June 2026 (1.8x fp16 at this context); the November 2026 kernels, not yet in an app release, reach 194.5 tok/s with CUDA graphs over a 512-token generation on the new exporter's file, with energy not re-measured. The Python-to-Rust gap is the overhead cost made visible.
- Accuracy/memory curve. fp16 (PPL 5.28, 13.5 GB) to Trapetum 4-bit (PPL 5.92, 3.8 GB) to AQLM 2-bit (PPL 7.01, 2.1 GB). Nothing beats dense on accuracy.
- Mean GPU power (Rust row): 209 W, measured via pynvml over the 512-token window.
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.
| 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) | Vocab | Plain tok/s | K=1 | K=2 | K=3 | Verdict |
|---|---|---|---|---|---|---|
| Llama-2-7B + llama-160m | 32k | 124.6 | 1.19x | 1.32x | 1.24x | WIN, 164 tok/s |
| R1-Distill-Qwen-7B + Qwen2.5-1.5B | 152k | 100.6 | 0.76x | 0.76x | 0.76x | parity |
| Llama-3.1-8B + Llama-3.2-1B | 128k | 117.9 | 0.85x | 0.90x | 0.80x | parity |
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
- Lossless. Every emitted token is the target's own greedy choice; a rejected draft is corrected in the same step. The sequence equals plain decode exactly, verified for K = 1, 2, 3 and for adversarial drafters.
- Wall-clock: measured, lost, diagnosed, fixed, now a win. The real two-model harness first LOST 3x (0.34x) despite 0.91 acceptance: profiling traced it to the batched-verify kernel taking M at runtime, which spilled its accumulators to local memory (M=2 cost 5.5x an M=1 forward). Templating the kernel on M fixed it: verify M=2 now costs 1.11x of M=1, and spec-dec measures 1.32x wall-clock at K=2 (164 vs 125 tok/s), lossless, on the Llama-2-7B + 160m pair, against the 1.50x projected ceiling. In that harness, large-vocab pairs (Qwen2.5 152k, Llama-3.1 128k) sat at 0.9x parity: each drafter forward read back 500+ KB of logits for a host argmax. For Llama-3.1, PARD now measures 2.51x (November 2026 update); Qwen2.5 was not re-measured. Same integration law as the kernel story, now closed at the loop level too.
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.
| Model | Arch | fp16 VRAM | 4-bit VRAM | Wikitext PPL (fp16 → 4-bit) | J/token net (fp16 → 4-bit) | Runtime |
|---|---|---|---|---|---|---|
| Qwen2.5-7B #1 local | GQA | 15.23 GB | 5.49 GB | 7.24 → 8.22 | 4.66 → 2.45 | ~101 tok/s (verified) |
| DeepSeek-R1-Distill-Qwen-7B | GQA | 15.23 GB | 5.49 GB | 25.78 → 30.73 | 4.63 → 2.42 | runs |
| Llama-2-7B | MHA | 13.48 GB | 3.81 GB | n/a | 3.87 → 2.13 | 194.5 tok/s, HF 16/16 |
| Llama-3.1-8B | GQA | 16.06 GB | 5.64 GB | 7.35 → 8.45 | 4.96 → 2.38 | 173.5 tok/s |
| Mistral-7B | GQA | 14.50 GB | 4.07 GB | 5.77 → 6.19 | 4.68 → 2.22 | runs |
| DeepSeek-V2-Lite (16B) MLA+MoE | MLA + 64-expert MoE | n/a | ~7 GB | n/a | n/a | coherent, ~10 tok/s |
| Gemma-2-9B GeGLU+softcap | GeGLU + softcaps + 4-norm | ~18 GB | ~6 GB | n/a | n/a | coherent, 74 tok/s |
| Phi-4 (14B) torch-free | Phi3, fused qkv/gate_up | n/a | ~7.9 GB | n/a | n/a | coherent, 79.9 tok/s |
| DeepSeek-R1 671B inversion | MLA q_lora + 256-expert MoE | 1.34 TB (bf16) | 350 GB, one node | n/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-bit | fp16 dense | Speedup |
|---|---|---|---|
| 4096 x 4096 | 0.279 ms | 0.431 ms | 1.55x |
| 4096 x 11008 | 0.959 ms | 1.588 ms | 1.66x |
| 11008 x 4096 | 0.809 ms | 1.771 ms | 2.19x |
End-to-end decode on the M4
| Model (4-bit) | Size on disk | Decode throughput |
|---|---|---|
| Llama-3.2 1B | 1.2 GB | ~28.5 tok/s |
| Llama-2 7B | 3.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
- Accuracy ceiling. The Trapetum scalar does not beat fp16 on accuracy and never will. It is behind vector and trellis methods (AQLM, QuIP#, QTIP) at low bits. Three calibration and VQ attempts to close the gap are documented negative results. A later recipe, E[x2]-weighted k-means with an 8-bit LM head, cut the gap to fp16 of the runtime's 4-bit file from 0.96 to 0.27 PPL (30-window harness; 5.72 against fp16 5.47 on the full test set), still behind AWQ (5.60). The value of this scheme is mainly memory and kernel speed.
- Bandwidth law. The kernel speedup is largest on bandwidth-limited GPUs (RTX 4090 x2.20, A40 x2.34). On a bandwidth-rich H100, fp16 GEMV is already near roofline and the kernel only ties (x0.99). The win lives on the hardware most users actually run.
- Batch-1 decode only. Batched throughput needs a GEMM path. (Since November 2026 a tensor-core GEMV verifies a speculative draft of 2, 4, 8 or 16 tokens at 1.05x to 1.21x the cost of one decode step at short context; that serves speculative decoding, not batched serving.) The Tensor-Core prefill experiments (prof10/11.cu) reached ~0.21x cuBLAS. Batched throughput is an open problem, not a quick reuse of the decode kernel.
- Synthetic data for microbenchmarks. Codebooks and indices are random (mt19937(0)). Real weight and index distributions may differ.
- Single shape. Kernel microbenchmarks at IC = OC = 4096. Numbers will differ at other shapes; re-run benchmark.py.
- Partial comparison against kernel peers. The headline baseline is cuBLAS fp16. On AQLM the accuracy is matched by construction (the additive kernel decodes a real Llama-2-7b-AQLM-2Bit-2x8 checkpoint), but the op-level timing against AQLM's own kernel is unfinished, and no matched-accuracy speed claim is made against Marlin, FLUTE or VQ-LLM. See the baselines section of the paper.
- Prefill stays at ~0.21x cuBLAS. Trapetum buys memory in the compute-bound regime, not speed. Closing the remaining gap requires production-Marlin engineering (ldmatrix + shared-memory swizzling).
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