perf(cuda): batched-decode MMQ follow-ups — Q6_K MMQ kernel, Q/O split-K, cp.async staging (#201 leftover headroom: ~317 -> ~450+ t/s at N=8)

[pekkah]

Umbrella tracker for the batched-decode headroom left after #201 / PR #202. Work items are split into child issues:

• perf(cuda): Q6_K decode MMQ kernel — per-16 int8 scales vs the m16n8k32 fragment (#203 item 1, ~5 ms/step at N=8) #204 — Q6_K decode MMQ kernel (per-16 int8 scales vs the m16n8k32 fragment; ~5 ms/step)
• 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) #205 — n16 MMQ kernel efficiency: split-K / BM=32 for the [4096×*] shapes + cp.async staging (+ small K/V & attn_v shapes; ~3-4 ms/step)
• perf(cuda): evaluate SHARPI_BATCH_DECODE_MMQ default-on for batched serving (#203 item 5) #206 — evaluate SHARPI_BATCH_DECODE_MMQ default-on for batched serving (N=5-7 cells + quality soak)

Context

#201 / PR #202 landed the opt-in decode MMQ (SHARPI_BATCH_DECODE_MMQ=1): N=8 aggregate 246 -> 317 t/s. The N=8 step is now ~25.2 ms against a ~9.7 ms whole-step weight-streaming floor. Per-shape (nsys, 40-step capture, Qwen3-8B Q4_K_M @ 4070 Ti, N=8):

shape	ms/step	vs floor	path
q6k ffn_down [4096x12288]	4.77	~3.2x	WS-sw (bit-exact) -> #204
q6k lm-head [151936x4096]	2.76	~2.7x	WS-sw -> #204
q4k gate/up [12288x4096]	6.89	1.7x	MMQ n16 -> #205 (cp.async)
q4k Q/O + q4k-half ffn_down [4096x*]	4.47	~2.7x	MMQ n16, grid 64 blocks -> #205 (split-K)
q4k K/V [1024x4096]	0.91	~6.6x	WS (rows<2048 fallback) -> #205 (low prio)
q6k attn_v [1024x4096]	0.77	~16x	WS-sw -> #205 (low prio)

Combined estimate if #204 + #205 land near floor: ~450-500 t/s aggregate at N=8.

Measured dead ends recorded in #201 / PR #202 (do not retry without new evidence): token-warp WS geometry; SoA-Q8_1-act WS reads (register-bound); sub-block-sized MMQ K-step (latency-serial).

🤖 Generated with Claude Code

[pekkah]

#206 landed via PR #262 — SHARPI_BATCH_DECODE_MMQ is now default-on for multi-user continuous-batching decode (single-user spec-verify stays bit-exact via allowDecodeMmq: false).

Re-bench on Qwen3-8B Q4_K_M @ 4070 Ti confirmed the win vs the bit-exact WS path:

N	WS	MMQ (now default)	Δ
5	204.9	227.6	+11.1%
6	225.4	262.7	+16.8%
7	242.3	297.3	+22.8%
8	249.7	318.9	+27.6%

The N=5 cell (the dispatch threshold) is a solid win, not a wash — the existing N≥5 gate holds.

Remaining child levers #204 (Q6_K decode MMQ, ~5 ms/step at N=8) and #205 (n16 split-K + cp.async, ~3–4 ms/step) are validated as genuine >3–4% kernel wins and stay open. With the MMQ path now default-on they compound directly into batched-serving throughput; their per-shape estimates put N=8 aggregate at ~450–500 t/s if both land.

Updated checklist:

• perf(cuda): evaluate SHARPI_BATCH_DECODE_MMQ default-on for batched serving (#203 item 5) #206 — evaluate SHARPI_BATCH_DECODE_MMQ default-on (done: default-on for multi-user decode)
• perf(cuda): Q6_K decode MMQ kernel — per-16 int8 scales vs the m16n8k32 fragment (#203 item 1, ~5 ms/step at N=8) #204 — Q6_K decode MMQ kernel
• 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) #205 — n16 MMQ split-K + cp.async staging

[pekkah]

#205 item 1 (BM=32 retile) measured +2.5–2.9% e2e — below the 3–4% bar standalone (full data on #205). Parked on branch perf/205-decode-mmq-bm32, not merged alone. The bigger lever is #204 (Q6_K decode MMQ) — the Q6_K WS-sw shapes are ~8.3 ms/step at N=8 vs the ~2.8 ms the [4096×*] Q4_K shapes waste. Next: implement #204 and bundle the parked BM=32 so the combined PR clears the bar.

Revised checklist:

• perf(cuda): evaluate SHARPI_BATCH_DECODE_MMQ default-on for batched serving (#203 item 5) #206 — decode MMQ default-on for multi-user (shipped PR perf(cuda): default-on int8 decode MMQ for multi-user batched serving (#206) #262, +11–28%)
• perf(cuda): Q6_K decode MMQ kernel — per-16 int8 scales vs the m16n8k32 fragment (#203 item 1, ~5 ms/step at N=8) #204 — Q6_K decode MMQ kernel (the bigger lever; bundle BM=32 from 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) #205)
• 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) #205 — BM=32 done (+2.5%, parked); cp.async untested

[pekkah]

Shipped PR #263 (master 520288c) — N=8 aggregate now ~342 t/s, +6.6% over the post-#206 baseline. Two orthogonal levers stacked:

N	baseline	+BM=32	+argmax	+both
6	264.3	271.1	273.6	281.7 (+6.6%)
7	298.6	306.4	310.0	317.3 (+6.3%)
8	321.1	328.9	332.9	342.3 (+6.6%)

• GPU-side argmax tail (the bigger lever, +3.5–3.8% alone): BatchForwardMultiArgmax reuses the perf(cuda): GPU-side greedy argmax — every decode token downloads the full-vocab logits (0.6–1 MB) with a blocking stream sync #219 ArgmaxRows kernel to elide the full N×vocab logits D2H (4.86 MB/step at N=8) when every active seq is greedy. This was NOT in the original perf(cuda): batched-decode MMQ follow-ups — Q6_K MMQ kernel, Q/O split-K, cp.async staging (#201 leftover headroom: ~317 -> ~450+ t/s at N=8) #203 item list — it came out of the e2e time-budget audit, and it beats every MMQ-tile lever on confidence×delta/effort.
• BM=32 tile (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) #205 item 1, +2.4–2.6% alone): grid-starvation fix for the [4096×*] shapes.

Closed #205 (item 1 shipped bundled; item 2 cp.async dropped — adversarial panel + the #176/#152 prefill precedent rate it ~0–1%/e2e-neutral on decode).

Remaining: #204 (Q6_K decode MMQ) is the last big chunk (~8.3 ms/step of WS-sw at N=8). It's now a stack on top of this (+3–5% est.), large-effort/low-confidence (bespoke per-16-scale mma + Q6_K SoA repack). Updated checklist:

• perf(cuda): evaluate SHARPI_BATCH_DECODE_MMQ default-on for batched serving (#203 item 5) #206 — decode MMQ default-on for multi-user (PR perf(cuda): default-on int8 decode MMQ for multi-user batched serving (#206) #262, +11–28%)
• 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) #205 item 1 (BM=32) + GPU-argmax tail (PR perf(cuda): GPU-side argmax for batched decode + BM=32 MMQ tile (#205, #203) #263, +6.6%)
• perf(cuda): Q6_K decode MMQ kernel — per-16 int8 scales vs the m16n8k32 fragment (#203 item 1, ~5 ms/step at N=8) #204 — Q6_K decode MMQ kernel (remaining; the cheap dequant→cuBLAS probe is a known regression, do NOT try it — bespoke kernel only)

[pekkah]

#204 built + validated bit-exact, but DEFERRED on VRAM (full detail on #204). The Q6_K decode-MMQ kernel is correct and a real +13% win at N=6, but the SoA companion (the AoS Q6_K weight can't be freed — single-token/GEMM-N/WS readers still need it) ~doubles the eligible Q6_K weights (~1.6 GB) and tips the 12 GB card into WDDM spill at N≥7 → −48% at N=8. Not merged (worktree branch kept as reference). VRAM-neutral fix = port all Q6_K readers to SoA so the AoS can be freed — larger follow-up.

This exhausts the umbrella's actionable >3-4% levers under the 12 GB constraint. Shipped: #206 (PR #262, +11-28%), #205-item1 BM=32 + GPU-argmax tail (PR #263, +6.6%). Dropped/deferred with evidence: #205-item2 cp.async (e2e-neutral on decode), #204 (VRAM-prohibitive companion; needs the reader-port follow-up). Closing the loop here pending the Q6_K-readers-to-SoA effort.

[pekkah]

Umbrella complete — N=8 batched decode went from ~250 → 382 t/s this arc (+53% cumulative on Qwen3-8B Q4_K_M @ 4070 Ti):

• perf(cuda): evaluate SHARPI_BATCH_DECODE_MMQ default-on for batched serving (#203 item 5) #206 (PR perf(cuda): default-on int8 decode MMQ for multi-user batched serving (#206) #262): decode MMQ default-on for multi-user → ~250 → ~320 (+11–28% N=5–8).
• 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) #205 item 1 (BM=32) + GPU-argmax batched tail (PR perf(cuda): GPU-side argmax for batched decode + BM=32 MMQ tile (#205, #203) #263): ~320 → ~342 (+6.6%).
• perf(cuda): Q6_K decode MMQ kernel — per-16 int8 scales vs the m16n8k32 fragment (#203 item 1, ~5 ms/step at N=8) #204 Q6_K decode MMQ via SoA reader-port (PR perf(cuda): Q6_K decode MMQ via SoA reader-port (free AoS) — +17% N=8 (#204) #264): ~342 → 382 (+16.8% over its baseline), VRAM 11394/12282 (no spill). The Q6_K trunk (lm-head + ffn_down-half) now runs the int8 TC tile; all Q6_K readers ported to SoA + AoS freed; bit-exact (incl 48-token single-user greedy parity).

Dropped/not-pursued (validated below the >3–4% bar): #205 item 2 cp.async (e2e-neutral on decode, per the #176/#152 prefill precedent); the Q6_K companion + AoS-direct approaches (companion spilled 12 GB at N=8; AoS-direct −7…−11% from uncoalesced staging — the reader-port was the answer).

The original ~450–500 t/s target was optimistic; at N=8 the step is increasingly per-seq-attention-bound (#194), not Q6_K/Q4_K weight-matmul-bound, so the remaining gap has no validated >3–4% kernel lever. Closing the umbrella — the batched-decode MMQ arc is done.