RTX 5090 Decode Bandwidth & 25% Penalty

The 25% decode penalty imp reports between the RTX 5090 and RTX 5090D directly validates a bandwidth-bound decode model, and the arithmetic implied by that figure puts imp’s territory in sharp relief against everything else in the inference landscape. Working backward from imp’s stated 1.34 TB/s ceiling for the 5090D and the claimed ~25% throughput gap, the RTX 5090’s effective decode bandwidth during inference is approximately 1.79 TB/s—roughly 26× the 68 GB/s ceiling we measured on Jetson Orin Nano.

What imp claims, and the conditions attached

imp reports a +42–48% decode advantage over llama.cpp on dense GGUF, tied to a specific llama.cpp commit (b9976, 2026-07-12). The vLLM comparison—“at-or-ahead on NVFP4”—is explicitly on NVFP4 inference, and imp qualifies it further: vLLM’s relevant path requires tcgen05, an opcode family absent on consumer Blackwell (sm_120a), so vLLM either falls back to Marlin or fails outright on the RTX 5090. The comparison is less “imp beats vLLM” and more “vLLM’s fast path doesn’t compile for this hardware.”

The Qwen3.6-35B decode figures tell a more granular story. imp reports 224 → 267–272 tok/s on GGUF after extending an FP8 SSM-projection sidecar to UD quants (post #962), and 257 → ~320 tok/s on the NVFP4 checkpoint path after the same sidecar (post #949). The GGUF gain is described as +18% ahead of llama.cpp; the +42–48% headline applies to dense GGUF generally. imp also discloses that cuBLAS autotuning introduces up to 2.6× variance in prefill numbers across container restarts—this appears in the documentation as a known issue, not a footnote.

Backend	Hardware	Format	Decode	Conditions
imp	RTX 5090	dense GGUF	+42–48% vs llama.cpp	b9976, 2026-07-12
llama.cpp	RTX 5090	dense GGUF	baseline	b9976, 2026-07-12
imp	RTX 5090	Q8_0/Q6_K → NVFP4 cache	267–272 tok/s	Qwen3.6-35B, post #962
imp	RTX 5090	NVFP4 checkpoint	~320 tok/s	Qwen3.6-35B, post #949
imp	RTX 5090D	any	~25% below RTX 5090	BW-bound; 1.34 TB/s stated
vLLM	sm_100 (B200/B300)	NVFP4	at-or-behind imp	requires tcgen05; inapplicable to sm_120a

Analysis

Decode throughput in memory-bandwidth-bound regimes is governed by:

$$\text{tok/s} \approx \frac{\text{BW}_\text{mem}}{W_\text{bytes}}$$

where $W_\text{bytes}$ is the active weight footprint moved per token. If the relationship is proportional—and imp explicitly invokes it when explaining the 5090D gap—then the ratio of decode speeds equals the ratio of bandwidths. Substituting the supplied figures:

$$\text{BW}_\text{5090} = \frac{\text{BW}_\text{5090D}}{1 - 0.25} = \frac{1.34\,\text{TB/s}}{0.75} \approx 1.787\,\text{TB/s}$$

This bound anchors the +42–48% llama.cpp gap. If imp’s NVFP4 path is moving 4-bit weights where llama.cpp is moving Q8_0 (8-bit) weights over the same bus, the theoretical bandwidth ceiling is a 2× advantage. A realized ~45% gain—within the stated 42–48% range—is consistent with that ceiling after accounting for routing, attention, and kernel-launch overhead. The gap is not implausibly large; it fits the bandwidth model without invoking any deeper magic.

For comparison, the Jetson Orin Nano’s 68 GB/s—from our edge AI benchmark—sits 1,340 / 68 ≈ 19.7× below the 5090D’s stated floor, and 1,787 / 68 ≈ 26.3× below the derived RTX 5090 ceiling. These ratios define what separates these performance tiers, independent of any software choice.

The NVFP4 decode cache is architecturally worth examining separately from the throughput claims. imp converts Q8_0/Q6_K weights to NVFP4 at init time and holds the result in VRAM—a one-time compute cost exchanged for a persistent halving of bytes-per-token on every subsequent step. For agentic loops where a 30B-class model is queried hundreds of times per session, the amortization is essentially immediate. The NVFP4 KV cache (--kv-nvfp4) extends the same logic to attention memory, claiming 4× context compression at decode parity.

The SWA-aware KV sizing is less discussed but operationally significant. imp reports allocating only the sliding-window portion of KV for models that use it—~5–6× fewer KV tokens for gemma-3/4, ~2× fewer for gpt-oss. On a 32 GB card running a model that fits by a narrow margin, that allocation reduction is the difference between a long agentic session fitting in memory and OOM mid-loop.

The ISA gap deserves explicit treatment because it is structural, not contingent on software versions. sm_120a is missing tcgen05, TMEM, wgmma, and TMA warp-specialized grouped GEMM. These aren’t secondary features—wgmma is the primary GEMM instruction for dense matmul on Hopper and datacenter Blackwell, and tcgen05 gates FlashInfer’s MoE dispatch and cuBLASLt’s FP4 kernels. cuBLASLt has no FP4 implementation for sm_120. imp reports that CUTLASS v4.6.1 using mma.sync mxf4nvf4 is therefore the only viable route for FP4 on this hardware. The uniqueness claim follows from the ISA, not from a competitive survey.

What cannot be concluded

The +42–48% figure is a snapshot against llama.cpp b9976 from 2026-07-12; llama.cpp’s RTX 5090 support has been under active development and the gap could compress or widen depending on what merged after that commit. The “at-or-ahead of vLLM” framing is structurally difficult to evaluate on consumer Blackwell because vLLM’s NVFP4 code path doesn’t execute there—the comparison cannot be run on the same hardware. The studio has not measured imp, and both the 5090D bandwidth figure (1.34 TB/s) and the 25% performance gap are imp’s own reported values; the derived 1.787 TB/s ceiling is a bound that inherits whatever measurement uncertainty those figures carry.

What this means for itria

itria runs llama.cpp with Metal acceleration on iPhone and iPad. The RTX 5090’s 1.79 TB/s ceiling is not directly operational for itria, but two design patterns transfer.

The init-time weight conversion—pay a one-time compute cost to reformat weights into the fastest inference format for the target hardware—is exactly the question itria should be pressing on its Metal path. On Apple Silicon, the equivalent is whether weights are held in the format that maximizes the GPU’s memory access efficiency for the specific SoC, or whether they remain in the download-friendly GGUF format by default. The logic imp uses for NVFP4 cache construction at startup applies wherever format and bandwidth interact.

The 2.6× prefill variance from cuBLAS autotuning is a concrete warning for any engine relying on a tuning cache that may be cold at process start. itria’s persistent local process model reduces this exposure compared to a containerized backend, but any Metal kernel warm-up that runs on first inference rather than at install time has the same structural exposure on cold launch.

The SWA-aware KV allocation and NVFP4 KV compression strategies are directly relevant to future itria sessions on long-context tasks within the constrained VRAM budgets of mobile devices. The principle—allocate only the KV tokens the attention pattern actually reads—applies regardless of whether the backend is CUDA or Metal. For related discussion of where software overhead sits in inference loops, see the agentic DSP verification note.

FAQ

Is the +42–48% decode gap reproducible on non-5090 Blackwell GPUs?

imp ships a compute_120f PTX fallback for the RTX 5080 and 5070 Ti, but PTX is not SASS—it will JIT-compile on first run without the instruction-level tuning that generates the raw sm_120a SASS imp emits for the 5090. The gap figure is tied to the RTX 5090 measurement conditions (b9976, 2026-07-12); imp makes no corresponding performance claim for the fallback path, and the architectural tuning that produces the headline number is 5090-specific.

Does the NVFP4 decode cache degrade model quality?

imp doesn't supply quality benchmarks in the available claims. The GGUF cache path converts Q8_0/Q6_K weights to NVFP4 at init time and is characterized as getting "the bandwidth benefit of sub-byte weights"—throughput is the stated goal. Quality implications of converting from the source quantization to FP4 are not addressed in the supplied documentation, and the studio has not evaluated this.

How serious is the 2.6× prefill variance in production?

imp explicitly discloses that cuBLAS autotuning causes up to 2.6× variance in prefill numbers across container restarts. For agentic workloads, prefill cost maps to time-to-first-token for every tool call that includes a long context; a 2.6× range makes TTFT SLAs unreliable until the tuning cache has stabilized. This is disclosed as a known structural property of cuBLAS autotuning on this hardware, not an intermittent anomaly.

Does any of this apply to the Metal/llama.cpp stack?

The bandwidth-bound decode model applies universally—decode speed is ultimately constrained by how fast weights can be moved from memory relative to compute. The specific FP4 ISA (mma.sync mxf4nvf4), CUTLASS kernels, and CUDA-specific paths are not portable. The practical transfer is the design pattern: prefer weight formats that minimize bytes-per-token at inference time, and pay format conversion costs at init rather than at every decode step.