Why llama.cpp's Blackwell NVFP4 Prefill Was Starving Tensor Cores

A fast matrix-multiply instruction does not guarantee a fast matrix multiply.

That sounds obvious, but it is easy to forget when looking at Blackwell and NVFP4. NVIDIA has extremely fast low-precision tensor-core machinery. If a local inference runtime already emits the right FP4 MMA instructions, it is tempting to assume that the remaining performance gap must be small.

A new llama.cpp pull request is a useful counterexample.

PR #28572, opened September 7, targets one narrow path: NVFP4 matrix multiplication on Blackwell during prompt processing. The patch does not introduce a new quantization format or change model math. It changes how quantized tiles reach the tensor cores and how the kernel overlaps memory movement with computation.

The contributor reports Qwen3.8-27B NVFP4 prefill on an RTX 5090 increasing from 6,146.5 ± 9.8 tokens/s to 7,009.7 ± 9.1 tokens/s, a 14.0% improvement, with llama-bench -ub 1024 -p 16384 -n 0 -r 3. Those are contributor measurements from the PR, not RAMGPT benchmarks, and the PR is still open as of September 8.

The interesting part is not the percentage. It is what had to change to get it.

The tensor cores were waiting for their food

The PR description is unusually direct: NVFP4 prefill was “starving the tensor cores.”

The old path made threads participate in moving pieces of a tile into shared memory and then synchronize at a block-wide barrier before the MMA work could proceed. That is functional, but it leaves a powerful execution pipeline waiting on data movement and synchronization.

The patch attacks the feeding problem in three related ways:

1. overlap tile loads with MMA work;
2. reduce register spilling in the accumulation path;
3. use Blackwell’s available asynchronous bulk/tensor-copy machinery so threads spend less effort manually shepherding data.

This is a classic GPU optimization problem wearing an LLM-inference costume. Arithmetic throughput is only useful when the pipeline can keep operands ready.

NVIDIA’s CUDA documentation describes exactly this copy/compute pattern: data moves from global memory into shared memory, computation consumes the shared-memory tile, and asynchronous mechanisms allow the movement to proceed without tying up the threads that will later do the arithmetic. TMA extends that model with hardware-supported bulk and multidimensional copies.

The llama.cpp patch is interesting because we can see that hardware concept colliding with the ugly details of a real quantized format.

NVFP4’s 36-byte block makes the load path awkward

The patch adds a four-byte cp.async helper specifically because an NVFP4 block is 36 bytes and therefore does not naturally satisfy the wider alignment assumptions that make many tiled load paths pleasant.

That detail matters.

Optimization discussions often treat “4-bit” as if it describes the whole storage problem. It does not. A quantized block includes values plus scaling metadata, and the exact block layout determines which memory instructions are convenient, how much address work threads perform, and how easily data can be staged for the MMA instruction layout.

In the earlier profiling discussion #28514, the contributor reported that the stock kernel’s tensor pipeline was only about 42% active and identified stalls around four-byte asynchronous copies. Again, this is the contributor’s Nsight Compute observation, not an independently reproduced RAMGPT result.

The new PR reorganizes the path so raw NVFP4 rows can be staged and consumed differently. It introduces asynchronous copy grouping and waiting, bulk asynchronous copies, mbarrier synchronization, and a 2D tensor-copy helper using cp.async.bulk.tensor.

NVIDIA documents TMA as a mechanism for moving multidimensional tiles between global and shared memory while offloading address-generation work that would otherwise be performed by CUDA threads. That is almost a perfect description of the opportunity here: stop spending so much thread machinery on delivering the next tile.

This is not just “use TMA”

The patch is more instructive if we resist reducing it to a feature checklist.

“Blackwell supports TMA” does not mean replacing a load instruction with a TMA instruction automatically yields 14%.

A pipelined kernel needs a producer/consumer protocol. One tile is being consumed while another is prepared. Completion has to be tracked correctly. Shared-memory buffers cannot be overwritten before their consumers finish. The next MMA cannot read a tile before the asynchronous copy completes.

That is why the patch adds mbarrier operations and explicit parity-based waiting around bulk copies. NVIDIA’s programming guide describes asynchronous barriers as a split arrive/wait synchronization mechanism; TMA operations can signal completion through memory barriers rather than forcing every thread through the old load-and-block pattern.

The practical change is architectural: data movement becomes a pipeline stage instead of a pause between compute stages.

The patch also addresses register spills in the accumulation step. That matters because a theoretically faster memory path can still lose if additional live state pushes registers beyond what the compiler can keep resident. Once temporary values spill into local memory, the optimization starts paying memory traffic to save memory traffic.

This is why kernel performance work rarely has a single knob.

The +14% patch came out of a much larger +44.7% experiment

PR #28572 is actually the conservative part of a larger story.

One day earlier, in llama.cpp discussion #28514, the same contributor posted a proof of concept comparing stock llama.cpp, a three-part experimental patch set, and NInfer on an RTX 5090 with Qwen3.8-27B NVFP4.

Their reported pp16384 results were:

The full PoC was therefore reported at +44.7% over stock. But that number combined several independent ideas, including chunking a gated-delta-net prefill path, fusing operations around GEMMs, improving the NVFP4 MMQ tile pipeline, and experimenting with the MTP hand-off.

That is not a number we should casually attach to llama.cpp itself. The PoC was explicitly presented as experimental, the author disclosed AI assistance in the kernel work, and some verification remained incomplete. The current PR extracts only the NVFP4 MMQ pipeline work into a reviewable change and reports +14.0% for that narrower patch.

From a production-engineering perspective, that separation is the most important part of the story.

A 45% prototype speedup is exciting. A smaller patch with a bounded scope, correctness checks, a control benchmark, and an understandable mechanism is much easier to reason about upstream.

The verification is better than the headline

PR #28572 includes three useful checks.

First, the contributor reports WikiText-2 perplexity of 7.1927 ± 0.0469 on master and 7.1957 ± 0.0469 with the patch. The values are close enough that the reported performance improvement is not accompanied by an obvious quality collapse in that test.

Second, CUDA backend operation tests pass for the touched matrix-multiply paths: 1288/1288 MUL_MAT and 880/880 MUL_MAT_ID.

Third, a Q5_K_S control on the same RTX 5090 is essentially unchanged: reported pp4096 is 3199.5 versus 3199.9 t/s, while tg64 is 71.1 versus 70.8 t/s.

That last check is particularly useful. The code is gated to the Blackwell NVFP4 path. A conventional quant showing no meaningful movement supports the claim that this is a targeted kernel-path optimization rather than a broad benchmark-environment change.

It still does not turn one contributor’s RTX 5090 measurements into a universal 14% promise. Power limit, CUDA version, model, prompt length, ubatch size, compiler behavior and future code changes can all move the result.

But this is the kind of upstream performance evidence worth paying attention to: mechanism, patch, test command, baseline commit, hardware, correctness checks and a negative control are all visible.

There is another bottleneck hiding behind MTP

The earlier discussion contains a second result that deserves attention even though it is not part of PR #28572.

The contributor reports that enabling MTP speculative drafting cost substantially more prefill performance in their llama.cpp experiment than in the NInfer comparison. Profiling attributed part of that gap to hidden-state hand-off traffic through the host on every ubatch. Their experimental device-side hand-off reduced measured GPU idle time under profiling.

This is a different problem from tile loading, but it points to the same systems lesson.

A GPU can have fast tensor cores and a fast speculative-decoding algorithm and still lose throughput because data crosses the wrong boundary at the wrong frequency.

Local inference optimization increasingly looks less like “make GEMM faster” and more like remove bubbles from the whole execution graph: global-to-shared movement, recurrent state, kernel fusion, host/device transfers, synchronization and scheduler behavior.

Why this matters beyond RTX 5090 owners

Most llama.cpp users do not own Blackwell hardware and many do not run NVFP4 models. The patch is still useful because it shows where the next layer of inference performance is coming from.

Quantization originally gave local inference an obvious win: move fewer weight bytes. Then specialized kernels learned to compute directly on quantized representations. Now, as low-precision tensor cores become extremely capable, the bottleneck moves again.

The runtime has to feed those units efficiently.

That means the physical representation of a quant block, shared-memory layout, asynchronous copy mechanism, register footprint and synchronization strategy can determine whether impressive peak hardware numbers appear in real prompt processing.

This is also why comparing runtimes only by supported model format is becoming less informative. Two runtimes can both “support NVFP4” while using very different pipelines to reach the same tensor-core instruction.

Support is binary. Performance engineering is not.

What I would watch before calling this done

PR #28572 is open, so the first question is whether maintainers accept the implementation in its current form or reshape it during review.

The second is whether the remaining ideas from discussion #28514 become separate PRs. The contributor specifically points to gated-delta-net prefill and GLU/MMQ fusion as future work. Those changes attack different bubbles and should be evaluated independently rather than bundled into the original +44.7% PoC number.

The third is MTP. If speculative drafting requires repeated host-mediated hidden-state movement, optimizing that hand-off could matter well beyond this exact NVFP4 kernel.

The broader lesson is already clear enough to publish without waiting for those patches:

Blackwell’s low-precision tensor cores were not the whole performance story. In this llama.cpp path, getting quantized tiles to them efficiently was worth a reported 14% by itself.

That is a much more useful result than another peak-TOPS number. It tells us where the software stack was actually leaving performance on the table.