diff --git a/python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py b/python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py index ac8ddff9f..073935744 100644 --- a/python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py +++ b/python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py @@ -493,6 +493,21 @@ class TgvGemmCuteExtKernel: d_layout, ) + if cutlass.const_expr(self.use_2cta): + # Cluster-wide exit barrier (sgl-project/sglang#32907): cluster + # sync previously existed only at kernel entry, so one CTA could + # retire while its peer still had in-flight peer-redirected + # mbarrier arrives / multicast tcgen05.commit / TMEM dealloc + # targeting it -> CUDBG_EXCEPTION_CLUSTER_BLOCK_NOT_PRESENT + # (Xid 13 "CTA Not Present"). All 8 warps (256 threads) of both + # cluster CTAs arrive here after their dispatch branch returns, + # so no CTA can exit until the peer's tail work has landed. + # relaxed arrive (matches the entry barrier at L378): this fence is + # only about CTA lifetime, not inter-CTA smem visibility -- every + # cross-CTA datum above already flows through mbarrier phases. + cute.arch.cluster_arrive_relaxed() + cute.arch.cluster_wait() + # ==================================================================== # DMA_A WARP — TMA-loads A tiles into sA[..., stage], one per K-iter. # 1-CTA: SM90_TMA_LOAD into local sA, arrives on local bar_full.