From 2cdef395c3d6922bef07d6e52fb6af1490ca20d9 Mon Sep 17 00:00:00 2001 From: Matthew Bonanni Date: Tue, 13 Oct 2026 13:03:53 -0400 Subject: [PATCH 1/2] [Bugfix] Sync cluster before DSMEM stores in ll_bf16 split-K router GEMM The split-K epilogue stores partials into peer CTAs' shared memory via st.shared::cluster before any cluster barrier, so a peer that has not started yet raises Xid 13 (CTA Not Present). Co-Authored-By: Claude Opus 5.5 Signed-off-by: Matthew Bonanni --- .../model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py b/vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py index b44a5f6e2c94..5b2b898b3863 100644 --- a/vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py +++ b/vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py @@ -357,6 +357,10 @@ def kernel( consumer_group=consumer_group, ) + # Peer CTAs must be running before the epilogue's DSMEM stores. + cute.arch.cluster_arrive_relaxed() + cute.arch.cluster_wait() + # Round-robin split-K: split z handles tiles z, z+split_k, ... K_total = cute.size(mA, mode=[1]) k_tile_count = cute.size(gA, mode=[2]) From ac5c06ee4b14c0cd78a2fd6f5abb6c5730facb10 Mon Sep 17 00:00:00 2001 From: Matthew Bonanni Date: Tue, 3 Nov 2026 12:14:35 -0500 Subject: [PATCH 2/2] Split the leading cluster barrier out of the PDL prologue Arrive at kernel start; DMA warps wait after priming the first stage (past griddepcontrol_wait), MMA warps wait before the mainloop. Co-Authored-By: Claude Opus 5.5 Signed-off-by: Matthew Bonanni --- vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py b/vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py index 3ea7a7a4bbe3..1a67b398e6ae 100644 --- a/vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py +++ b/vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py @@ -359,7 +359,6 @@ def kernel( # Peer CTAs must be running before the epilogue's DSMEM stores. cute.arch.cluster_arrive_relaxed() - cute.arch.cluster_wait() # Round-robin split-K: split z handles tiles z, z+split_k, ... K_total = cute.size(mA, mode=[1]) @@ -406,6 +405,7 @@ def kernel( pred=tApA, ) mainloop_pipeline.producer_commit(producer_state) + cute.arch.cluster_wait() producer_state.advance() for k_tile in cutlass.range( @@ -432,6 +432,7 @@ def kernel( mainloop_pipeline.producer_tail(producer_state) else: + cute.arch.cluster_wait() # MMA warps with k-phase interleaving cute.arch.setmaxregister_increase(232) # large MMA fragments