From 7d3766ce72b7f2272d2febe5feb4fc16d33c6bcc Mon Sep 17 00:00:00 2001 From: PanJason Date: Mon, 22 Jun 2026 17:56:01 +0200 Subject: [PATCH 1/5] [bugfix]: keep sync barrier single-queued - restore the sync loop exit so BARRIER remains posted through one representative queue - keep QUIET handling separate from BARRIER semantics for CXI timeout debugging --- ep/include/uccl_ibgda.cuh | 1 + 1 file changed, 1 insertion(+) diff --git a/ep/include/uccl_ibgda.cuh b/ep/include/uccl_ibgda.cuh index f774eddff..b8d46cdb6 100644 --- a/ep/include/uccl_ibgda.cuh +++ b/ep/include/uccl_ibgda.cuh @@ -409,6 +409,7 @@ __forceinline__ __device__ void nvshmem_sync_with_same_gpu_idx( } } #endif + break; } // Then wait for each proxy’s barrier to complete From b20d747f9f0edcf2910354f21d115151573dac23 Mon Sep 17 00:00:00 2001 From: PanJason Date: Mon, 22 Jun 2026 17:56:16 +0200 Subject: [PATCH 2/5] [bugfix]: preserve fifo control wr ids - widen quiet and barrier wr id storage so FIFO ring bits are not truncated - keep the existing negative sentinel while preserving high 32-bit queue indices --- ep/include/proxy_ctx.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ep/include/proxy_ctx.hpp b/ep/include/proxy_ctx.hpp index 85f12d4ef..ebb2eb6ef 100644 --- a/ep/include/proxy_ctx.hpp +++ b/ep/include/proxy_ctx.hpp @@ -201,10 +201,10 @@ struct ProxyCtx { // Async-barrier state (single inflight assumed) bool barrier_inflight = false; uint64_t barrier_seq = 0; - int barrier_wr = -1; + int64_t barrier_wr = -1; bool quiet_inflight = false; - int quiet_wr = -1; + int64_t quiet_wr = -1; // Rank-0 bookkeeping std::vector barrier_arrived; // size = num_ranks; 1 if arrival seen From 9c470bf18f84474953a3c04232509f152509b263 Mon Sep 17 00:00:00 2001 From: PanJason Date: Mon, 22 Jun 2026 17:57:02 +0200 Subject: [PATCH 3/5] [bugfix]: quiet all CXI proxy queues - post QUIET to every CXI D2H queue so sibling proxy queues are drained - track the posted queue index and wait on the exact queue for completion --- ep/include/uccl_ibgda.cuh | 26 ++++++++++++++++++-------- 1 file changed, 18 insertions(+), 8 deletions(-) diff --git a/ep/include/uccl_ibgda.cuh b/ep/include/uccl_ibgda.cuh index b8d46cdb6..f24f3039b 100644 --- a/ep/include/uccl_ibgda.cuh +++ b/ep/include/uccl_ibgda.cuh @@ -325,16 +325,23 @@ __device__ static __forceinline__ void nvshmemi_ibgda_quiet( EP_DEVICE_ASSERT( num_d2h_channel_addrs % kChannelPerProxy == 0 && "num_d2h_channel_addrs must be multiple of kChannelPerProxy"); - /* NOTE(MaoZiming): This is sent to all proxy threads. Since each proxy - * thread manages kChannelPerProxy ring buffers, we just need to post a quiet - * command to one out of the kChannelPerProxy ring buffer per cpu thread. */ EP_DEVICE_ASSERT(num_d2h_channel_addrs % kChannelPerProxy == 0); EP_DEVICE_ASSERT(num_d2h_channel_addrs / kChannelPerProxy == kNumProxyThs); - // First, atomically commit QUIET to one ring per proxy - uint64_t slots[kNumProxyThs]; + // First, atomically commit QUIET to all CXI rings. Each proxy thread owns + // multiple D2H rings, so posting to one ring per proxy is not enough to fence + // sibling rings before RDMA buffer reuse. +#if defined(USE_LIBFABRIC_CXI) + constexpr int kQuietStride = 1; + constexpr int kMaxQuietPosts = kNumProxyThs * kChannelPerProxy; +#else + constexpr int kQuietStride = kChannelPerProxy; + constexpr int kMaxQuietPosts = kNumProxyThs; +#endif + uint64_t slots[kMaxQuietPosts]; + int posted_d2h_channel_idxs[kMaxQuietPosts]; int num_posted = 0; for (int d2h_channel_idx = 0; d2h_channel_idx < num_d2h_channel_addrs; - d2h_channel_idx += kChannelPerProxy) { + d2h_channel_idx += kQuietStride) { auto* h = reinterpret_cast( static_cast(d2h_channel_addrs[d2h_channel_idx])); #ifdef USE_MSCCLPP_FIFO_BACKEND @@ -343,7 +350,9 @@ __device__ static __forceinline__ void nvshmemi_ibgda_quiet( TransferCmd cmd{}; cmd.cmd_type = CmdType::QUIET; h->atomic_set_and_commit(cmd, &slot); - slots[num_posted++] = slot; + slots[num_posted] = slot; + posted_d2h_channel_idxs[num_posted] = d2h_channel_idx; + ++num_posted; } #else while (true) { @@ -356,6 +365,7 @@ __device__ static __forceinline__ void nvshmemi_ibgda_quiet( cmd.cmd_type = CmdType::QUIET; h->atomic_set_and_commit(cmd, &slot); slots[num_posted] = slot; + posted_d2h_channel_idxs[num_posted] = d2h_channel_idx; ++num_posted; break; } @@ -366,7 +376,7 @@ __device__ static __forceinline__ void nvshmemi_ibgda_quiet( // Then wait for all QUIET commands to complete for (int i = 0; i < num_posted; ++i) { auto* h = reinterpret_cast( - static_cast(d2h_channel_addrs[i * kChannelPerProxy])); + static_cast(d2h_channel_addrs[posted_d2h_channel_idxs[i]])); wait_until_cmd_consumed(h, slots[i], nvl_rank, CmdType::QUIET); } } From e38d77c6240a4d7b234fb50048e45819ea81e47e Mon Sep 17 00:00:00 2001 From: Ziming Mao Date: Tue, 21 Jul 2026 09:18:28 -0700 Subject: [PATCH 4/5] Adding comment on the "break" in "sync" Co-authored-by: Ziming Mao --- ep/include/uccl_ibgda.cuh | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ep/include/uccl_ibgda.cuh b/ep/include/uccl_ibgda.cuh index f24f3039b..774583f07 100644 --- a/ep/include/uccl_ibgda.cuh +++ b/ep/include/uccl_ibgda.cuh @@ -419,6 +419,8 @@ __forceinline__ __device__ void nvshmem_sync_with_same_gpu_idx( } } #endif + // NOTE: the `break` here is intentional (only post to proxy thread 0): unlike + // QUIET, one proxy thread per GPU suffices since one GPU proxy thread is enough to form a complete barrier across all ranks. break; } From 54ef06429dde1c2debd3dbfb565530020650f7fd Mon Sep 17 00:00:00 2001 From: Ziming Mao Date: Tue, 21 Jul 2026 09:26:55 -0700 Subject: [PATCH 5/5] Fixing formating check --- ep/include/uccl_ibgda.cuh | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/ep/include/uccl_ibgda.cuh b/ep/include/uccl_ibgda.cuh index 774583f07..40b48e005 100644 --- a/ep/include/uccl_ibgda.cuh +++ b/ep/include/uccl_ibgda.cuh @@ -419,8 +419,9 @@ __forceinline__ __device__ void nvshmem_sync_with_same_gpu_idx( } } #endif - // NOTE: the `break` here is intentional (only post to proxy thread 0): unlike - // QUIET, one proxy thread per GPU suffices since one GPU proxy thread is enough to form a complete barrier across all ranks. + // NOTE: the `break` here is intentional (only post to proxy thread 0): + // unlike QUIET, one proxy thread per GPU suffices since one GPU proxy + // thread is enough to form a complete barrier across all ranks. break; }