perf(cuda): n16 decode-MMQ kernel efficiency — split-K for [4096×*] shapes + cp.async staging (#203 items 2-3, ~3-4 ms/step at N=8)

[pekkah]

Child of #203. Two efficiency gaps inside llm_mmq_q4k_soa_acts_n16 (PR #202), measured on Qwen3-8B Q4_K_M @ 4070 Ti, N=8:

1. Grid starvation on the [4096 x cols] shapes (~2 ms/step): Q/O + the Q4_K ffn_down half run at BM=64 -> grid = 64 blocks on 60 SMs (avg 49.7 us vs 18.7 us floor for Q/O). Split the K-range across 2-4 blocks per row-strip with a deterministic second-pass reduce (split-K scratch + tiny reduce kernel — NO fp atomics; run-to-run determinism is load-bearing for the test oracles). Alternative: BM=32 tile (128 blocks, 4 warps) — simpler, no reduce pass, weight still read once.

2. cp.async staging (~1-1.5 ms/step): gate/up sits at 1.7x floor (95.7 us vs 56) and the K-loop is staging-latency-limited — the super-block K-step (8 mma/barrier pair) was the difference between losing to the WS matvec and +28% aggregate, but the global->register->shared round-trip still bounds the loop. The Track B helpers (sharpi_cp_async16/commit/wait) and a proven double-buffered kernel shape (llm_mmq_q8_0_soa_acts_cpa) already exist; the SoA quants are 16-B-aligned and cp.async-eligible. Note the Track B precedent: it was kernel-level +10-15% but e2e-NEUTRAL on prefill — decode N=8 is a different regime (the kernel IS 49% of the step), so measure e2e before defaulting.

Also fold in (low priority, ~1.7 ms combined): K/V [1024x4096] Q4_K (6.6x floor, WS fallback below the rows>=2048 threshold) and q6k attn_v (16x floor, 0.77 ms) — a BM=32 tile or a fused QKV matmul would cover them.

Expected combined: N=8 ~317 -> ~380-400 t/s with #204 still pending.

🤖 Generated with Claude Code

[pekkah]

Validation — item 1 (BM=32 retile) implemented and measured: +2.5–2.9% e2e, below the 3–4% bar on its own.

Built the grid-starvation fix as the issue's "simpler, no-reduce-pass" BM=32 tile: emit llm_mmq_q4k_soa_acts_n16 at a second row-tile size (BM=32, 128 threads, 4 warps) from one template, and route shapes with ceil(rows/64) < 2·SM (the [4096×cols] Q/O proj + Q4_K ffn_down half — 64 blocks on a 60-SM 4070 Ti) there; high-row shapes (gate/up 12288, lm-head 151936) keep BM=64. Output is bit-identical (same fragments/mma/accumulation; only the block→row mapping changes) — validated via CudaDecodeMmqTests (both tiles vs the fp32 CPU reference) and the N=6 MMQ oracle.

A/B on Qwen3-8B Q4_K_M @ 4070 Ti (2 runs each, default-on MMQ, BM=32 routing off vs on):

N	BM=64	BM=32	Δ
5	228.1	233.6	+2.4%
6	264.2	271.2	+2.6%
7	297–298	305–307	+2.6–2.9%
8	320.0	328.0	+2.5%

The nsys per-shape estimate (~2 ms/step ≈ ~8%) did not translate e2e. The [4096×*] Q4_K shapes overlap with larger concurrent work and aren't as dominant in the real step as the isolated-kernel profile implied; the Q6_K WS-sw shapes (#204) are the bigger chunk (ffn_down half 4.77 + lm-head 2.76 + attn_v 0.77 ≈ 8.3 ms/step).

Decision (per the >3–4% threshold): BM=32 is correct and bit-identical but doesn't clear the bar alone, so it is not being merged standalone. It's parked on branch perf/205-decode-mmq-bm32 (kill-switch SHARPI_DECODE_MMQ_BM32=0) to bundle with #204 (Q6_K decode MMQ) — together they clear the bar decisively (umbrella #203 target ~450–500 t/s). Item 2 (cp.async) remains untested on decode; given the prefill precedent was e2e-neutral (#176/#152), #204 is the higher-confidence lever to pair with.