[Fridge003]

Motivation

test/registered/e2e/models/test_kimi_k3_b300_low_latency.py::TestKimiK3B300LowLatency::test_bs_1_speed (NV Extra, extra-b-test-8-gpu-b300) fails deterministically on PRs tested against current main: the accept length is 3.4594594594594597 against the 4.0 threshold, on both retries and on every PR.

PR runs of this test, by the main commit CI merged each PR onto:

main base	result	PR
b147f4aa10	pass	deps/transformers-5-21
1305b1830a	3.4595	codex/pp-arrival-match-main
89b2641381	3.4595	fix-mamba-copy-from-temporal-oom, feat/hicache-mamba-request-end, fix/hicache-prefetch-swa-match-boundary
3fa242f27c	3.4595	deps/transformers-5-21, codex/pp-arrival-match-main
c03b6b3df5	3.4595	fix-mamba-copy-from-temporal-oom
06c60f412a	3.4595	#48715

The only PR that still passes on these bases is #48468 (FlashInfer 0.7.3), which changes the MXFP4 MoE runner Kimi-K3 uses and so changes its numerics anyway.

Of the main commits in b147f4aa10..1305b1830a, only #47683 reaches Kimi-K3's TP8 + DSpark decode path:

• [Fix] KDA: Compile the CuTe DSL decode kernel for the caller's dt_bias dtype #47866 (CuTe DSL KDA decode) only runs with --linear-attn-decode-backend cutedsl; this test uses triton.
• [Perf] [NVIDIA] dflash: Fill decode-prep length buffers in bulk #47678 (DFlash decode prep) fills the same buffers.
• [Perf] [NVIDIA] qwen3: Fuse prefill QK norm and RoPE with contiguous Q #47785 (Qwen3 QK norm) and [Feat] dsa: Exact multicast top-k merge for DCP decode #48588 (DSA DCP top-k) do not reach Kimi-K3.
• The rest are docs, NPU, AMD, CI and diffusion.

#47683 makes the tiny GEMM switch from the SIMT to the MMA kernel at a batch size picked per shape. Each kernel alone is batch invariant, but a row's result depends on the batch through that switch. Kimi-K3 runs its KDA gate and router projections through kimi_k3_tiny_gemm:

(N, K)	MMA from m =	bs=1 decode (m = 1)	8-token verify (m = 8)
144 × 7168	2 (cost model, B300)	SIMT	MMA
896 × 7168	2 (cost model, B300)	SIMT	MMA
1536 × 128	8 (tuned)	SIMT	MMA

So the same token's gates round differently in decode and in speculative verify.

Modifications

Pin the three Kimi-K3 shapes to SIMT at every m in _TUNED_MMA_MIN_M, as 128x512 already is. Other tiny GEMM callers (DeepSeek routers, DeepSeek-V4.1, beam, dots3) use other shapes and are unchanged.

Validation

Draft until confirmed: I had no free B300 to bisect on, so the root cause above comes from elimination plus the dispatch rule. /rerun-test of the failing test on this PR decides it.

• If it passes, this goes ready for review.
• If it fails, the next step is to check [Perf] [NVIDIA] gemm: Tiny GEMM as SIMT + mma.sync templates with batch-size dispatch #47683 directly.

🤖 Generated with Claude Code

---

CI States

Latest PR Test (Base): ❌ Run #42121183518
Latest PR Test (Extra): ❌ Run #42121183956
Latest PR Test (AMD ROCm 10): ❌ Run #42121183357

Tip

Base CI only covers the core cases, and its workload may not exercise your change. Use /rerun-test to run the tests your change touches, or copy this prompt for an agent:

Find the registered tests this PR affects and run them with /rerun-test, following https://github.com/sgl-project/sglang/blob/main/.agents/skills/rerun-test/SKILL.md

[Fridge003]

/rerun-test registered/e2e/models/test_kimi_k3_b300_low_latency.py::TestKimiK3B300LowLatency.test_bs_1_speed registered/kernels/ops/gemm/test_tiny_gemm.py

[github-actions]

🚀 8-gpu-b300 (1 test): ✅ View workflow run

cd test/ && python3 registered/e2e/models/test_kimi_k3_b300_low_latency.py TestKimiK3B300LowLatency.test_bs_1_speed

🚀 1-gpu-h100 (1 test): ✅ View workflow run

cd test/ && python3 registered/kernels/ops/gemm/test_tiny_gemm.py

[claude]

Claude Code Review

This repository is configured for manual code reviews. Comment @claude review for a one-time review, or @claude review always to subscribe this PR to a review on every future push.

_{Tip: disable this comment in your organization's Code Review settings.}