perf(ops): shared Q5 K-split MMA for small-batch linear and linear_add (+11% MTP3 decode)

[giveen]

Problem and scope

Related Issue: #293

The Q5 small-batch linear paths leave throughput on the floor in two distinct
ways.

First, Q5 had no shared K-split MMA implementation, so the small-T linear_add
routes fell back to MmaResidualR64C16/R64C24 kernels that Q4 had already
moved past.

Second, the Q5 plain-linear shape dispatchers served all of T<=96 from a
single Capacity=32 K-split rung. Capacity is compile time and sizes the
accumulator fragments and kItemsPerSplit, so one wide rung makes small T stage
and accumulate columns it never fills. The cost was flat and directly visible:
at N=5120/K=6144, T=2 took the same 28.7us as T=17. Q4 has laddered its rungs
since it landed; only Q5 was left on one.

What changes as a result: MTP3 speculative decode gains 11.4% end to end.
Ordinary (non-speculative) generation moves about a percent, which is close to
this machine's run-to-run noise. This is not a general "Q5 generation is faster"
change and is not presented as one.

Implementation

Q5 K-split MMA (src/ops/linear/q5/q5_ksplit_mma.cuh, new) — shared
exact-shape kernel. Capacity is the compile-time column capacity of one tile;
MaxColumns > Capacity tiles along blockIdx.y, streaming the weights once per
tile.

Q5 linear_add — small-T routes now go through the K-split MMA
(KSplitMmaResidual) instead of the superseded MmaResidualR64C16/R64C24
kernels, with the crossover at T=60 (K=6144) and T=64 (K=17408). T=1 stays on
the split2 SIMT kernel: the K-split route leaves most of a 32-column tile idle
for a single row, and collapsing T=1 into the batched band cost ordinary decode
measurably. This mirrors what Q4 linear_add already does.

Q5 plain linear shape dispatchers — laddered to Capacity 4/8/16 before the
existing column-tiled Capacity=32 instance. No new kernels: every rung was
already being instantiated from the same template on the linear_add side, so
this is wiring.

Q4 K-split MMA (src/ops/linear/q4/q4_ksplit_mma.cuh) — double-buffered
staging pipeline.

Two boundaries were chosen by measurement rather than by copying Q4:

• A 24-rung was tried and rejected — it lost to the column-tiled
  Capacity=32 instance across T=17..24.
• T=1 moves to the 4-rung only on the two N=5120 shapes, where it beats the
  dedicated split4_c1 GEMV by 8-11%. N=6144 and N=7168 keep that GEMV, which
  is still 7-12% faster for them.

Limitations and tradeoffs:

• The large op-level win at T=2..16 does not show up end to end. This
  model's decode path does not spend enough time in those four Q5 plain-linear
  shapes for a ~27% per-call improvement to surface. It is included because it
  is strictly faster per call at no cost, not because it moves throughput.
• All numbers come from a single RTX 5090. The route boundaries are tuned
  against that one GPU and may want re-checking elsewhere.
• launch_q5_split4_c1_k6144 and launch_q5_split4_c1_k17408 no longer have
  callers. Left in place rather than deleted.

Verification

End to end, ninfer_bench, qwen3.8-27b (q5_g64_fp16), RTX 5090, fp8 KV, r=5.
Both builds compiled and run back to back in one session:

./build/bench/ninfer_bench --weights models/qwen3_8_27b.ninfer \
    -pg 2048,512 --kv-dtype fp8
./build/bench/ninfer_bench --weights models/qwen3_8_27b.ninfer \
    -pg 2048,512 --kv-dtype fp8 --spec mtp --draft-tokens 3 --lm-head-draft

	master	this branch	change
MTP3 decode	197.26 ± 0.93 tok/s	219.68 ± 0.46 tok/s	+11.4%
ordinary decode	76.12 ± 0.29 tok/s	77.20 ± 0.16 tok/s	+1.4%
MTP3 prefill	2910.89 ± 7.31 tok/s	2883.44 ± 15.71 tok/s	−0.9%

Speculative acceptance is identical on both sides (93.32%, 675 rounds, 0
fallbacks), so the MTP gain is time per round rather than a change in what the
draft head accepts. Prefill sits at the noise floor.

Op level, cold cache, comparing min_us against a build of the parent commit:

./build/bench/ninfer_linear_bench --qtype Q5 --n 5120 --k 6144 \
    --sweep 1:32:1 --repeat 100

T=17..32 is an unchanged route in both builds and is reported as a control band
for per-run drift:

shape	T=1	T=2..16	T=17..32 (control)
5120x6144	−11.3%	−27.5%	+0.0%
5120x17408	−7.8%	−26.5%	−0.0%
6144x5120	−0.4%	−29.4%	−0.1%
7168x5120	−0.2%	−33.4%	+6.8%

The 7168 row's control band drifted +6.8%, so treat its small-T figure as the
softest of the four.

Correctness:

• ctest: 119/120 pass. The one failure is ninfer_chat_templates_test, which
  fails on a missing jinja2 Python module in my environment and is unrelated
  to these changes. The qwen3_5_*_real tests skip for lack of local artifacts.
• tests/ops/linear/test_q5_a16.cpp and tests/ops/linear_add/test_q5_a16.cpp
  extended; existing route coverage already exercised the new rung boundaries
  (4/5, 8/9, 16/17) on all four shapes.
• Greedy generation is byte-identical between the ordinary and MTP3 paths, and
  output is coherent on both.