SM103 serving hang caused by TRTLLM_GEN_BMM artifact regeneration in #3708 (batched_gemm-dd6d23e-721ae60, 0.6.14): FP4 batched-GEMM clusters stuck in mbarrier phase wait

## Summary During sustained serving of Qwen3.5-397B-A17B-NVFP4 (MoE, TP4, `trtllm_fp4_block_scale_moe` via sglang) on GB300 (SM103), a TRTLLM_GEN_BMM kernel from the flashinfer-cubin 0.6.14 artifact intermittently never completes: a few thread-block clusters remain stuck in an mbarrier phase-wait loop after the rest of the grid has finished. The affected rank never reaches the next TP allreduce, peer ranks wait indefinitely in `ncclDevKernel_AllReduce`, and the NCCL watchdog eventually aborts the serving group. **We attribute this regression to #3708 — specifically to the regenerated TRTLLM_GEN_BMM artifact it ships (`batched_gemm-dd6d23e-721ae60`), not to its host-side code changes.** The isolation is three-layered: 1. **Version window**: 0.6.13 (artifact `896b90b`) does not reproduce; 0.6.14 (artifact `dd6d23e-721ae60`) reproduces in 9 of 10 runs. The only changes to this artifact and its launch path in that window are #3708 and #3588. 2. **#3588 (PDL wiring) excluded**: forcing `device_support_pdl` to False (verified at runtime) still hangs. 3. **#3708 split-tested**: keeping flashinfer 0.6.14 fully intact — including #3708's host-side selection changes in `csrc/trtllm_batched_gemm_runner.cu` / `KernelRunner.h` — and substituting only the artifact contents with the previous generation (`batched_gemm-91e0ba0-da44fdf`, forcing a runtime JIT rebuild against its metainfo) eliminates the hang. The host half of #3708 is therefore exonerated; the artifact half is the trigger. GPU core dumps of all four TP ranks, captured while wedged, are available. ## Environment - GB300 (SM103), CUDA 13.0, torch 2.11.0+cu130 - sglang, disaggregated prefill, TP4, `moe-runner-backend: flashinfer_trtllm` - flashinfer trio 0.6.14 (`flashinfer_python` / `flashinfer-cubin` / `flashinfer-jit-cache`); also reproduces with the wheels shipped in public sglang cu13 nightly images (20260711 and later) - Workload: multi-turn 256k-context serving; token-per-expert counts vary per batch (dynamic-batch BMM variants) ## Evidence (from the core dumps) Kernels resident on each GPU at capture time (3+ minutes after the last completed batch; all scheduler threads blocked on their next stream sync): | rank | resident kernel | state | |---|---|---| | 0, 1, 3 | `ncclDevKernel_AllReduce_Sum_bf16_RING_LL`, grid (32,1,1)×(544,1,1) | waiting for peer | | 2 | `bmm_Bfloat16_E2m1E2m1_Fp32_Ab16_Bb16tokFp32_t128x128x256_s6_et128x128_m256x128x64_c2x1x1_rM_TN_transOut_schPd2x1x2x3_biasFp32M_bN_rgTma_clmp_dynB_sm100f`, grid (32,1631,1)×(384,1,1), cluster dim (2,1,1) | stuck | Residual state of the stuck kernel: 10 of 52,192 blocks (5 cluster pairs) still resident, the rest completed: ``` BlockIdx To BlockIdx Count State * (2,0,0) (3,0,0) 2 running (0,1,0) (1,1,0) 2 running (20,1,0) (21,1,0) 2 running (16,2,0) (17,2,0) 2 running (24,3,0) (25,3,0) 2 running ``` All warps of the inspected block are in the same wait loop (SASS at `$pc`): ``` => @!P0 NANOSLEEP.SYNCS 0x989680 @!P0 SYNCS.PHASECHK.TRANS64 P0[R5+URZ+0x33890],R3 @!P0 BRA <back to the poll> ``` The mbarrier phase never advances. With the NCCL watchdog enabled this surfaces as `Watchdog caught collective operation timeout: ... OpType=ALLREDUCE ... 600000ms` on the waiting ranks. Onset is intermittent: minutes to about 45 minutes of sustained serving (observed across roughly 30 occurrences; about one run in eight completes without hanging). ## Isolation experiments Single-variable substitutions; same image, config, and hardware for every row. "Hang rate" counts full sustained-serving runs (about 4,300 prefill batches each when they complete). | flashinfer trio | TRTLLM_GEN_BMM artifact in use | PDL | outcome | |---|---|---|---| | 0.6.12 | `batched_gemm-91e0ba0-da44fdf` | default | 0 hangs in 20+ runs | | 0.6.13 | `batched_gemm-91e0ba0-896b90b` | default | 0 hangs in 2 runs | | 0.6.14 | `batched_gemm-dd6d23e-721ae60` | default (on) | hang in 9 of 10 runs | | 0.6.14 | `batched_gemm-dd6d23e-721ae60` | forced off (`device_support_pdl` patched to False, verified at runtime) | hang in 1 of 2 runs — **PDL is not the trigger** | | 0.6.14 | **`batched_gemm-91e0ba0-da44fdf` (substituted)** | default (on) | **0 hangs in 2 runs** | Notes on the substitution row, since it is the load-bearing one: - The prebuilt `fused_moe` modules in `flashinfer-jit-cache` (which carry a build-time artifact path) were removed, forcing a runtime JIT rebuild; the new artifact directory was deleted from `flashinfer-cubin` so any residual reference would fail loudly rather than silently fall back. - Verified post-hoc: the rebuilt module's compile flags point at the old artifact's include directory, its binary contains none of the 504 kernel names that exist only in the new metainfo, and no runtime re-download of the new artifact occurred. - The stuck kernel's name exists in both artifact generations with different cubin SHA-256, so the substituted runs exercise the same-named kernel compiled from the older generation. ## Root-cause attribution **The regression entered with #3708's regeneration of the TRTLLM_GEN_BMM artifact (`91e0ba0-896b90b` → `dd6d23e-721ae60`), shipped in 0.6.14.** Since the artifact is a pregenerated cubin package, the defective change itself lives in the kernel-generator side that produced `dd6d23e`; within the flashinfer repo, #3708 is the terminal attribution. Supporting chronology: - #3153 (May) bumped the artifact `91e0ba0-da44fdf` → `91e0ba0-896b90b`. That bump was reverted on `release-v0.6.12` (15886abd) but kept on main, so 0.6.13 ships `896b90b` — and 0.6.13 does **not** reproduce this hang. The `896b90b` generation is healthy; the break is specific to the `dd6d23e-721ae60` regeneration. - #3708's host-side bias config selection changes (`csrc/trtllm_batched_gemm_runner.cu`, `KernelRunner.h`) are exonerated by the substitution experiment, which ran them unchanged against the older artifact with no hangs. ## Reproduction Runnable recipes for this serving setup are in NVIDIA/InferenceMAX#117 (`benchmarks/multi_node/srt-slurm-recipes/sglang/qwen3.5/gb300-fp4/agentic/`). The hang reproduces on the disaggregated TP4↔TP4 agentic points (e.g. `disagg-gb300-1p1d-tp4-tp4-c64-mtp-convaware-jid2214439.yaml`, 2×GB300 nodes) with flashinfer 0.6.14 obtained either way: - run the recipe image `lmsysorg/sglang:nightly-dev-cu13-20260709-074bb928` and pip-install the trio at 0.6.14 (`flashinfer_python==0.6.14` from PyPI, `flashinfer-cubin==0.6.14` from https://flashinfer.ai/whl, `flashinfer-jit-cache==0.6.14` from https://flashinfer.ai/whl/cu130), or - use any 20260711-or-later cu13 nightly image (ships 0.6.14 natively). Dataset is public (`semianalysisai/cc-traces-weka-062126-256k` on HF). Keep `--mamba-scheduler-strategy extra_buffer` as in the recipes. Expected within minutes to about 45 minutes of sustained load: prefill batch lines stop, then the NCCL watchdog fires ALLREDUCE timeouts on the peer ranks. The 0.6.12 or 0.6.13 trios, aggregated (non-disaggregated) serving, or `no_buffer` do not reproduce it. ## Suggested change - **Short term**: revert the TRTLLM_GEN_BMM artifact pointer to the previous generation (`91e0ba0-896b90b`, as shipped in 0.6.13) on main and the next release, following the precedent of the `release-v0.6.12` revert (15886abd). Our isolation shows both the `da44fdf` and `896b90b` generations run clean under the exact load that wedges `dd6d23e-721ae60`, and no host-side change is needed. - **Root fix**: locate the defect in the generator that produced `dd6d23e-721ae60` — the failure mode is a cluster-pair mbarrier wait that never satisfies (TMA transaction/arrival accounting) in the dynamic-batch bias variants — and regenerate. We can turn around validation of candidate artifacts on the serving reproducer quickly; core dumps and cuda-gdb transcripts are available on request.