perf(cuda): README vs llama.cpp CUDA benchmark gaps — full results + tracking

[pekkah]

Tracking issue for the CUDA performance gaps surfaced by a head-to-head of the README's recorded CUDA speeds vs llama.cpp CUDA on the same models and the same box. Records the full results, the methodology, and links the per-finding sub-issues so they aren't re-derived.

Methodology

• Hardware: RTX 4070 Ti (12 GB) + Ryzen 9 7900X (Zen 4). Same box for both engines.
• llama.cpp: CUDA build b9529 (C:\p\llama.cpp-cuda\bin\llama-bench.exe).
• README numbers: taken verbatim from the per-model tables in README.md (prefill = warm ~2K-token working-context rate; decode = near-zero-ctx headline, ≈ llama.cpp tg128).
• llama.cpp — full-offload models: live llama-bench -ngl 99 -p 2048 -n 128, so pp2048 matches the README's ~2K prefill convention and tg128 matches its decode convention.
• llama.cpp — CPU-MoE hybrids: documented -ngl 99 -ncmoe 99 (pp512/tg128) from prior on-box runs (live-running these is heavy CPU-expert inference). Prefill is pp512 (not 2K) so those prefill gaps are approximate; the hybrid decode gaps carry thermal/MTP caveats.
• Cross-check: the live numbers land on the README's own inline llama.cpp citations — Gemma4-E4B Q8 8577 / 78.2 vs README's "~8475 / ~78", Gemma4-12B 57.2 vs "57", Qwen3-8B 6010 vs README's pp1008 5764. Methodology is consistent.

Full results

Full GPU offload (live llama-bench, -ngl 99, pp2048 / tg128):

Model	README pf / dec	llama.cpp pf / dec	Prefill gap	Decode gap
OLMoE-1B-7B Q4_K_M	97.4 / 126.0	17 137 / 426.1	176×	3.4×
SmolLM2-1.7B Q4_K_M	194 / 268.0	20 192 / 333.6	104×	1.24×
Qwen3-8B Q4_K_M	2 287 / 74.5	6 010 / 91.1	2.63×	1.22×
Gemma4-E4B q4_0	3 666 / 100.4	9 236 / 124.4	2.52×	1.24×
Gemma4-12B q4_0	1 714 / 54.1	4 201 / 57.2	2.45×	1.06×
Gemma4-E4B Q8	4 145 / 70.5	8 577 / 78.2	2.07×	1.11×

CPU-MoE hybrids (documented -ngl 99 -ncmoe 99, pp512 / tg128):

Model	README pf / dec	llama.cpp pf / dec	Prefill gap	Decode gap
Carnice 35B-A3B-MTP-ft	522 / 26.5	755 / 58.3	~1.45×	~2.2× *
Qwen3.6-35B-A3B-UD	475 / 24.5	537 / 49.9	~1.13×	~2.0× *
Qwen3.6-35B-A3B-MTP	480 / 33.3	537 / 49.9	~1.12×	~1.5× *

* Decode carries thermal/MTP caveats — README notes plain decode is ~80% of llama.cpp warm-vs-warm (~1.25×); the 2× figures use a cooler reference run. See [project_decode_carnice_cpumoe] analysis.

No clean on-box baseline (not benchmarked): Qwen3-Coder-30B-A3B, Qwen3.6-27B-MTP, Llama-4 Scout (Scout not present locally; the two MoE/GDN ones would need a CPU-heavy -ncmoe run).

The four biggest differences

1. OLMoE — prefill 176× (97.4 → 17 137 t/s) → perf(cuda): batched int8-MMQ trunk prefill gated off for MoE — OLMoE prefill 97 vs llama.cpp 17,137 t/s (176×) #406
2. SmolLM2 — prefill 104× (194 → 20 192 t/s) → perf(cuda): batched int8-MMQ trunk prefill gated off for dense NORM-RoPE & QKV-bias models — SmolLM2 105×, VibeThinker/Qwen2 159× #407
3. OLMoE — decode 3.4× (126 → 426 t/s); only model lagging badly on both axes → perf(cuda): full-offload MoE decode ~3.4× behind llama.cpp — OLMoE 126 vs 426 t/s (on-GPU experts, kernel-bound) #408
4. Qwen3-8B — prefill 2.6× (2 287 → 6 010), narrowly ahead of a 2.0–2.5× dense cluster (Gemma4 family) → perf(cuda): dense int8-MMQ prefill ~2–2.6× behind llama.cpp (cp.async MMQ + flash attn) — Qwen3-8B / Gemma4 #409

Root cause (verified in code)

OLMoE and SmolLM2 both prefill below their own decode rate — the signature of a per-token prefill loop. CudaForwardPass.IsBatchedPrefillSupported() (src/SharpInference.Engine/CudaForwardPass.cs:2618) gates the fast int8-MMQ batched-trunk prefill:

if (_isMoE || _tqEnabled || _hasAttnBias || !_hp.IsNeoxRope) return false;
if (_hasQkNorm && _hp.UseL2QkNorm) return false;

• OLMoE → _isMoE trips it → per-token loop (97.4 t/s). → perf(cuda): batched int8-MMQ trunk prefill gated off for MoE — OLMoE prefill 97 vs llama.cpp 17,137 t/s (176×) #406
• SmolLM2 (Llama arch, NORM / non-NeoX RoPE) → !_hp.IsNeoxRope trips it → per-token loop (194 t/s). → perf(cuda): batched int8-MMQ trunk prefill gated off for dense NORM-RoPE & QKV-bias models — SmolLM2 105×, VibeThinker/Qwen2 159× #407

Qwen3-8B and Gemma4 satisfy the gate, get the batched MMQ trunk, and sit at the residual ~2–2.6× kernel gap vs llama.cpp's cp.async MMQ + flash-attn prefill. → #409

Sub-issues

• perf(cuda): batched int8-MMQ trunk prefill gated off for MoE — OLMoE prefill 97 vs llama.cpp 17,137 t/s (176×) #406 — batched int8-MMQ trunk prefill gated off for MoE (_isMoE) → OLMoE prefill 155× (un-gate + batched expert FFN)
• perf(cuda): batched int8-MMQ trunk prefill gated off for dense NORM-RoPE & QKV-bias models — SmolLM2 105×, VibeThinker/Qwen2 159× #407 — batched int8-MMQ trunk prefill gated off for dense NORM-RoPE (!IsNeoxRope) & QKV-bias (_hasAttnBias) → SmolLM2 105×, VibeThinker/Qwen2 159×
• perf(cuda): full-offload MoE decode ~3.4× behind llama.cpp — OLMoE 126 vs 426 t/s (on-GPU experts, kernel-bound) #408 — full-offload MoE decode ~3.4× behind llama.cpp → OLMoE 126 vs 426 (on-GPU experts, kernel-bound)
• perf(cuda): dense int8-MMQ prefill ~2–2.6× behind llama.cpp (cp.async MMQ + flash attn) — Qwen3-8B / Gemma4 #409 — dense int8-MMQ prefill ~2–2.6× behind (cp.async MMQ + flash attn) → Qwen3-8B / Gemma4 (stretch)
• perf(cuda): CPU-side prefill not batched on the hybrid passes — 27B-MTP 62–91×, Coder-30B 19× behind llama.cpp #410 — CPU-side prefill not batched on the hybrid passes → 27B-MTP 62–91×, Coder-30B 19× (batched CPU dense-FFN / expert GEMM in prefill)

Priority: #406/#407 are 100×+ and only need the existing GPU batched path un-gated → highest value. #410 is the CPU-side sibling (also 19–91×, needs a batched CPU FFN/expert GEMM). #408 needs profiling. #409 is deep kernel work (stretch).

[pekkah]

Fresh full run recorded (2026-06-26)

Re-ran both engines warm, same-session, same box (RTX 4070 Ti, llama.cpp b9529), now covering every CUDA + CUDA-hybrid pass including the CPU-MoE hybrids (-ncmoe 99) and dense-GDN 27B (-ot ffn=CPU). Full report + raw data committed at docs/benchmarks/cuda-vs-llamacpp-2026-06-26.md. SharpInference numbers reproduced the README within noise.

Two findings beyond the original four:

1. VibeThinker-1.5B (Qwen2) is a third instance of the prefill gate — prefill 148 vs llama.cpp 23,552 = 159×. Cause is the _hasAttnBias clause of the same IsBatchedPrefillSupported() gate (Qwen2 has QKV bias). So the gate disables batched MMQ prefill for three families: _isMoE (perf(cuda): batched int8-MMQ trunk prefill gated off for MoE — OLMoE prefill 97 vs llama.cpp 17,137 t/s (176×) #406, OLMoE), !IsNeoxRope (perf(cuda): batched int8-MMQ trunk prefill gated off for dense NORM-RoPE & QKV-bias models — SmolLM2 105×, VibeThinker/Qwen2 159× #407, SmolLM2/Llama), and _hasAttnBias (Qwen2/VibeThinker). Worth folding the attn-bias case into perf(cuda): batched int8-MMQ trunk prefill gated off for dense NORM-RoPE & QKV-bias models — SmolLM2 105×, VibeThinker/Qwen2 159× #407's scope (or a sibling).

2. CPU-side prefill isn't batched — a second prefill bottleneck distinct from the GPU-MMQ gate:

   • Qwen3.6-27B-MTP: prefill 9.9 (Q4) / 6.2 (Q5) vs llama.cpp 615 / 566 = 62× / 91× — dense FFN runs per-token on CPU during prefill; llama.cpp batches it. But we win decode 9.9 vs 4.46 (2.2×) via MTP self-spec + faster CPU-FFN.
   • Qwen3-Coder-30B: prefill 31 vs 588 = 19× — SLRU expert-stream-bound; never got the GPU MoE op-offload (Make GPU op-offload MoE prefill default-on — after fixing the pinned-buffer decode regression #390) the 35B GDN models did.

Updated biggest-gap ranking (ratio): VibeThinker pf 159× · OLMoE pf 155× · SmolLM2 pf 105× · 27B-MTP-Q5 pf 91× · 27B-MTP-Q4 pf 62× · Coder pf 19× · OLMoE dec 3.4×.

Hybrid decode raw ratios (1.6–2.2×) overstate the real gap: llama-bench tg128 is a tight L2-resident loop; sustained-vs-sustained our plain decode is ~80% of llama.cpp (~1.25×) — experts are RAM-floored/equal, the gap is trunk + per-layer coordination (see project_decode_carnice_cpumoe).

[pekkah]

Tracking update — CPU-side prefill batching (#410):

• perf(cuda): CPU-side prefill not batched on the hybrid passes — 27B-MTP 62–91×, Coder-30B 19× behind llama.cpp #410 closed: both work items implemented and merged in PR perf(cuda): batched CPU-side prefill for the hybrid passes (#410, opt-in) #415 (9e877b9) — batched CPU dense-FFN prefill for the 27B-MTP config (SHARPI_BATCHED_CPU_FFN) and batched CPU-MoE routed-expert prefill for Coder-30B (SHARPI_HYBRID_BATCHED_MOE, with the Q8_KS int8 tier). Both opt-in / default-off; decode untouched; bitwise parity oracles in the tree.
• Validate #410 batched CPU prefill (PR #415) on the desktop — parity + A/B bench, then decide flag defaults #416 opened: desktop validation — parity runs + the interleaved A/B vs llama-bench b9529, then the default-flip decision (or the escalation levers listed there: fused R×4 kernels, dense int8, Carnice prefill toward llama.cpp 676 t/s — per-layer rewrite levers (op-offload follow-up to #387) #388 GPU-router port, gate tuning). The rows in this issue's table for 27B-MTP and Coder-30B prefill should be re-measured with the flags ON once Validate #410 batched CPU prefill (PR #415) on the desktop — parity + A/B bench, then decide flag defaults #416 runs.

No README table changes yet — the shipped defaults leave the measured behavior unchanged.