perf(cuda): close remaining Qwen3-8B Q4_K DECODE gap to llama.cpp — non-matvec cost (prefill handled by #167; kernel-efficiency in #149/#152)

[pekkah]

Follow-up to #156/#158/#160 (PR #161). Qwen3-8B Q4_K_M on RTX 4070 Ti is now prefill 432 / decode 74.7 t/s; llama.cpp b8585 is pp1008 ~5764 / tg128 ~78–91. Two distinct gaps remain, each needing its own approach.

Decode (74.7 → ~78–91) — the lever is NOT the matvec

The Q4_K decode matvec is now ~89% of HBM peak after the #160 SoA repack (was 74%), vs Q8_0's ~90%. Per CudaDecodeMatvecQ4KRooflineProbe, that's near the realistic ceiling — squeezing the last ~1pp (the min-correction's 2 extra dp4a + 0x01010101 sum, the d8/scale float mults, or occupancy) is low-yield. Confirm with the probe before investing.

The real signal: the matvec gained +13–15% BW but e2e decode only moved +7%, so ~half the per-token cost is elsewhere. Mirror the Gemma win (#142, which turned out to be the sampler, not kernels):

• Run SHARPI_CUDA_PROFILE=1 per-phase decode breakdown on Qwen3-8B and identify the dominant non-matvec cost (attention / RoPE / norms / KV append / sampler).
• Check the sampler path specifically — Qwen3-8B greedy --temp 0 may still hit a full-vocab artifact like Gemma's bench --verbose-prompt LINQ-sort (see project_gemma4_decode_sampler_142).
• Fix the top item; argmax-stable; same-session A/B; README.

Prefill (432 → ~5764) — cp.async-pipelined MMQ

The ~13× gap is llama.cpp's mul_mat_q cp.async double-buffered weight load, which hides the per-token-tile weight re-read that our llm_mmq_q4k/llm_mmq_q4k_soa (and cuBLAS, via L2) pay. The SoA layout (#160, quants 16B-aligned) is the right substrate for cp.async.

• Prototype cp.async.cg.shared.global (sm_80+) double-buffered tile load in llm_mmq_q4k_soa; overlap the next tile's global→shared copy with the current tile's mma.
• Bit-identical / argmax-stable vs current MMQ; A/B at ~100-tok and ~1K prompts; README.

Notes

• Decode item is the higher-confidence win (Gemma precedent); prefill cp.async is larger but riskier.
• All knobs in place: SHARPI_CUDA_PROFILE, SHARPI_Q4K_SOA, SHARPI_PREFILL_MMQ, SHARPI_CUDA_GRAPH. A/B harness: scripts/bench-q4k-soa-156.ps1. Model: models/Qwen3-8B-Q4_K_M.gguf.

[pekkah]

Test-hardening follow-ups (from PR #161 review)

The PR #161 pr-test-analyzer pass surfaced coverage gaps worth closing alongside the perf work above:

• Default-on A/B test — no automated test toggles SHARPI_Q4K_SOA=0 vs unset on a real loaded Qwen3-8B and asserts argmax-equality. The maxAbs==0 kernel tests prove repack+dispatch == interleaved for synthetic weights, and the manual same-session A/B (74.7 vs 69.7 t/s) was run by hand, but the default-flip itself has no E2E gate.
• MoE-skip negative test — the !_isMoE repack gate (CudaForwardPass.cs:2646) is load-bearing (MoE Q4_K readers aren't SoA-aware) but untested. Assert an MoE Q4_K upload's handle is not marked SoA.
• cols % 256 guard test — Assert.Throws on a misaligned RepackQ4KSoa/DispatchMatVecQ4KSoa to document the contract (low value; every real hidden dim is 256-aligned).
• N2 SoA reader is not yet on a live path — llm_matvec_q4k_n2_soa is bit-identity-tested in isolation but unreachable in production (the dense-MTP CUDA pass CudaHybridGdnForwardPass doesn't repack to SoA). If/when SoA is wired into the GDN pass for dense-MTP decode/prefill perf, add an E2E CUDA byte-parity test for the N2 path (the existing MtpDecoder_GreedyParity_LlamaCpp is CPU-only and does not cover it).

CI caveat: all the CUDA SoA tests no-op without a GPU + local .gguf, so none of these gate the green check — they only run on a CUDA dev box.

[pekkah]

Worked both items (PR #163). Decode: no easy win. Prefill: shipped a long-prompt fix + redirected the core gap.

Decode (74.7 → ?) — no deletable non-matvec hotspot

SHARPI_CUDA_PROFILE=1 per-phase on Qwen3-8B (36L × 192 tok):

phase	share	what
ffn	57.8%	gate/up/down matvecs
qkv	14.4%	q/k/v matvecs
final+download	8.7%	lm_head matvec + logits
o-proj+res	8.0%	o-proj matvec
kv+attn	6.1%	KV append + SDPA
rope+qknorm	5.0%

Matvecs are ~89% of forward time; non-matvec compute is ~11%. The greedy --temp 0 sampler is a plain O(vocab) argmax — not the Gemma-#142 full-vocab-sort artifact. Decode is genuinely BW-bound on the already-SoA matvecs. The roofline note in the issue holds: pushing the matvec past ~89% is low-yield. No easy decode win — the only lever is the matvec kernel itself.

Prefill — long-prompt fallback fixed; core gap is the cold-L2 kernel

Two distinct things:

1. Shipped in perf(cuda): chunked batched prefill for >4096-token prompts (#162) #163: prompts >4096 tokens were silently dropping to the per-token loop (49.9 t/s). The 4096 cap only protects the non-flash AttentionBatched kernel; flash streams KV. Added chunked batched prefill (4096-token windows, flash attends prior KV) → 415.5 t/s (8.3×), argmax-identical, dense non-SWA. Gemma 4 SWA chunking is deferred (window-ring KV across chunk boundaries unvalidated) — that's the remaining sub-item here.

2. The cp.async-MMQ item = the same cold-L2 kernel-efficiency gap as perf(cuda): split-K + better-tiled int8 MMQ GEMM — 23-34% → 60-80% of TC peak (#141 follow-up) #149/perf(cuda): MMQ tiling rewrite — the remaining Gemma 4 prefill lever (split-K & cp.async ruled out) (#141/#149 follow-up) #152. New CudaQ4KPrefillMatmulProbe: isolated kernels do the trunk in ~0.5s, but cycling the full 5GB working set the same kernels take ~4s (100% SM, low DRAM, ~5× cold-L2 penalty). cuBLAS C1 ties the MMQ, so it's not just our kernel's tiling — it's the memory-locality regime a mul_mat_q-style cp.async/tiling rewrite addresses. Tracking under perf(cuda): split-K + better-tiled int8 MMQ GEMM — 23-34% → 60-80% of TC peak (#141 follow-up) #149/perf(cuda): MMQ tiling rewrite — the remaining Gemma 4 prefill lever (split-K & cp.async ruled out) (#141/#149 follow-up) #152; cp.async was already ruled out for Q8_0 in perf(cuda): MMQ tiling rewrite — the remaining Gemma 4 prefill lever (split-K & cp.async ruled out) (#141/#149 follow-up) #152, so this needs the full tiling rewrite, not a quick patch.

Remaining open sub-items on this issue: (a) Gemma 4 SWA chunked prefill >4096; (b) the Q4_K prefill MMQ tiling rewrite (≡ #149/#152).

[pekkah]

Prefill half: resolved by a different root cause than this issue assumed (PR #167, merged)

This issue framed the prefill gap as kernel efficiency / cp.async-pipelined MMQ. Profiling under #167 showed that was not the dominant cost. A cold-regime probe (CudaMmqColdProbe: round-robin distinct weights, working set ≫ L2) reproduced the trunk in ~0.6 s — only ~1.2× the warm probe, not the ~8× of real prefill — so the previously-theorized "cold-L2 ~5× kernel penalty" was wrong.

The real cause was mixed-quant fallback: Qwen3-8B-Q4_K_M keeps 37 tensors in Q6_K (half of ffn_down + attn_v + output), and Q6_K/Q5_K had no compute-bound prefill path — they fell to the per-token GEMM-N matvec that re-streams the whole weight once per token (~94 GB of weight reads for one N=2290 ffn_down). That memory-bound fallback was the dominant large-N prefill cost.

Fix (PR #167): dequant→fp16→cuBLAS GEMM for Q6_K/Q5_K (llm_dequant_q6k/q5k_to_f16), routed in all three CUDA forward passes (all-GPU + both hybrids; the hybrids also got Q8_0/Q4_K → MMQ), default-on, argmax-stable.

Result: Qwen3-8B prefill 432 → 2431 t/s @1008 (5.6×), matmul 4972 → 627 ms at N=2290. README updated. MTP acceptance unchanged (92%/92%).

What remains

• Prefill (~2.4× to llama.cpp's 5764): now genuinely the warm MMQ kernel efficiency (~16% of int8 TC peak) — the cp.async-pipelined / better-tiled MMQ rewrite already tracked in perf(cuda): split-K + better-tiled int8 MMQ GEMM — 23-34% → 60-80% of TC peak (#141 follow-up) #149 / perf(cuda): MMQ tiling rewrite — the remaining Gemma 4 prefill lever (split-K & cp.async ruled out) (#141/#149 follow-up) #152. Nothing new here; this issue's prefill scope folds into those.
• Decode (74.7 → ~78–91): unchanged and still open (the non-matvec-cost investigation in this issue's body). Decode path was untouched by perf(cuda): compute-bound Q6_K/Q5_K prefill closes the Qwen3-8B gap (#162) #167.

Suggest re-scoping this issue to the decode half and tracking the residual prefill lever in #149/#152. Spun off the broader-fallback guardrail and the test-coverage gap as separate issues.