CUDA: NVFP4 prefill on Blackwell leaves ~1.45x performance on the table. Comparison + PoC

[jesdga95]

DISCLAIMER:

The patches were written with the help of Fable 5.1, doing the profiling and kernel work, with me measuring. PoC only; I'm a software developer but this is not my area of expertise.

---

RTX 5090 (sm_120a, 600 W), CUDA 13.3, Qwen3.8-27B NVFP4, llama-bench -ub 1024 -p 16384 -r 3, no drafter.

build	pp16384 t/s
stock llama.cpp	6,122 ± 8
stock + PoC patch	8,860 ± 16 (+44.7%)
NInfer (hand-written W4A4), speculation off	8,466

Same checkpoint, same card, so the headroom is real independent of the patch.

What the patch does (three CUDA commits, gated to NVFP4 MMQ and Qwen3.5-style models):

• Chunked gated delta net prefill, 64-token chunks instead of per-token recurrence (×1.27). ncu: stock was issue-bound at one warp per scheduler.
• Fusions around the GEMMs: GLU output quantized directly into the MMQ activation format, gate/up sharing one quantized src1, transposed concat, batched-row rms_norm.
• NVFP4 MMQ tiles fed by TMA with mbarriers, FP4 MMA accumulated in place (×1.16). ncu: stock tensor pipe 42% active, stalled on 4-byte cp.async because the 36-byte NVFP4 block is only 4-byte aligned.

MTP finding. --spec-type draft-mtp cost 23% prefill (8,499 → 6,518 engine-reported) vs NInfer's MTP at −9%. nsys: the hidden state goes through the host every ubatch, four 21 MB copies per 2048 tokens, GPU idle 15%. Three more commits keep the hand-off on device (device tensor + backend event); idle 15% → 7.7%, 7,065 → 7,550 under the profiler, output is identical.

Verified: test-backend-ops passes for every op touched (MUL_MAT 1302/1302, MUL_MAT_ID 880/880, GATED_DELTA_NET 42/42, new CONCAT/CPY/NORM/GLU-fusion cases). Temp-0 output token-count-stable across builds. Not done: perplexity vs stock. tg128 82.8 → 78.9, not yet A/B tested.

Patches:
01-cuda-chunked-gdn-and-graph-fusions.patch
02-cuda-cp-async-pipelined-nvfp4-mmq.patch
03-cuda-tma-pipelined-nvfp4-mmq.patch