fix(gin): de-alias Anvil signals across logical contexts

projects/rccl/CHANGELOG.md

@@ -75,6 +75,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * Fixed a 15 to 30% `ncclScatter` bandwidth loss from 16 MB upward on multi-node AINIC jobs, caused by the move to two network channels per peer (for example on 4 nodes of 8 GPUs). Scatter now uses and connects the most channels per peer whose flows still fit the P2P channel pool. Point-to-point operations and the other collectives keep their channel count.
 * Fixed intermittent data corruption in multi-node collectives, seen as out-of-bounds values in 8-node AllToAll, on hosts that use GPU Direct RDMA through peermem. Since RCCL 2.30.4 the GDR flush scratchpad was registered only through dma-buf, so on peermem hosts the flush degraded to a single RDMA read of the receive buffer, which does not order earlier relaxed-ordering writes. The scratchpad is now also registered through peermem, restoring the strictly ordered RDMA write that precedes the flush read.
 * Fixed `ncclCommAbort` and `ncclCommRevoke` hanging when a DDA collective (IPC on gfx942/gfx950, fabric LL/LL128/VMM on gfx1250) was waiting on a peer that never joined. DDA kernels now observe the communicator abort flag while waiting. A parent communicator shrunk with `NCCL_SHRINK_ABORT` no longer uses DDA for later collectives.
+* Fixed Anvil SDMA (`NCCL_GIN_TYPE=7`) GIN contexts sharing signal cells across logical contexts on a connection, which let a consecutive `ncclGinAllContexts` barrier return early on a stale count. Each logical context now has its own signals.
 
 ### Known issues
 * The FP8 ReduceSum and ReduceCopy device APIs are not exported in the LLVM bitcode library.

projects/rccl/docs/userguide/source/api/device_gin.rst

@@ -269,9 +269,6 @@ ncclGinBarrierSession
       Same as the single-context constructors, but arrival signaling and fencing iterate every GIN context on the comm.
       Signaling on each context preserves ordering when puts and signals use different network queue pairs. Use this when
       puts or gets were sharded across ``ginContextCount`` contexts and a ``Put`` or ``Get`` fence must drain all of them.
-      On the Anvil SDMA backend those contexts currently share one signal array, so a second AllContexts barrier in the
-      same kernel can observe the first barrier's counts (AICOMRCCL-2339). That is a plugin follow-up; the constructors
-      here still iterate every context.
 
    .. cpp:function:: void sync(Coop coop, cuda::memory_order order, ncclGinFenceLevel fence = ncclGinFenceLevel::Put | ncclGinFenceLevel::Get)
 

projects/rccl/src/dev_runtime.cc

@@ -1960,6 +1960,7 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom
   size_t ginSignalShadowsOffset = 0;
   size_t ginAnvilNetSignalsOffset = 0;
   int nGinContextsTotal = 0;
+  bool ginDevCommSetupDone = false;
   void* outDevCommPreserve = nullptr;
   struct ncclDevComm outDevCommTmp;
   cudaStreamCaptureMode captureMode = cudaStreamCaptureModeRelaxed;
@@ -2186,6 +2187,7 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom
     reqs->ginSignalCount = ginSignalTotal;
     reqs->ginCounterCount = ginCounterTotal;
     NCCLCHECKGOTO(ncclGinDevCommSetup(comm, reqs, outDevComm, deviceCodeVersion), ret, fail);
+    ginDevCommSetupDone = true;
     // SDMA signal binding is deferred until the resource window is created
     // (see ncclGinAnvilBindResourceWindowSignals call below).
   }
@@ -2296,6 +2298,13 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom
 fail_stream:
   CUDACHECKIGNORE(cudaStreamDestroy(stream));
 fail:
+  // Destroy the GIN contexts this call created. Without this, a failure between
+  // ncclGinDevCommSetup and the signal bind leaves them registered with the
+  // backend (Anvil keeps them pending), and a retried create binds the stale
+  // contexts alongside the new ones.
+  if (ginDevCommSetupDone && ncclGinDevCommFree(comm, outDevComm) == ncclSuccess) {
+    ncclGinDevCommClearFields(outDevComm);
+  }
   CUDACHECKIGNORE(cudaThreadExchangeStreamCaptureMode(&captureMode));
   return ret;
 }

projects/rccl/src/gin/gin_host.cc

@@ -430,18 +430,22 @@ static ncclResult_t ginDevCommSetupWithBackend(struct ncclComm* comm, struct ncc
       if (ginStateDevComm->ginCtx[commIdx]) backend->ncclGin->destroyContext(ginStateDevComm->ginCtx[commIdx]);
     }
     free(ginStateDevComm);
-    devComm->backendIndex = 0;
-    devComm->ginConnectionCount = 0;
-    devComm->ginContextCount = 0;
-    devComm->ginConnectionStride = 0;
-    devComm->ginConnectionStride_rcp32 = 0;
-    devComm->ginContextStride = 0;
-    memset(devComm->ginNetDeviceTypes, 0, sizeof(devComm->ginNetDeviceTypes));
-    memset(devComm->ginHandles, 0, sizeof(devComm->ginHandles));
+    ncclGinDevCommClearFields(devComm);
   }
   return ret;
 }
 
+void ncclGinDevCommClearFields(struct ncclDevComm* devComm) {
+  devComm->backendIndex = 0;
+  devComm->ginConnectionCount = 0;
+  devComm->ginContextCount = 0;
+  devComm->ginConnectionStride = 0;
+  devComm->ginConnectionStride_rcp32 = 0;
+  devComm->ginContextStride = 0;
+  memset(devComm->ginNetDeviceTypes, 0, sizeof(devComm->ginNetDeviceTypes));
+  for (auto& handle : devComm->ginHandles) handle = nullptr;
+}
+
 ncclResult_t ncclGinDevCommSetup(struct ncclComm* comm, struct ncclDevCommRequirements const* reqs,
                                  struct ncclDevComm* devComm, uint32_t deviceCodeVersion) {
   struct ncclGinState* ginState = &comm->sharedRes->ginState;