perf(cuda): evaluate SHARPI_BATCH_DECODE_MMQ default-on for batched serving (#203 item 5)

[pekkah]

Child of #203. The #201 decode MMQ ships opt-in (SHARPI_BATCH_DECODE_MMQ=1): N=8 aggregate 248 -> 317 t/s, argmax-stable (same contract as the prefill MMQ and the SHARPI_BATCH_DECODE_GEMM toggle), with per-tensor fallback to the bit-exact WS path for Q6_K / rows<2048 / N<5.

Before flipping the default for batched serving (SHARPI_MAX_BATCH / ContinuousBatchingEngine users):

1. N=5-7 cells: the dispatch threshold N>=5 was set from N=2 (123 vs 139, loses), N=4 (wash), N=8 (+28%); the 5-7 range is interpolated, not measured. Bench those cells (SHARPI_BENCH_BATCH_N=5,6,7) and adjust the threshold if 5/6 are washes.
2. Quality soak: argmax-stable is a kernel-level contract; run a multi-sequence generation comparison (greedy, ~100+ tokens per seq, several prompts) MMQ=1 vs default and confirm token-level agreement or benign divergence only.
3. Decision: default-on for BatchForwardMulti only (single-user Forward is untouched either way), keeping SHARPI_BATCH_DECODE_MMQ=0 as the bit-exact kill-switch — or keep opt-in and document the knob in the server/CLI docs.
4. Re-evaluate after 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/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 land — if Q6_K joins the MMQ path the default-on win grows from +28% to an estimated +60-80% at N=8.

🤖 Generated with Claude Code