nvfp4 experts cuda: decode kernels stall on load latency; a cp.async ring makes GLM-5.3 decode up to 24% faster, same bits #2265

Component: CUDA NVFP4 expert decode kernels (zig/kernels/cuda/nvfp4_experts.cu: nvfp4_expert_kernel, nvfp4_expert_waves, nvfp4_expert_half). GLM-5.3-Flash FP4-Local, Qwen3.6 ModelOpt NVFP4, Nemotron and Flash Next all reach them.

Metric: decode tokens/s.

What ncu shows: in GLM-5.3-Flash FP4-Local decode (RTX PRO 6000 Blackwell, TP=2), the gate-up expert kernel is the largest single cost, about 3.75 ms of the 11.8 ms each token takes. Each warp spends about 110k cycles stalled on long scoreboard, yet DRAM is only about 30% busy. The register double buffer keeps roughly one k32 group's loads in flight, and the activation loads hit L1 only 4% of the time.

Proposal: each warp keeps a ring of D = 4 shared-memory stages, filled by cp.async with its weight words, block scales and both rows' inputs. Each group is summed in the same k order as today. Every output byte stays identical, so each kernel is a drop-in under its existing contract registration.

Measured on a branch off main (bit-identical; kernel contract green on sm_120, sm_89 and sm_86):

• GLM-5.3-Flash FP4-Local, 2x RTX PRO 6000 Blackwell, TP=2: plain decode 83.8 -> 104.1 tok/s (+24%). Drafted: code 212.7 -> 226.5, chat 163.2 -> 180.3. token_sha is identical to main, and drafted == plain.
• Qwen3.6-35B-A3B-NVFP4 on one RTX 6000 Ada: decode +4.7% to +8.5%; 4 streams +5.7% to +6.5%.
• Qwen3.6-35B-A3B-NVFP4 on one RTX A6000: decode +5.6% to +7.7%; 4 streams +5.2% to +5.6%.
• token_sha is identical between base and candidate on every prompt, and concurrent == alone == serial.

Related:

• glm cuda: TF_GLM_L2PF reads a decode window's next weights into L2 on a side stream, on by default on a GB10 #2162 (TF_GLM_L2PF) prefetches the next window's weights into L2. This is complementary: it hides latency inside a kernel.
• glm cuda: windows to 32 rows, priced by rounds timed at load #2257 moves GLM decode gate-up onto the factored tiles. This change still covers GLM's decode until then, and afterwards its down projection (experts_waves) and every other NVFP4 model's decode.