[MatthewBonanni]

Purpose

The ll_bf16 split-K router GEMM (#42562, M >= 8) stores partials into peer CTAs' shared memory (st.shared::cluster) before any cluster barrier. If a peer CTA hasn't started yet, the GPU raises Xid 13 CTA Not Present, which surfaces as CUDA error: unspecified launch failure.

Fix: all threads cluster_arrive_relaxed() after pipeline init. DMA warps cluster_wait() after committing stage 0, and MMA warps wait at the start of their branch. So every peer is resident before the epilogue's st.shared::cluster, and the wait stays out of the PDL prologue (split suggested by @LopezCastroRoberto).

Reproduction

This uses @Skyseaee's single-GPU harness (comment). It runs the router GEMM in a CUDA graph, concurrent with non-cluster copy kernels that write to mapped pinned host memory:

nvcc -O3 -shared -Xcompiler -fPIC -arch=sm_103a copy.cu -o libhostcopy.so
python -u router_probe.py --tag <tag> --mode direct --seconds 120

B300, fresh process per run	Result
main (a12e122206)	unspecified launch failure in the first batch, 4/4 runs
this PR (ac5c06ee4b)	120 s clean, 4,912,896 router calls, outputs correct

The harness does not trigger the fault on H200. In serving, the same Xid hit GLM-5.3 HiSparse P/D decode (2x 8xH200, TP4) within minutes. The first revision of this fix (arrive+wait at kernel start) ran 6/6 passes clean there; the current revision has not been run in serving.

No regression test: the race depends on hardware CTA scheduling, so it can't be made deterministic. test_ll_bf16_gemm.py covers numerics only.

Test results

All on ac5c06ee4b.

pytest tests/kernels/test_ll_bf16_gemm.py: 144 passed on B300 and on H200.

Latency is µs per call: the median of 30 replays of a CUDA graph holding 200 back-to-back PDL ll_bf16_gemm calls, N=256.

GPU, K	M=8 / 12 / 16, main	M=8 / 12 / 16, this PR
B300, 6144	2.45 / 2.74 / 2.71	2.47 / 2.81 / 2.77
B300, 7168	2.78 / 2.79 / 2.76	2.80 / 2.80 / 2.76
B300, 4096	2.59 / 2.46 / 2.47	2.61 / 2.47 / 2.47
H200, 6144	2.62 / 2.68 / 2.72	2.66 / 2.72 / 2.75
H200, 7168	2.79 / 2.78 / 2.85	2.84 / 2.82 / 2.88
H200, 4096	2.40 / 2.61 / 3.00	2.44 / 2.66 / 2.94

🤖 Generated with Claude Code

[LopezCastroRoberto]

Perf verified on B300. Cost is ~0.06 µs/call on this setup with with PDL enabled.

nit: splitting the barrier so it stays out of the PDL prologue seems to recover about half of that (+0.035 vs +0.063 µs vs main), and all of it at:

cute.arch.cluster_arrive_relaxed()   # keep; drop the cluster_wait() here
...
    mainloop_pipeline.producer_commit(producer_state)   # first stage
    cute.arch.cluster_wait()  # after B prefetch + griddepcontrol_wait
...
else:
    cute.arch.cluster_wait()  # MMA warps idle until stage 0 lands anyway
    cute.arch.setmaxregister_increase(232)

Note: I couldn't reproduce the bug you report. Can you check this alternative also fixes it?

[Skyseaee]

@MatthewBonanni We have an independent single-GPU reproduction that may help with validation here. It reproduces without model weights, a connector SDK, RDMA, or a P/D deployment.

The trigger is running the ll_bf16 split-K router concurrently with ordinary CUDA kernels that write directly to mapped pinned host memory on another stream. The copy kernels do not use clusters, and their buffers are separate from the router's buffers.

Environment:

• B300 / SM103, driver 580.159.04
• vLLM 0.30.1.dev15+ga6bb06c0f, torch 2.13.0+cu130
• CUDA toolkit 13.0.88, CUTLASS DSL 4.7.1, cuda-bindings 13.4.3
• PDL enabled; CUDA_LAUNCH_BLOCKING unset

The unpatched _ll_bf16_splitk.py is byte-identical to upstream main at fa75d34 (SHA256 f282d0c1c4649d3affb19d694b7edf66cecf7764946b4effe57737830ec7a3ad). The rest of the environment is not a pristine-main build.

Reproduction procedure (sources and logs are in the attached ZIP):

1. Use BF16 A[8,7168] and W[896,7168], FP32 output; dispatch selects split_k=6, num_stages=4.
2. Capture 256 ll_bf16_gemm calls in one CUDA graph.
3. Allocate a GPU source of shape [27,4096,576] bytes and a pinned-host destination of shape [4096,27,576]. Obtain the destination's device pointer with cudaHostGetDevicePointer.
4. On a separate stream, enqueue four groups of 27 non-cluster uint4 copy kernels (one launch per layer, 4096 blocks x 128 threads), each group writing 60.75 MiB directly to mapped host memory. Replay the router graph on the main stream, then synchronize both streams and repeat.

From the extracted directory, in an existing vLLM environment:

mkdir -p results
nvcc -O3 -shared -Xcompiler -fPIC -arch=sm_103 copy.cu -o libhostcopy.so
unset CUDA_LAUNCH_BLOCKING
export CUTE_DSL_DISABLE_FILE_CACHING=1
export CUDA_DEVICE_WAITS_ON_EXCEPTION=0
.venv/bin/python -u router_probe.py --tag baseline --mode direct --seconds 10
# Apply this PR's patch, then use a fresh process:
.venv/bin/python -u router_probe.py --tag pr61463 --mode direct --seconds 120

Please select an isolated GPU with CUDA_VISIBLE_DEVICES. The README includes environment and patching instructions. Compare dmesg before/after each run for new Xid lines.

On 2026-10-16 we tested exactly the patch in this PR, with no other router changes:

Kernel	Result
Unpatched	Failed in the first batch: CUDA unspecified launch failure and new Xid13 / CTA Not Present
This PR's exact patch	4,887,808 router calls over a 120-second stress loop; router output and copied-byte checks passed; no additional Xid

The final recorded elapsed time is 130.383 seconds, including correctness checks after the loop. Generated PTX confirms the new arrive/wait precedes the first st.shared::cluster. This validates our isolated reproduction; we have not yet run a full-model serving soak test with this exact patch. Detailed results are in VALIDATION.txt in the attached ZIP.

This is a timing-sensitive reproduction established on B300 with the environment above; we have not established reproduction on Hopper. Hopefully it provides a smaller validation workload than a full P/D deployment.

llbf16-pr61463-repro.zip

[claude]

Claude Code Review

This pull request is from a fork — automated review is disabled. A repository maintainer can comment @claude review to run a one-time review.

[LopezCastroRoberto]

/ci run

[github-actions]

✅ Triggered Buildkite CI #99931 for commit 4d15a19372c6.