# Serving 30B models at 300 tok/s on a single RTX 5090 > Why no existing inference engine fully exploits consumer Blackwell, what NVFP4 changes, and the numbers from building one that does. Source: https://rfriedmann.de/blog/serving-30b-models-rtx-5090/ Published: 2026-06-02 · Track: log · Level: Advanced A single RTX 5090 decodes 30B-total MoE models (~3B active parameters) at around 300 tokens per second, single-stream, batch 1. Not in a lab, not extrapolated, but measured, on my desk, reproducible down to the commit SHA. The surprising part isn't the number. It's that no existing engine has a native-NVFP4 decode path on `sm_120`, so none matches imp's single-stream FP4 decode on this card. That gap is the whole reason [imp](https://github.com/kekzl/imp) exists. This post is about what's inside it, and why a $2,000 graphics card was leaving most of its own silicon switched off. ## A 5090 is not a small data centre GPU On paper the 5090 looks like a shrunk B200: same Blackwell architecture, FP4 tensor cores, 32 GB of fast memory. In practice the consumer chip (`sm_120a`) is a different machine wearing the same family name. It lacks the async-compute path that data centre Blackwell kernels are built around (tcgen05/TMEM, wgmma); it keeps the async-copy engine (TMA, cp.async).

Datacenter Blackwell - B200 (sm_100)

Consumer Blackwell - RTX 5090 (sm_120a)

That mismatch splits the ecosystem cleanly, and not in the 5090's favour: - **vLLM** is built for data centre Blackwell, where it's excellent. Its NVFP4 path gates on `tcgen05`, an instruction the 5090 doesn't have, so on a consumer card it falls back to slower kernels or won't load that path at all. Different target, not a flaw. - **llama.cpp** runs on everything, on purpose, which is a real strength and the reason it's the backbone of local inference. Its decode is genuinely well-built (quantized integer kernels, no FP16 detour). But staying universal means it can't lean on one chip's quirks, so the 5090's native FP4 tensor cores stay idle. So the one consumer card with native FP4 hardware had no engine that spoke native FP4. These 30B-MoE models fit comfortably in 32 GB once quantized, and the memory bandwidth says they *should* decode north of 250 tok/s. Nobody was delivering that. So I built the thing that would. ## What NVFP4 actually buys you NVFP4 is the native weight format for Blackwell's FP4 tensor cores: 4-bit floating point with a per-block scale, exported straight from NVIDIA's Model Optimizer or llm-compressor. Here's the one fact that makes it matter. Decode speed on a single GPU is almost entirely a memory-bandwidth problem: every token you generate has to drag the active weights through the chip. Halving the bytes per weight (FP4 vs FP8), or quartering them (FP4 vs FP16), moves the ceiling by the same factor. Fewer bytes to move, more tokens per second. That's the entire trick. The catch is the part nobody hands you: there are no ready-made FP4 matrix kernels for `sm_120`. If you want this, you write the GEMM yourself, in imp's case `mma.sync mxf4nvf4` with FlashAttention-2-style block scaling, built on CUTLASS, because the Hopper and B200 kernel designs simply cannot run on this chip. ## The numbers Single RTX 5090, greedy decoding, CUDA 13.3, CUDA Graphs on. Every figure here is dated and commit-anchored in [BENCHMARKS.md](https://github.com/kekzl/imp/blob/main/BENCHMARKS.md), with the exact command to reproduce it.
NVFP4 decode throughput - single RTX 5090 (tok/s, single-stream, batch 1, tg256)
[diagram omitted — see the page for the chart]
Native-NVFP4 decode on consumer Blackwell, single-stream (batch 1), greedy, CUDA 13.3. tg256 = 256 generated tokens at short context. Measured per BENCHMARKS.md.
There's no "vs." column on that chart for a simple reason: on `sm_120` there's nothing to compare against. vLLM's NVFP4 path needs `tcgen05`, and llama.cpp has no native NVFP4 support at all. Native-NVFP4 decode on consumer Blackwell is, at the moment, a field of one. Where there *is* a fair comparison, namely dense GGUF, which is llama.cpp's home ground, imp decodes Qwen3-8B Q8_0 at about 269 tok/s, **at least 37% faster than llama.cpp** (b8445+, full offload, flash attention on), and up to 72% depending on the model. The trick is a small one: GGUF weights get converted to FP4 once at load, so the decode hot path streams sub-byte weights while prefill keeps the full-precision source. ## Where imp loses Benchmarks that only show wins are advertising. So here's the honest list, which, pleasingly, has grown shorter since I first drafted this post, because prefill was the weak spot and prefill got the attention: - **GGUF prefill used to be a clear loss; now it's roughly even.** A new INT8 prefill path (default since #617) lifted throughput by +36% to +151% per model. imp now edges llama.cpp on MoE (Qwen3-30B-A3B, +7%) and dense (Qwen3-14B, +1.5%), and trails only on Gemma-4-26B (1.20×) and one hybrid model (1.55×). - **NVFP4 prefill is basically at parity.** Reworking the attention kernel (#687) put MoE long-context prefill *ahead* of vLLM and cut the dense gap from 1.27× down to 1.04×. To be precise about what that beats: on `sm_120` vLLM runs a fallback path, since its tcgen05 NVFP4 kernels don't run on consumer Blackwell, so this is ahead of vLLM's fallback, not its optimised data centre path. The one survivor is dense long-context prefill, losing by about 4%. - **Qwen3.6-35B GGUF decode still loses ~31%** to llama.cpp, a structural cost on that model's particular attention design that FP4 weights can't recover. - **Batched, aggregate throughput under concurrency is not imp's axis.** Every number here is single-stream, batch 1; for many concurrent requests, engines like vLLM lead and that's where they're built to win. - It's **experimental, single-GPU, single-author**. llama.cpp is still the right choice if you want mature and universal; vLLM if you batch on data centre cards. imp answers exactly one question: how fast can a single 5090 go? ## The part I still find slightly absurd imp is roughly 97k lines of C++20 and CUDA (GGUF and SafeTensors loaders, a BPE tokenizer, a paged KV cache, the CUTLASS NVFP4 dispatch, the attention kernels, MoE routing, an OpenAI- and Anthropic-compatible server) and every line of it was written by [Claude Code](https://claude.ai/claude-code). My job is the architecture, the model choices, the benchmarking, and the direction. The agent writes the code. We're 1,000+ commits in. How that workflow holds up at this scale is a story for its own post. If you've got a 5090 and want to watch it run: everything ships as Docker images on [GHCR](https://github.com/kekzl/imp/pkgs/container/imp), no local CUDA toolkit required. The repo is [github.com/kekzl/imp](https://github.com/kekzl/imp): MIT, experimental, and very much a single-architecture bet.