From 32c6e33cb33017195e5f6e97b537371fd17ef649 Mon Sep 17 00:00:00 2001 From: yuzho-amd Date: Thu, 23 Jul 2026 04:03:16 -0400 Subject: [PATCH] fix allreduce refill issue --- csrc/include/custom_all_reduce.cuh | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/csrc/include/custom_all_reduce.cuh b/csrc/include/custom_all_reduce.cuh index 9e635209cf6..ee304c132c9 100644 --- a/csrc/include/custom_all_reduce.cuh +++ b/csrc/include/custom_all_reduce.cuh @@ -2044,6 +2044,11 @@ __global__ void __launch_bounds__(1024, 1) *reinterpret_cast(output + out_idx) = zero_pack; } } + // Pair start_sync with a final barrier before any rank can reuse the + // registered input buffer or this block's signal slot for the next fused + // collective. Without it, a faster rank may advance while a peer is still + // reading the previous invocation, causing silent cross-request corruption. + end_sync(sg, self_sg, rank); } // Per-group quant variant of the 1-stage fused allreduce+rmsnorm kernel. @@ -2127,6 +2132,7 @@ __global__ void __launch_bounds__(1024, 1) acc, weight_p, hidden_dim, eps, idx, tidx, padded_block_size, group_size, output, scale_out, active, bf16_output, token_num); } + end_sync(sg, self_sg, rank); } template @@ -2236,6 +2242,7 @@ __global__ void __launch_bounds__(1024, 1) acc, weight_p, hidden_dim, eps, idx, tidx, padded_block_size, output, scale_out, active, bf16_output); } + end_sync(sg, self_sg, rank); } template