Benchmarks

End-to-end Pareto: speed, memory, accuracy, energy

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 weights as the Python row, with no Python at runtime. You can test these numbers yourself in the live compare demo, 4-bit versus fp16 on two identical RTX 4090s.

* 4.73 GB measured peak (3.5 GB weights + KV cache + activations). The short-context Rust throughput reaches 135 tok/s.

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 weights 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 weights in Rust reach 81 tok/s (1.8x fp16 at this context). 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.

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

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

The 32k-vocab pair wins as projected (1.32x measured vs the 1.50x ceiling). The large-vocab pairs sit at parity because every drafter forward currently reads back the full logits (500+ KB) for a host argmax; a device-side argmax is the next lever.

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. Large-vocab pairs (Qwen2.5 152k, Llama-3.1 128k) sit at 0.9x parity: each drafter forward reads back 500+ KB of logits for a host argmax, the next lever (device-side argmax). 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.

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.

The point: the memory win (a 7B in ~3.5 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)

End-to-end decode on the M4

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.

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.

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.

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

Dequantization bandwidth (A40)

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.

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).

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.

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 shipping pure-Rust runtime, which reaches 135 tok/s at short context and 81 tok/s at a 512-token context.

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

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

PPL protocol: custom codebook eval harness (ablations); 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 calibration and VQ attempts to close the accuracy gap are documented negative results. The best accuracy at 4-bit is the trained additive VQ (6.13 PPL), which also decodes at 2.39x (A40). The value of this scheme is memory and kernel speed, not accuracy.

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).

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.

Reproducing the numbers

All code, raw JSON results, and figures are public. 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. The value of this scheme is 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. 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).