perf(cuda/quant): block output rows per warp in qmv below Ampere

[inureyes]

Follows #1678, which has merged, so this is six commits against main touching qmv.cu and the environment variable reference. It replaces #1680, which GitHub closed when #1678's branch was deleted out from under it.

What lands

Two changes to qmv, both on the pre-Ampere decode path, both with an environment override, and both byte-identical to the stock kernel at --temp 0.

Row blocking. qmv gives each warp one output row and has that warp read the whole activation vector, so a launch moves n * k * 2 bytes of activations against n * k / 2 bytes of 4-bit weights: four bytes of activation for every byte of weight. Giving each warp R consecutive output rows and loading the activation tile once for all of them divides that term by R and leaves weight traffic alone. The per-row arithmetic is untouched, which is what keeps the result bit-identical.

A pinned resident-block floor. The R sweep turned out to be a curve in occupancy rather than in loads, and it tracks the register column exactly: R = 1 at 40-48 registers and 4-5 blocks per SM gives 44.7 tok/s, R = 2 at 64-72 and 3-4 gives 56.0, R = 3 at 80-110 and 2-3 gives 53.4, R = 4 at 96-142 and 1-2 gives 45.2, R = 6 at 128 and 2 gives 42.7. Transactions per output row fall at every step and past R = 2 every step also costs a resident block, and the occupancy wins each time. So __launch_bounds__(256, min_blocks) hands ptxas the register budget directly instead of trying to reach it through the row count.

Measured on a V100 at 4 bits, decode as the slope between -n 60 and -n 200, all defaults in place, generated text byte-identical in every arm:

arm	gemma-4-12B-it-4bit	qwen3.8-27B-4bit
default (R = 2, floor 3)	58.3 tok/s	26.3 tok/s
R = 3, floor 3	57.1	26.7
floor 1 (no launch bounds)	54.7	26.3
R = 1 (stock kernel)	45.3	20.8

Defaults are R = 2 with a floor of 3, on sm_70 only, read from the compute capability the encoder already carries. Turing stays at R = 1 with no floor: it has a different register file per SM, no sm_75 device was available to measure, and the R = 4 reversal is a direct demonstration that this tradeoff does not extrapolate. MLXCEL_QMV_ROWS_PER_WARP=1 is the rollback.

What is here as an instrument

MLXCEL_QMV_PROBE_NO_SCALE reads the scale and the bias from group 0 rather than the group the step is in, so the output is wrong by construction. It exists because two of the four loads a step issues per output row are the scale and the bias, four bytes of metadata, and interleaving them into a single 32-bit word at load time would have cost a device-side packing cache, about 11% more VRAM, and an extension to the quantized.cpp overlay. The probe prices the loads before any of that: removing both is worth 7.3% on the 12B and 4.4% on the 27B, so the interleave that removes one is not being built.

The probe stays because ncu cannot run here. The box is a Docker container whose root holds no CAP_SYS_ADMIN, /proc/driver/nvidia/params is read only, and it reports RmProfilingAdminOnly: 1, so Nsight Compute returns ERR_NVGPUCTRPERM even under sudo. Nsight Systems works, but only with --trace=cuda: the default trace set breaks cuModuleLoadDataEx on MLX's JIT PTX with an empty JIT log, which looks like a corrupt cache and is not one.

What was tried and rejected

Each of these is worth a line because the next person to read qmv.cu will have the same ideas.

• __restrict__ on the row-blocked kernel. No change to register counts, no change to throughput.
• Batching the scale and bias through a warp shuffle. Four consecutive steps span exactly 32 groups, one per lane, so a single coalesced load plus __shfl_sync can serve all four and turn 8 transactions per four steps into 2. It lost everywhere: 60.1 to 51.5 at R = 2 and 58.8 to 46.7 at R = 3, at identical register counts and identical residency. The shuffles and the pack and unpack cost more than the transactions they save.
• Staging the activation tile in shared memory. Every warp in a block walks the same indices, so a block issues the same 1024 bytes per step eight times; staging four steps once per block turns that into one cooperative load plus an LDS per consumer, with no added arithmetic. Four interleaved repeats put it at +3.9% on the 27B (t = 2.06, and it more than doubled the run-to-run spread) and +1.2% on the 12B (t = 0.46), against doubling the instantiation count and taking the build from 23 minutes to 50.
• Widening the rms_norm block when the grid is one row. Landed and reverted in this branch on purpose, so the attempt and its numbers stay in the history: see the revert commit. rms_norm_small is 7.7% of the decode window at 336 launches per token, and upstream only ever gives it 64 or 128 threads, which is four warps on one SM of eighty. Going to 640 threads on the 27B with the aligned fast path intact moves throughput 3.4% at t = 1.94, and the 12B 1.3% at t = 0.52.

The staging result is the useful one: cutting the activation term's global transactions eightfold for almost nothing is what established that this kernel is bound by the load instructions its warps issue rather than by the bytes they move.

Not covered

sm_75 is untested, by choice, and defaults to the stock path.

Decode targets are not met. The 12B is at 60 tok/s against a 110 target and the 27B at 27 against 50. qmv is 68.6% of the decode window, so halving it yields 34% of the 45% the 12B target needs; the targets require qmv roughly three times faster. One structural idea remains unexplored: rewriting the dot product as sum_g [ scale * sum(x * w_int) + bias * sum(x) ] accumulates the raw integer product per group and applies the scale and bias once per group, which turns the per-row accumulator from sixteen elements into one scalar and reopens the register budget that pins R at 2. It changes the summation order, so the byte-identity check that validated everything in this PR stops applying and a perplexity gate has to replace it.

Refs #1543