perf(cuda): dense int8-MMQ prefill ~2–2.6× behind llama.cpp (cp.async MMQ + flash attn) — Qwen3-8B / Gemma4

[pekkah]

Part of #405 (README vs llama.cpp CUDA benchmark sweep). Lower priority — this is the residual kernel gap after the batched MMQ prefill is already engaged (unlike #406/#407, where the batched path is gated off).

Measured (RTX 4070 Ti, llama.cpp CUDA b9529, full offload, pp2048)

Model	README prefill	llama.cpp pp2048	gap
Qwen3-8B Q4_K_M	2 287	6 010	2.63×
Gemma4-E4B q4_0	3 666	9 236	2.52×
Gemma4-12B q4_0	1 714	4 201	2.45×
Gemma4-E4B Q8	4 145	8 577	2.07×

Decode for all four is within ~1.05–1.25× (near ceiling); only prefill trails. These models do use the batched int8-MMQ trunk + flash-attn prefill, so this is the genuine kernel-efficiency gap, consistent with the README's own note for Qwen3-8B ("gap is its cp.async MMQ + flash attn").

Where llama.cpp's prefill edge comes from

1. cp.async double-buffered MMQ — llama.cpp's mmq.cu pipelines global→shared loads with cp.async so the int8 tensor-core math overlaps weight DMA; our llm_mmq_* (Gemma 4 GPU prefill is ~70x slower than llama.cpp — adopt cuBLAS GEMM for prefill #141) does a simpler load→compute. This is the biggest single piece.
2. Fused flash-attention prefill — ours is good (perf(cuda): full tensor-core (mma.sync) flash-attention prefill for d=512 — beat the half2 kernel (#141 follow-up) #146/perf(cuda): multi-warp / d-split TC flash prefill — fix the single-warp occupancy limit (#146 follow-up) #147 TC flash) but llama.cpp's fattn is more aggressively tiled/pipelined.
3. SoA weight layout (perf(cuda): split-K + better-tiled int8 MMQ GEMM — 23-34% → 60-80% of TC peak (#141 follow-up) #149) helps; a fuller vectorized/aligned load story would help more.

Scope / priority

This is a stretch optimization (2–2.6×, deep kernel work) vs #406/#407 which are 100×+ and just need the existing batched path un-gated. Recommend tackling #406/#407 first. Track here so the dense-prefill ceiling isn't re-investigated from scratch; pursue cp.async MMQ when the dense trunk becomes the priority.

References

• src/SharpInference.Engine/CudaForwardPass.cs (PrefillMmqEnabled :372, PrefillFlashTcEnabled :392, SoA _mmqSoa :394)
• MMQ kernels: Gemma 4 GPU prefill is ~70x slower than llama.cpp — adopt cuBLAS GEMM for prefill #141; SoA repack: perf(cuda): split-K + better-tiled int8 MMQ GEMM — 23-34% → 60-80% of TC peak (#141 follow-up) #149; TC flash prefill: perf(cuda): full tensor-core (mma.sync) flash-attention prefill for d=512 — beat the half2 kernel (#141 follow-up) #146/perf(cuda): multi-warp / d-split TC flash prefill — fix the single-warp occupancy limit (#146 follow-up) #147
• llama.cpp reference: ggml/src/ggml-cuda/mmq.cu, fattn*.cu