perf(dsfp4): pick the grouped-prefill tile from route density on sm_89

[gdevenyi]

The Hopper assumption

routed_experts_fp4_prefill uses one static tile config, and its comment says why:

# One static config for every density (no autotune): the kernel is
# dequant-floor-bound, so per-expert padding at BLOCK_M=64 costs the same
# as tighter tiles while keeping the wgmma-wide M tile on sm_90.
cfg = dict(BLOCK_SIZE_M=64, BLOCK_SIZE_N=64, BLOCK_SIZE_K=64, GROUP_SIZE_M=8,
           num_warps=8, num_stages=1)

wgmma exists only on sm_90. On sm_89 nothing requires M to be 64, and once that requirement is gone the padding a 64-row tile pays stops being free. DeepSeek-V4-Flash has 256 experts and top-6 routing, so at 1536 routes each expert receives about 6 rows. A 64-row M tile is then mostly padding.

Measurements

RTX 6000 Ada, sm_89, 142 SMs, 96 MiB L2, 960 GB/s. DSV4-Flash geometry at TP=2. Measured on main.

288 configs per shape. Every config's output was checked against the sm_90 config's output before its time counted. No config was rejected on accuracy.

routes	rows per expert	shape	shipped	best on sm_89	gain
1 536	6	gate_up	4.233 ms	BM16 BN64 BK128 w4	2.376 ms, +43.9%
1 536	6	down	2.298 ms	BM16 BN128 BK128 w4	1.121 ms, +51.2%
6 144	24	gate_up	5.305 ms	BM32 BN128 BK128 w8	3.375 ms, +36.4%
6 144	24	down	2.665 ms	BM32 BN128 BK128 w4	1.444 ms, +45.8%
24 576	96	gate_up	11.098 ms	BM128 BN128 BK128 w8	6.587 ms, +40.6%
24 576	96	down	5.648 ms	BM128 BN128 BK64 w8	3.420 ms, +39.5%

The best M tile tracks the rows per expert: 16, then 32, then 128. That is the padding term, so M is now chosen from route density rather than fixed. BLOCK_K moves from 64 to 128 nearly everywhere, which halves the number of block-scale loads per output element on a kernel limited by dequantization.

What is not the mechanism

I expected num_stages=1 to be the problem, because sm_89 has no TMA and setting one stage disables the cp.async pipeline that is its only way to hide load latency. It is not. One stage wins or ties at four of the six shapes. The gain comes from tile shape. num_stages stays at 1.

Scope

sm_90 and above returns exactly the constants it returns today, so Hopper is unchanged.

The new branch covers everything below sm_90, but the density rule is measured only on Ada. If you would rather gate it to (8, 9) exactly until someone measures sm_80 or sm_86, say so and I will narrow it. I have no other pre-Hopper card to test.

Effect on a running server

On the machine this was measured on, the change is not visible end to end. That deployment runs --moe-backend offload, which puts routed experts on the CPU and leaves the GPU 13% busy, so this kernel is not the limiting factor there. A GPU-resident deployment spends a much larger share of its time in it: 57.9% of GPU kernel time in an nsys trace of 16-concurrent long-prompt serving on main.

Testing

tests/moe on main with this PR: 94 passed, 6 skipped, 1 failed. The failure is test_cpu_moe_q4_0.py::test_cpu_decode_q4_0_matches_ggml_mmvq, which also fails on main without this PR.

This PR adds no test. The config picker is a pure function of route count and expert count, and a small test that pins its three branches would be cheap. Tell me if you want it before merge.