qwentin

Qwen3.6-27B at 256k context on a single RTX 5090 — 119-131 tok/s at short context, ~119 at 128k and ~110 at 200k, vLLM-class throughput at higher quality, with cold prefill 2–4× faster by default (2026-07, measured on the 5090 itself).

qwentin is a from-scratch CUDA inference engine that runs the 27B hybrid-attention Qwen3.6
tower on one 32 GB consumer GPU (Blackwell / SM120), using hand-written tensor-core kernels
(FP6 weights, 4-bit KV cache) + MTP speculative decoding. It exposes an OpenAI-compatible API,
so any client — including an opencode/continue-style coding agent — can use it.

27B params + 256k context            →  one 32 GB RTX 5090   (steady-state ~30.8 GiB)
single-stream decode                 →  119-131 tok/s @ short · ~119 @ 32k-128k · ~110 @ 200k · ~98 @ 245k
cold prefill (wide+MMA, default ON)  →  2-4× faster at any length · 1960 tok/s @ 32k · 1245 @ 128k
quality (tf-top1 vs bf16)            →  91.3  (the highest that still fits 256k on 32 GB)

Why it's interesting

On a single RTX 5090, qwentin is the only setup that serves 27B at the full 256k context
window while keeping near-FP8 quality — and it does so at interactive speed.

vLLM's 4-bit path is ~15% faster on raw decode (fewer weight bytes) — but at −5.5 points of
quality and it cannot reach 256k on a 32 GB card. qwentin also ships a 4-bit weight mode
(E2M1) that matches NVFP4 quality (86.5) at ~134 tok/s with the full 256k window, if you want
those bytes. The interactive lever qwentin leans on instead is speculative decoding (MTP),
which is lossless w.r.t. quality. (Quality ladder, same protocol: FP8 95.94 > FP6 91.30 >
E2M1 86.46 ≈ NVFP4 85.78.)

Highlights

• Fits 27B + 256k in 32 GB. FP6 (E2M3) block-scaled tensor-core weights (~20 GiB) + a
  4-bit-K / E4M3-V KV cache. Needle-in-a-haystack retrieval: 4/4 @ 239k (depths 2k–235k).
• MTP speculative decoding. A multi-token-prediction covering tree + batched FP6 verify
  (each weight read once for the whole tree) → real single-stream speedup at accept-length ~2.6–3.0.
• Wide prefill. A dedicated N-wide prefill path — wide FP6 GEMM + chunkwise-parallel
  gated-DeltaNet + tensor-core wide attention — makes cold prompts 2–4× faster at any
  length (works at Q4-KV/256k; the old 16k cap is gone). ON by default in the server, fully
  integrated with spec-decode and the prefix cache.
• Hybrid attention. Qwen3.6 mixes 48 gated-DeltaNet linear-attention layers (O(1) state,
  context-independent) with 16 full-attention layers — the engine implements both.
• Custom SM120 tensor-core kernels. Hand-written inline-PTX Blackwell block-scaled MMA
  (mma.sync ...kind::mxf8f6f4.block_scale) + FP6 ldmatrix ...b6x16 unpack, built with the stock
  CUDA 13 nvcc/ptxas toolchain — no external assembler or precompiled cubins.
• OpenAI-compatible server with cross-turn prefix/KV caching, tool-calling, and a
  reasoning/answer split.
• Batched / multi-client throughput tier. A paged-KV + continuous-batching server
  (serve_batched.py) serves N concurrent requests on one engine — 1159 tok/s @ N=32, and
  with the opt-in 4-bit (E2M1) weights ~4 clients @128k (beating vLLM's ~2-seq long-context cap
  on the same dense 27B/5090). Default-off; the FP6 single-stream path is untouched.

Research engine. The target is single-stream latency/quality on one RTX 5090, not portability —
it is SM120-only and wired for the Qwen3.6-27B layout.

Performance

Current build (2026-07) — Qwen3.6-27B, FP6 weights + Q4 KV, single stream, MTP
spec-decode, cold prefill on the default wide+MMA path. Measured on an RTX 5090
(170 SM), 256k ship config (TQ_KV_Q4=1 TQ_EMBED_FP8=2). Decode tok/s moves a few
percent with the accept-length at that text offset; ms/round is the hardware truth:

Steady-state VRAM @256k ≈ 30.8 / 31.4 GiB (the TQ_EMBED_FP8=2 6-bit embed table
is what makes 256k fit — without it the 32 GB card OOMs past ~230k).

The decode column is nearly flat: a 200k-deep conversation decodes at ~83% of the
short-context speed (27.2 vs 21.7 ms/round). Short contexts (<8k) are bit-identical
to the 2026-06 build; the long-context gains come from a producer/consumer group
attention kernel (one 512-thread CTA: 8 warps score a whole kv group's K read ONCE
per super-tile while 8 warps run the previous tile's P·V from a double-buffered
smem slab), fused Q4 scale/code loads, key-split prefill attention and a
context-gated standalone attention path — needle retrieval on this card: 4/4 @120k
and 4/4 @239k (ship config; @24k 4/4 with the bf16 embed table). End-to-end on
the server: a cold 10.8k-token first turn drops from ~15 s to 5.5 s; a follow-up
turn hits the prefix cache (10752 tokens reused) and prefills only the new suffix in
0.074 s.

<details>
<summary>RTX PRO 6000 Blackwell (188 SM, same GB202/SM120 class, same build)</summary>

Short-context rows and cold prefill are within a few percent of the 5090 (same
silicon class). These rows predate the one-wave grid sizing (measured at the old
default chunk count), so the bigger die has a little headroom left on top.
</details>

Batched / multi-client (serve_batched.py — paged KV + continuous batching). Aggregate decode
throughput scales with concurrency N (FP6; measured on the 2026-06 build, whose single-stream
reference was 70 tok/s):

The opt-in 4-bit (E2M1) weights push N=32 to 1351 tok/s and raise capacity from ~2 to ~4
clients @128k (~98 short clients). vLLM on the same dense 27B/5090 caps at ~2 concurrent @128k —
the hybrid architecture's O(1) DeltaNet state (~155 MiB/client) is the short-client floor (a
long-context advantage, a fixed tax at short ctx). Native ragged batched prefill keeps the
end-to-end HTTP path close to the steady-state ceiling.

How it works

HuggingFace Qwen3.6  ──convert_qwen_tqf.py──▶  model.tqf  (FP6 E2M3, block-scaled, + MTP head)
                                                   │
                          libforward_qwen.so  ◀────┘   (one CUDA translation unit, SM120)
                                   │
                       tools/serve_openai.py  ──▶  OpenAI /v1/chat/completions

• Weights — FP6 E2M3, block-scaled, QMMA.SF. 6 bits/param on the tensor cores with 128-wide
  block scales; ~20 GiB for 27B.
• KV cache — 4-bit K (rotated int4 + Hadamard) + E4M3 V. This is what buys 256k in 32 GB. A
  fp32-KV mode also exists (caps ~32–40k); wide prefill runs against either.
• Speculative decoding — MTP tree. Covering-tree build → batched k-split FP6 verify over the
  tree → dense-argmax descent → single-path commit that advances the real
  DeltaNet/conv/full-attn-KV state and extends the draft trunk.
• Wide prefill. Projections become one wide FP6 GEMM (weight read once for N tokens); the 48
  DeltaNet layers run a chunkwise-parallel gated-delta kernel; the 16 full-attn layers run a
  tensor-core MMA wide attention that works against the Q4 KV cache at any length (up to 256k)
  — no length gate. The server enables the whole path by default.
• SM120 kernels. Block-scaled tensor-core MMA (mma.sync ...mxf8f6f4.block_scale), FP6
  ldmatrix ...b6x16 unpack, fused DeltaNet chains, and the persistent-MLP GEMV are hand-written
  as inline PTX, compiled by the stock CUDA 13 nvcc/ptxas.

Requirements

• RTX 5090 (SM 12.0 / Blackwell GB202, 32 GB) — kernels are compute_120f-only. (Other SM120
  Blackwell parts, e.g. RTX PRO 6000, also work.)
• CUDA Toolkit 12.x/13.x with a driver new enough for SM120.
• Python 3.10+ with torch, transformers, numpy, safetensors.
• A Qwen3.6-27B (or Qwen3.5) HuggingFace checkpoint to convert.

Build

git clone [email protected]:kacper-daftcode/qwentin.git
cd qwentin
export PATH=/usr/local/cuda/bin:$PATH          # so cmake finds nvcc
cmake -B build-qwen -DCMAKE_CUDA_ARCHITECTURES=120
cmake --build build-qwen --target qwentin-forward-qwen -j
# -> build-qwen/libforward_qwen.so

Convert a model

# FP6 (E2M3) block-scaled weights + MTP head — the production format (~22.6 GB).
# ALL of --block-scaled always, --block-layout qmma-e2m3 and TQ_EMIT_MTP=1 are
# required: --block-scaled defaults to `auto`, which on a BF16 checkpoint silently
# falls back to plain global-scale FP8 (ignoring --block-layout), and without
# TQ_EMIT_MTP=1 the MTP head is dropped (no spec-decode).
# Optional: TQ_GPU_PACK=1 quantizes/packs on the GPU via the built
# libforward_qwen.so (much faster than the ~13 min numpy path).
TQ_EMIT_MTP=1 python3 tools/convert_qwen_tqf.py /path/to/Qwen3.6-27B \
    -o /path/to/qwen3_6-27b-e2m3-mtp.tqf \
    --block-scaled always --block-layout qmma-e2m3 --block-scale-policy pow2

.tqf is qwentin's on-disk format: quantized weights in the QMMA fragment layout, the
non-quant tensors (embeddings, norms, conv1d) in bf16, and the MTP draft head for spec-decode.
Verify with python3 tools/inspect_tqf.py model.tqf before serving: a production file
reports block_scaled_e2m3: True and has_mtp_section: True (header flags 0x53d).
If you see flags: 0x1d / has_mtp_section: False (and ~28 GB instead of ~22.6), the
file is plain FP8 without the MTP head — reconvert with the exact flags above.

Serve (OpenAI API)

Single-user, latency-optimized (FP6 + MTP spec-decode)

# Production: FP6 + 4-bit KV, 256k context, MTP spec-decode (needs the MTP
# section in the .tqf — see "Convert a model"). The wide+MMA cold-prefill path
# is ON by default (2-3x faster first turns at any length; --no-wide-prefill or
# TQ_WIDE_PREFILL=0 reverts to the 16-token chunked baseline). TQ_EMBED_FP8=2
# (6-bit embed table) is required for the full 256k window on a 32 GB card.
# --model-dir is the original HF checkpoint dir (tokenizer + chat template);
# run from the repo root so the relative --lib path resolves.
CUDA_VISIBLE_DEVICES=0 TQ_CTX=262144 TQ_KV_Q4=1 TQ_EMBED_FP8=2 \
    python3 tools/serve_openai.py --port 8000 --no-thinking \
    --tqf /path/to/qwen3_6-27b-e2m3-mtp.tqf \
    --model-dir /path/to/Qwen3.6-27B \
    --lib build-qwen/libforward_qwen.so

curl localhost:8000/v1/chat/completions -d '{
  "messages": [{"role": "user", "content": "Write a haiku about tensor cores."}],
  "temperature": 0.0, "max_tokens": 64
}'

Per-request stats are returned under x_qwentin (prefill seconds, accept-length, gen tok/s,
prefix-cache hit). --no-thinking defaults enable_thinking=false (recommended for agents — it
keeps the prefix cache valid across turns).

Batch-optimized: many concurrent clients (paged KV + continuous batching)

# Ship config for one 32 GB card: 12 slots over a shared paged pool
# (--num-blocks 1536 x --page 128 = 196k pooled KV tokens next to the FP6
# weights), each client up to TQ_CTX tokens. Without TQ_CTX the engine caps
# every client at 8192 tokens.
CUDA_VISIBLE_DEVICES=0 TQ_CTX=131072 TQ_KV_Q4=1 \
    python3 tools/serve_batched.py --port 8000 \
    --max-slots 12 --num-blocks 1536 \
    --tqf /path/to/qwen3_6-27b-e2m3-mtp.tqf \
    --model-dir /path/to/Qwen3.6-27B

Add TQ_W_E2M1=1 in front for the opt-in 4-bit (E2M1) throughput tier — more clients
(~4 @128k) and ~17% more aggregate tok/s, at NVFP4-level quality (the FP6 default keeps
the quality ceiling).

serve_batched.py admits/decodes/detaches N streams against a shared paged KV pool;
serve_openai.py stays the latency-optimized single-stream path (FP6, the quality default).

Key environment flags

Verify

The verify tools take the same --tqf / --model-dir / --lib paths as the servers
(needle_check reads them from TQ_MODEL_TQF / TQ_MODEL_DIR / TQ_LIB):

# End-to-end MTP spec-decode (tok/s, accept-length, divergence vs greedy)
CUDA_VISIBLE_DEVICES=0 python3 tools/mtp_spec_smoke.py --prompt-tokens 1024 --gen 128 \
    --tqf /path/to/qwen3_6-27b-e2m3-mtp.tqf --model-dir /path/to/Qwen3.6-27B \
    --lib build-qwen/libforward_qwen.so

# Decode-only round benchmark: prefill once, time M production spec-rounds
# (ms/round, net tok/s, accept-length; --profile brackets the timed rounds
# for `nsys profile -c cudaProfilerApi`). TQ_EMBED_FP8=2 keeps TQ_CTX=262144
# inside 32 GB on a 5090.
CUDA_VISIBLE_DEVICES=0 TQ_KV_Q4=1 TQ_CTX=262144 TQ_EMBED_FP8=2 \
    python3 tools/bench_rounds.py --prompt-tokens 65536 --rounds 200 \
    --tqf /path/to/qwen3_6-27b-e2m3-mtp.tqf --model-dir /path/to/Qwen3.6-27B \
    --lib build-qwen/libforward_qwen.so

# Needle-in-a-haystack retrieval quality at long context
CUDA_VISIBLE_DEVICES=0 TQ_CTX=16384 \
    TQ_MODEL_TQF=/path/to/qwen3_6-27b-e2m3-mtp.tqf \
    TQ_MODEL_DIR=/path/to/Qwen3.6-27B \
    TQ_LIB=build-qwen/libforward_qwen.so \
    python3 tools/needle_check.py

Repository layout

src/forward_qwen.cu        the Qwen3.6 engine: all kernels + C ABI, one CUDA TU (~22k lines)
tools/serve_openai.py      single-stream OpenAI server (prefix cache, tools, reasoning split)
tools/serve_batched.py     multi-client OpenAI server (paged KV + continuous batching)
tools/mtp_spec_smoke.py    spec-decode harness (also exports prefill() used by the server)
tools/bench_rounds.py      decode-only spec-round benchmark (ms/round, net tok/s; nsys hook)
tools/accept_probe.py      teacher-forced accept probe (degeneration-free draft-quality A/B)
tools/persist_phases.py    phase-clock breakdown of the persistent kernel (work vs barrier tails)
tools/serve_smoke.py       2-turn cold + prefix-cache E2E smoke against a running server
tools/paged_smoke.py       batched/paged decode parity + throughput harness
tools/needle_check.py      long-context retrieval gate
tools/convert_qwen_tqf.py  HuggingFace Qwen -> .tqf converter (+ convert.py, sparse_pack.py)
tools/mtp_*.py             MTP draft-head training / accept-length eval / dump
tools/inspect_tqf.py       inspect a .tqf model file

License

MIT — see LICENSE.