v100-skinny README

The project's README from github.com/dnv2003/v100-skinny. As one markdown file: README.md.

v100-skinny

Blackwell gets NVFP4 support in silicon. Volta gets it from software.

v100-skinny runs Qwen3.8-27B's published mixed NVFP4/FP8 weight representation unchanged on four Tesla V100-SXM2-16GB. In the release's main decode and k≤7 verification regime, NVFP4 regions run through QPN2 (W4A16) and FP8 regions through QPN8 (W8A16), both hand-written around Volta's mma.sync.m8n8k4. No FP8→FP4 requantization. No persistent FP16 weight copy. FP16 activations and FP16 KV cache. OpenAI-compatible serving — a usable Qwen3.8 endpoint, not a kernel demo.

The GPUs cost about A$600 used. (That is the accelerators only — not the server, not the rest of the build.)

Quick start · Results · How QPN works · Reproduction · Limitations

Result

AIME 2026 Problem 1, five seeds, both engines run in our lab through one harness, each at its own best measured native-MTP depth. The RTX 5090 was running NInfer, a specialist engine built for maximum performance on this exact model family rather than a weak generic baseline. Sampling is NInfer's published profile applied identically to both: temperature 0.6, top-p 0.95, top-k 20, presence penalty 1.0, thinking on, reasoning_effort medium, seeds 1001–5005.

4× V100-SXM2-16GB (2017) 1× RTX 5090 (2025)
Engine v100-skinny on 1Cat-vLLM, k=7 NInfer, --draft-tokens 5
Decode 219.1 ± 5.9 tok/s 214.7 ± 9.2 tok/s
Tokens committed / round 5.89 4.27
Round latency 26.9 ms 19.9 ms
Time to correct answer 6.90 ± 0.30 s 6.56 ± 1.34 s
Completion tokens 1,513 ± 44 1,403 ± 253
Correct answer (277) 5/5 5/5
Published quantization RadixArk NInfer artifact, Unsloth-derived

This is a same-lab system comparison on the same Qwen3.8 base model, not a same-weight causal engine A/B: the two engines use different published NVFP4 artifacts. NInfer depth 5 is its maximum exposed depth and its best measured setting on this fixture; v100-skinny's best measured setting is k=7.

This is parity — 1.02×, intervals overlapping. We do not claim a win, and the result is not universal: their prefill is roughly 4× ours. The decode result is an exact cancellation — our round is 35% longer (26.9 vs 19.9 ms) and we commit 38% more tokens per round (5.89 vs 4.27); 1.38 / 1.35 = 1.02. QPN makes the wider k=7 verification batch cheap enough for the model's native MTP depth to pay.

Why depth is the lever. mma.sync.m8n8k4 issues an 8-row tile whether or not all 8 rows carry work, and QPN2 holds 71% of read roofline at M=8 — the k=7 verification width. On the measured stack, verification costs +0.383 ms per extra draft row against +0.817 ms per sequential drafter step. The 5090 completes a round much faster; Volta compensates by extracting more useful tokens from each round.

What v1.1 changes

The published mixed weight allocation is now directly executable on Volta. ModelOpt's native mixed FP8/NVFP4 path was gated above Volta — the FP8 regions had no SM70 path at all. v1.0 created the first practical bridge: it down-converted those protected FP8 layers to FP4 and served the resulting derivative through the new W4A16 stack. v1.1 adds QPN8, lowers the mixed-path gate to SM70, and retires the conversion that v1.0 had made necessary. The boot gate verifies the QPN2/QPN8 route census and zero reference fallbacks on every start. Across four ranks, all 512/512 protected FP8 module instances are eligible for QPN8, the checkpoint's own lm_head routes to QPN2, and reference fallback calls remain zero.

Also in v1.1:

End-to-end benchmarks

Serving configuration. TP4, GMU 0.88, FP16 KV, decode partition pinned, MT=2 on, MTP with greedy drafting and local-argmax reduction, --max-num-seqs 1 (single stream). NUMA-pinned to socket 0, application clocks pinned.

The domain table below was run at --max-model-len 32768; the AIME rows were run at 98,304. Before the decode-partition fix that difference changed short- context round latency. On the shipping path, declared MML is timing-neutral; the two values are retained here as measurement provenance, not as performance tuning. The same AIME configuration reads 206.8 tok/s without the pin and 219.1 tok/s with it.

Sampling — the domain table. Fully greedy and fully specified; no field is left for the server to fill from a model card:

temperature 0.0 · top_p 1.0 · top_k disabled
presence_penalty 0.0 · frequency_penalty 0.0 · seed 1001
thinking off (except the math cell, which is thinking-on)

Sampling — the AIME rows. ninfer's published profile, applied identically to both engines:

temperature 0.6 · top_p 0.95 · top_k 20
presence_penalty 1.0 · frequency_penalty 0.0
thinking on · reasoning_effort medium
five fixed seeds: 1001, 2002, 3003, 4004, 5005

reasoning_effort is set explicitly rather than defaulted: the chat template defaults to xhigh, which injects a system message and costs 4.0–11.0% at matched output length, while medium injects none and matches the prompt ninfer's converter produces.

Metric. Throughput is (generated − (τ+1)) / decode_seconds — we subtract a whole speculative round. The convention that subtracts a single token reads higher by exactly τ/generated: 0.4–0.7% on the cells below, more at k=15 where τ is larger. Every figure here uses ours, the more conservative one.

Cell k=3 τ k=7 τ k=15 τ best
doc2json (invoice → JSON) 195.1 3.00 309.5 6.83 392.1 13.88 k=15
log2json (logs → records) 195.6 3.00 307.1 6.81 391.4 13.94 k=15
json (generate records) 184.9 2.77 268.6 5.73 270.3 8.93 k=15 (tie)
mergesort 175.3 2.56 226.9 4.71 197.7 6.27 k=7
code (red-black tree) 173.1 2.52 197.1 3.93 155.0 4.67 k=7
math (thinking on) 161.6 2.30 187.0 3.75 147.0 4.46 k=7
prose 113.9 1.32 101.7 1.55 68.5 1.49 k=3

Round latency is nearly constant within each depth — 20.3–20.5 ms at k=3, 25.0–25.4 at k=7, 36.4–38.1 at k=15 — so cells differ almost entirely in how many drafted tokens survive verification. (The k=15 band is the loosest: the two extraction cells sit ~1.5 ms above prose there, which is the M=16 NVFP4 kernel gap below showing up end-to-end.)

The ~392 tok/s figures are workload-specific, not headline throughput. Extraction is the best case for speculation: at k=3 both extraction cells pin at τ = 3.00 — every draft accepted, the ceiling of that depth — and depth is what breaks the ceiling. Free-form prose runs at 113.9 and prefers k=3. The optimal depth is a property of the workload, not of the engine, which is why every column names its depth. Depth is a boot-time setting in v1.1.

Recommended release profiles (concrete commands in Quick start): k=7 for general short/medium-context work, k=3 for long-context serving (65k measured), and k=15 only for highly predictable structured/extraction workloads where its wider verification actually pays.

Declared context is free

--max-model-len used to cost real time: the decode partition geometry is derived from it, so declaring a large window taxed every speculative round even at short live context. With the pin, that is gone across the whole supported range — eleven boots, 4,096 to 262,144, a 64× range of declared context, ms/round:

MML mergesort json log2json
4,096 25.14 25.18 25.40
8,192 25.10 25.17 25.42
16,384 25.06 25.10 25.36
32,768 25.13 25.23 25.44
65,536 25.05 25.10 25.36
98,304 25.13 25.17 25.47
131,072 25.20 25.22 25.51
163,840 25.14 25.19 25.48
196,608 25.17 25.14 25.51
244,608 25.08 25.13 25.53
262,144 25.15 25.19 25.61

Flat to within 0.25 ms, with τ identical in every arm (4.71 / 5.73 / 6.81), so all eleven did byte-identical work. The only hint of a trend is log2json, the longest cell, drifting ~0.2 ms across the full 64× range — under 1%, and absent from the other two. Unpinned at the same 32,768 the figures are 25.84 / 26.50 / 28.03 — the pin recovers up to 2.59 ms (−9.2% on log2json), and the penalty it removes grows with live context.

There is no "fastest" --max-model-len. Declare the window you need within the KV capacity reported at boot; 244,608 is reliable on both measured memory profiles, while 262,144 is marginal (results/mml_speed_20260819.csv).

Declaring a large window is free; filling it is not. Every arm above generates at short live context — what is flat is the cost of having declared a large window, not the cost of having filled one. Decode does slow as the cache fills: at k=7, ~127 tok/s at short context against ~55 at 65k. Ordinary conversation stays in the fast regime, and long documents want the k=3 profile (Depth profiles), which is ~40% faster than k=7 at 65k. The measured curve is in Limitations.

Seconds to answer, AIME f01, decode only, n=5 — published including where it goes against us:

seconds completion tokens
ours, k=7 6.90 ± 0.30 1,513 ± 44
NInfer, draft 5 6.56 ± 1.34 1,403 ± 253
NInfer, draft 3 (their published default) 7.73 ± 1.40 1,464 ± 238

Their point estimate is 4.9% ahead at their best depth; we are 10.7% ahead of their published configuration; our latency variance is 4.5× tighter. This excludes prefill, which is theirs by ~4×, so a full end-to-end figure would favour them by more than 5%.

Quick start

Requires 4× SM70 GPUs, CUDA 12.8+, and the 1Cat-vLLM 1.2.2 wheel. The bootstrap pins tilelang==0.1.10 and apache-tvm-ffi==0.1.10 together. Earlier tilelang does not build on Volta and fails later inside GDN attention where it looks like a kernel bug; and tilelang does not pin its own apache-tvm-ffi, so a machine built today otherwise resolves a newer one that aborts on import (tvm::ffi::Error: TypeAttr __ffi_repr__ is already registered).

pip will print a red dependency-conflict block during this step, because the 1Cat-vLLM 1.2.2 wheel declares the 0.1.9 versions of both. That is expected and the bootstrap says so — the install is correct, and the bootstrap now fails loudly if import tilelang does not work afterwards.

# Reproducible path: no overrides. The wheel URL and its SHA256 are pinned in
# the script, and the digest is verified before installation.
bash scripts/bootstrap-sm70.sh

# The checkpoint is pinned to an immutable revision. Any other revision is a
# different set of weights and the published numbers do not describe it.
hf download RadixArk/Qwen3.8-27B-NVFP4 \
  --revision 554ebba9b5f1b79dc11246341960360e6ef05ef4 \
  --local-dir ./Qwen3.8-27B-NVFP4

bash scripts/serve-qwen38-native.sh ./Qwen3.8-27B-NVFP4

Supplying your own VLLM_WHEEL requires VLLM_WHEEL_SHA256 with it — the bootstrap fails closed on an unverified wheel rather than installing it. The serve script accepts any directory containing a config.json, which is a convenience for local experiments; it is not the reproducible path, and results from an unpinned checkpoint are not comparable to the ones here.

Depth profiles

Speculation depth is a boot-time setting in v1.1 (per-request selection is v1.2 work). K and MML are environment overrides the launcher already honours, so each profile below is a real command rather than a suggestion — and the boot gate verifies that the depth you asked for is the depth the engine served, so a profile that fails to take exits non-zero instead of quietly serving k=7:

CKPT=/path/to/Qwen3.8-27B-NVFP4

# general, short/medium context — the default, and what the benchmarks use
bash scripts/serve-qwen38-native.sh $CKPT

# long context — k=3 measures 76.3 tok/s at ~65k against k=7's 54.7 (+39%),
# and beats speculation-off (65.5) too. Declaring the larger window is free.
K=3 MML=196608 bash scripts/serve-qwen38-native.sh $CKPT

# structured extraction, short context only — k=15 wins on high-acceptance
# cells (doc2json 392.1 tok/s) and loses badly on prose. Do not use it as a
# general default, and do not use it at long context.
K=15 bash scripts/serve-qwen38-native.sh $CKPT

Pick k by context length, not by taste. The drafter runs k sequential steps per round and each one reads the whole KV cache, so a step costs ~0.817 ms at short context but ~3.65 ms at 65k — while the tokens it wins barely improve (1.55 accepted/round at k=3 against 1.63 at k=7). Deep speculation stops paying as context grows, which is why the long-context profile is shallower rather than off.

bootstrap-sm70.sh installs the pinned engine, deploys the fork patches with backups, and builds the kernels. It refuses a wheel that is not 1Cat-vLLM 1.2.2 — every published number here was measured on it — and verifies VLLM_WHEEL_SHA256. Overriding the wheel without a digest is refused.

serve-qwen38-native.sh starts an OpenAI-compatible server on port 8000, bound to 127.0.0.1. The server has no authentication, so reaching it from another machine should be an SSH tunnel:

ssh -N -L 8000:127.0.0.1:8000 you@your-box

Binding a public interface requires saying so explicitly (HOST=0.0.0.0 I_UNDERSTAND_THIS_IS_UNAUTHENTICATED=1 …), and even then it belongs behind a firewall — vLLM's own security guidance is that its API keys do not protect every endpoint. It aborts before boot if the GPUs are occupied, then refuses to report success unless the release gates pass: served depth equals requested k; lm_head is served from the checkpoint's own codes; no lm_head repack fallback occurs; the resolved KV storage is FP16; the fast XQA/tensor-core decode-attention route is observed with zero scalar_paged calls; and the expected QPN2/QPN8 decode routes are present. A mis-booted server exits non-zero rather than quietly producing unquotable numbers.

How QPN works

The naive approach to low-bit weights on old hardware is: dequantize a tile to FP16 in memory, then call a GEMM. That spends HBM bandwidth on the format you were trying to avoid.

The QPN decode and verification paths do not do that. Packed low-bit weights stay compressed through HBM and are decoded only at the point of consumption, directly into the FP16 register operands mma.sync.m8n8k4 requires. On those paths, the weight stream exists in FP16 only in registers for the instant it is consumed. Large-M FP8 prefill above the packed-kernel range is a separate, transient reconstruction path described below.

The mapping is the hard part. m8n8k4 is the one FP16 tensor-core op Volta exposes and it is notoriously difficult to feed. QPN puts all four quadpairs on the N dimension, sharing one activation tile: a single warp instruction issues four independent 8×8×4 MMAs against the same 8×4 activation fragment, so activation traffic per weight byte drops 4×. Weights are pre-permuted at load time into fragment order — nibble-interleaved so the decoder's output register pair is the B-fragment the instruction wants, with no shuffle in the inner loop. No shared memory in the main loop, one barrier total, 56 registers, zero spills (docs/qpn_race_notes.md).

QPN2 is the NVFP4 (e2m1 + FP8 group-16 scale) instantiation. QPN8 generalizes the same execution architecture to FP8 e4m3 — one contiguous byte stream, a simpler decoder, and the same quadpair-on-N geometry. MT=2 issues two 8-row tiles against a single weight pass, which is what makes the FP8 M 9–16 band viable for k=15. It exists for QPN8 only: there is no gemm_qpn2_mt2, so NVFP4 at M 9–16 falls back to the first-generation gemm_qpn.

One lesson worth recording: the first correct-looking version overflowed FP16 on real activation outliers. Only real-weight, real-activation tests exposed it. The fix folds the global scale into the group scales in-kernel with a per-16-element FP32 flush.

Native mixed-precision checkpoint path

The checkpoint is served as published, which requires the loader to do three things it previously would not:

  1. Dispatch per region, and the two regions have different reach.

    M NVFP4 regions (MLP trunk) FP8 regions (attention projections)
    1–8 qpn2 qpn8
    9–16 qpn qpn8-mt2 (two 8-row tiles, one weight pass)
    17–96 Marlin qpn8-chunked
    > 96 Marlin qpn8-prefill-reconstruct

    FP8 stays on packed QPN8 kernels through M=96; above that it remains under v100-skinny dispatch but uses transient reconstruction. NVFP4 hands off at M=17. The asymmetry is deliberate on one side and a consequence on the other. For FP8, chunking was measured to beat transient reconstruction all the way to M≈112 — by 3.1× on all three protected shapes — so the boundary sits at 96. Above it, weights are reconstructed transiently from the packed codes and never persisted.

    For NVFP4 the WMMA path that nominally covers 17–64 requires the checkpoint-native code stash, and production drops that stash (VLLM_SKINNY_DROP_CT=1) once the QPN prepack exists, to reclaim its memory. So M ≥ 17 falls to Marlin — a memory decision with a routing consequence, not a designed handoff at 17. Decode and speculative verify never enter that band.

    Every route is logged and counted, and the census is checked for fallbacks after every boot. The FP8 regions are the attention projections of every layer, two per layer, 128 modules per rank: 48 GDN layers contributing linear_attn.in_proj_qkvz and linear_attn.out_proj, and 16 full-attention layers contributing self_attn.qkv_proj and `self_att

v100-skinny (dnv2003/v100-skinny)

model-serving
Searching Database…

174 ·GitHub
V100-skinny
What is this?v100-skinny provides hand-written NVFP4 (W4A16) and FP8 (W8A16) CUDA kernels that let four Tesla V100-SXM2-16GB GPUs run Qwen3.8-27B's published mixed NVFP4/FP8 weights directly, serving the model through an OpenAI-compatible API.
Is this a general CUDA library or something more specific?It is a specific inference stack: two custom kernels (QPN2, QPN8) hand-written around Volta's mma.sync.m8n8k4 instruction, plus a serving layer on top of a patched 1Cat-vLLM — not a general-purpose quantization library.
How do I get it serving a model?Install CUDA 12.8+, then run the project's bootstrap script to install the pinned 1Cat-vLLM 1.2.2 wheel (with tilelang==0.1.10 and apache-tvm-ffi==0.1.10), then point the serve script at a checkpoint directory containing a config.json.
What exact software versions does it require?It needs the 1Cat-vLLM 1.2.2 wheel specifically; earlier tilelang versions fail to build on Volta, and an unpinned apache-tvm-ffi resolves to a version that crashes on import.
What are the main gaps or weak spots?The NVFP4 kernel still falls back to a first-generation, less-optimized path at the M=9-16 batch band, decode throughput drops sharply at long context (around 219 tok/s short-context vs ~55 tok/s at 65k), and prefill is roughly 4x slower than a modern GPU baseline.
Do I need a specific GPU, or will any NVIDIA card work?You need exactly four SM70 (Volta) GPUs such as the Tesla V100-SXM2-16GB — the kernels are hand-tuned for Volta's mma.sync.m8n8k4 instruction and are not a general multi-architecture kernel set.
What model does v100-skinny actually run?It runs Qwen3.8-27B using the model's own published mixed NVFP4/FP8 quantized weights, served unmodified rather than re-quantized.
How does four old V100s compare to a modern GPU?In the authors' own AIME benchmark, four V100s via v100-skinny matched an RTX 5090 running the specialist NInfer engine at about 219 vs 215 tokens/sec decode — parity, not a win, since the V100's prefill stage is roughly 4x slower.
Why bother with V100s instead of a newer card?The four V100 GPUs used in testing cost about A$600 total secondhand, so v100-skinny turns much cheaper retired hardware into a usable quantized-LLM serving endpoint.
What changed in the v1.1 release?v1.1 adds a new QPN8 FP8 kernel, extends the mixed FP8/NVFP4 path to SM70, switches the KV cache to FP16, makes speculative drafting greedy, and fixes a decode-partition bug worth up to 2.6ms per round.
What are similar projects?qwen3.8-Flash-DGX, dgx-spark-qwen38, HyperQwen, GLM-5.3-Flash-EXL3-2x-DGX-Sparks, TensorFold, Qwen3.8-Flash-Next-Dual-DGX-Sparks, TypeLLM, DeepSeek-v4-Flash-One-DGX-Spark
Which topics is it in?nvfp4, cuda-kernels, performance-optimization, gpu-programming, llm-serving, speculative-decoding, inference-engine, model-serving, llm-inference, vllm, qwen, quantization
Creator
dnv2003
Size
0.5 MB
Stars
174
Updated
CommitAug 19
Age
2mo
Commits · 90 Days