dlamd1dai · Oct 26, 2026 · Oct 10, 2026 · Oct 10, 2026 · Oct 12, 2026 · Oct 12, 2026
@@ -78,6 +78,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * 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.
+* Fixed GIN signal shadows being sized for the requested context count instead of the rounded-up `ginContextCount`, which let the last contexts write past the shadow region when the GIN connection count does not divide the requested context count.
 
 ### Known issues
 * The FP8 ReduceSum and ReduceCopy device APIs are not exported in the LLVM bitcode library.

@@ -1927,6 +1927,19 @@ void ncclDevCommDump(struct ncclDevComm* devComm) {
   printf(" GIN World Barrier signal0 %d\n", devComm->worldGinBarrier.signal0);
 }
 
+// The devComm's ginContextCount is ncclGinContextCount() for whichever backend
+// ncclGinDevCommSetup picks, and per-context resource regions are indexed by it.
+// Size for the worst case over every active backend: rounding to the widest
+// ginCommCount alone is not enough (ROUNDUP(4, 4) < ROUNDUP(4, 3)).
+static int devrGinContextsTotal(const struct ncclGinState* ginState, int nGinContextsRequested) {
+  int total = nGinContextsRequested;
+  for (int b = 0; b < ginState->numActiveBackends; b++) {
+    int ginCommCount = std::max(ginState->backends[b].ginCommCount, 1);
+    total = std::max(total, ncclGinContextCount(nGinContextsRequested, ginCommCount));
+  }
+  return total;
+}
+
 ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCommRequirements* reqs,
                                         struct ncclDevComm* outDevComm, bool isInternal, uint32_t deviceCodeVersion) {
   ncclResult_t ret = ncclSuccess;
@@ -1939,7 +1952,7 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom
   struct ncclDevrTeam* tmLsa;
   ncclTeam_t ucTeam, mcTeam;
   size_t bufSizeTotal;
-  int nGinContexts = reqs->ginContextCount;
+  int nGinContextsRequested = reqs->ginContextCount;
   int ginSignalTotal = 0, ginCounterTotal = 0;
   struct ncclDevResourceRequirements* resReqsHead = reqs->resourceRequirementsList;
   struct ncclDevResourceRequirements lsaBarReq;
@@ -2164,21 +2177,14 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom
       rr = rr->next;
     }
     bufSizeTotal = alignUp(bufSizeTotal, 128);
+    nGinContextsTotal =
+      devr->ginEnabled ? devrGinContextsTotal(&comm->sharedRes->ginState, nGinContextsRequested)
+                       : nGinContextsRequested;
     ginSignalShadowsOffset = bufSizeTotal;
-    bufSizeTotal += nGinContexts * ginSignalTotal * sizeof(uint64_t); // include signal shadows
+    bufSizeTotal += (size_t)nGinContextsTotal * ginSignalTotal * sizeof(uint64_t); // include signal shadows
     ginAnvilNetSignalsOffset = bufSizeTotal;
     if (devr->ginEnabled && ginSignalTotal > 0) {
-      // 2.31 moved ginCommCount into per-backend state; round up to the widest
-      // active backend so every backend's signal shadows stay in range.
-      struct ncclGinState* gs = &comm->sharedRes->ginState;
-      int ginCommCount = 1;
-      for (int b = 0; b < gs->numActiveBackends; b++) {
-        if (gs->backends[b].ginCommCount > ginCommCount) ginCommCount = gs->backends[b].ginCommCount;
-      }
-      nGinContextsTotal = ROUNDUP(nGinContexts, ginCommCount);
       bufSizeTotal += (size_t)nGinContextsTotal * ginSignalTotal * sizeof(uint64_t);
-    } else {
-      nGinContextsTotal = nGinContexts;
     }
     bufSizeTotal = alignUp(bufSizeTotal, devr->granularity);
   }

@@ -286,7 +286,7 @@ static ncclResult_t ginDevCommSetupWithBackend(struct ncclComm* comm, struct ncc
     // TODO: check if a shared devComm in the list could match our requirements.
   }
 
-  nContextsTotal = ROUNDUP(nContextsTotal, backend->ginCommCount);
+  nContextsTotal = ncclGinContextCount(nContextsTotal, backend->ginCommCount);
   devComm->ginContextCount = nContextsTotal;
   int nContextsPerComm = nContextsTotal / backend->ginCommCount;
   INFO(NCCL_INIT,

@@ -9,6 +9,7 @@
 #define _NCCL_GIN_HOST_H_
 
 #include "allocator.h"
+#include "bitops.h"
 #include "nccl.h"
 #include "nccl_gin.h"
 #include "nccl_device/core_tmp.h"
@@ -77,6 +78,13 @@ struct ncclGinState {
   struct ncclGinBackendState backends[NCCL_GIN_MAX_ACTIVE_BACKENDS];
 };
 
+// The ginContextCount a devComm gets from a backend with ginCommCount connections.
+// ncclDevrCommCreateInternal sizes per-context resource regions with it (via
+// devrGinContextsTotal), so the two must not diverge.
+static inline int ncclGinContextCount(int nGinContextsRequested, int ginCommCount) {
+  return ROUNDUP(nGinContextsRequested, ginCommCount);
+}
+
 extern int64_t ncclParamGinType();
 
 // Get the GIN type from comm. ginType is set to the GIN type that can be used

@@ -6168,17 +6168,60 @@ TEST_F(DevCommCreateTest, RequirementsFilterFails_ReturnsErrorWithoutQueueing) {
 }
 
 
+// ---------------------------------------------------------------------------
+// devrGinContextsTotal sizes the per-context resource regions for the largest
+// ginContextCount any active backend could round nGinContextsRequested up to.
+
+TEST(DevrGinContextsTotal, NoActiveBackends_ReturnsRequested) {
+  struct ncclGinState ginState{};
+  EXPECT_EQ(devrGinContextsTotal(&ginState, 4), 4);
+}
+
+TEST(DevrGinContextsTotal, SingleBackend_RoundsUpToConnectionCount) {
+  struct ncclGinState ginState{};
+  ginState.numActiveBackends = 1;
+  ginState.backends[0].ginCommCount = 3;
+  EXPECT_EQ(devrGinContextsTotal(&ginState, 8), 9);
+  EXPECT_EQ(devrGinContextsTotal(&ginState, 6), 6);
+}
+
+// Branch: the widest backend is not the one that rounds furthest.
+TEST(DevrGinContextsTotal, MixedBackends_TakesLargestRoundedCount) {
+  struct ncclGinState ginState{};
+  ginState.numActiveBackends = 2;
+  ginState.backends[0].ginCommCount = 4;
+  ginState.backends[1].ginCommCount = 3;
+  EXPECT_EQ(devrGinContextsTotal(&ginState, 4), 6);
+
+  // Mirrored order: the further-rounding backend must still win when it is
+  // visited first. Without the max-over-backends accumulator this would return 4.
+  ginState.backends[0].ginCommCount = 3;
+  ginState.backends[1].ginCommCount = 4;
+  EXPECT_EQ(devrGinContextsTotal(&ginState, 4), 6);
+}
+
+// Branch: a backend reporting no connections must not divide by zero.
+TEST(DevrGinContextsTotal, ZeroConnectionBackend_TreatedAsOne) {
+  struct ncclGinState ginState{};
+  ginState.numActiveBackends = 1;
+  ginState.backends[0].ginCommCount = 0;
+  EXPECT_EQ(devrGinContextsTotal(&ginState, 5), 5);
+}
+
+
 // ---------------------------------------------------------------------------
 // ncclDevrCommCreateInternal is where a queued create is actually carried out.
 // This suite covers its GIN request validation, which runs before any resource
-// is touched and rejects combinations the communicator cannot serve.
+// is touched and rejects combinations the communicator cannot serve, and the
+// resource-window sizing that follows it, up to the cuMemCreate call.
 //
-// The body past that gate builds the whole devcomm -- GIN activation, resource
-// windows, barriers -- and is not covered here.
+// The rest of the devcomm build -- symmetric-memory window creation, GIN
+// window registration, barriers -- is not covered here.
 
 class DevrCommCreateInternalTest : public NcclVersionCompatTest {
 protected:
   std::unique_ptr<ncclComm> commStorage;
+  std::unique_ptr<ncclSharedResources> sharedResStorage;
   ncclComm* comm = nullptr;
   ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
   ncclDevComm outDevComm{};
@@ -6195,9 +6238,23 @@ class DevrCommCreateInternalTest : public NcclVersionCompatTest {
     ncclDevCommCompat_v22902.maxVersion = NCCL_VERSION_CODE;
 
     commStorage = std::make_unique<ncclComm>();  // value-initialised: POD members zeroed
+    sharedResStorage = std::make_unique<ncclSharedResources>();
     comm = commStorage.get();
+    comm->sharedRes = sharedResStorage.get();
     comm->nRanks = 1;
     comm->devrState.lsaSize = 1;
+    comm->devrState.cftSize = 1;
+    comm->devrState.cftMcSize = 1;
+    comm->devrState.granularity = 1;
+  }
+
+  void TearDown() override {
+    while (comm->devrState.teamHead != nullptr) {
+      ncclDevrTeam* next = comm->devrState.teamHead->next;
+      free(comm->devrState.teamHead);
+      comm->devrState.teamHead = next;
+    }
+    NcclVersionCompatTest::TearDown();
   }
 
   ncclResult_t Create() {
@@ -6237,8 +6294,9 @@ TEST_F(DevrCommCreateInternalTest, FullGinRequestedOnRailOnlyComm_ReturnsInvalid
 // connection, so it fails the same support check rather than being ignored.
 //
 // The counterpart -- the same requirements with the flag clear, which requests
-// nothing and passes the gate -- is not tested: it continues into the devcomm
-// build, which this fixture does not set up for.
+// nothing and passes the gate -- is not tested: passing the gate is only
+// observable once the devcomm build completes, and this fixture stops that
+// build at the resource-window cuMemCreate.
 TEST_F(DevrCommCreateInternalTest, GinForceEnable_BehavesAsFullConnection) {
   reqs.ginForceEnable = true;
   reqs.ginConnectionType = NCCL_GIN_CONNECTION_NONE;  // overridden by the flag
@@ -6247,6 +6305,34 @@ TEST_F(DevrCommCreateInternalTest, GinForceEnable_BehavesAsFullConnection) {
   EXPECT_EQ(Create(), ncclInvalidArgument);
 }
 
+// Integration point: the resource-window allocation must use the rounded
+// context count, not merely the helper's result in isolation. Stop at
+// cuMemCreate after recording its requested size to avoid constructing the
+// unrelated symmetric-memory window.
+TEST_F(DevrCommCreateInternalTest, GinSignalRegions_UseLargestRoundedContextCount) {
+  reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL;
+  // 4 discriminates: ROUNDUP(4, 4) = 4 but ROUNDUP(4, 3) = 6.
+  reqs.ginContextCount = 4;
+  reqs.ginSignalCount = 1;
+  comm->globalGinSupport = NCCL_GIN_CONNECTION_FULL;
+  comm->sharedRes->ginState.numActiveBackends = 2;
+  comm->sharedRes->ginState.backends[0].ginCommCount = 3;
+  comm->sharedRes->ginState.backends[1].ginCommCount = 4;
+
+  size_t allocationSize = 0;
+  ScopedHook create(g_hipMemCreate,
+                    [&](hipMemGenericAllocationHandle_t*, size_t size, const hipMemAllocationProp*,
+                        unsigned long long) {
+                      allocationSize = size;
+                      return hipErrorInvalidValue;
+                    });
+
+  EXPECT_NE(Create(), ncclSuccess);
+  EXPECT_EQ(create.calls, 1);
+  // ROUNDUP(4, 3) = 6 contexts; shadow and Anvil regions each use 6 * 1 * 8 bytes.
+  EXPECT_EQ(allocationSize, 96u);
+}
+
 
 // ---------------------------------------------------------------------------
 // ncclDevCommCopyLsaData copies the LSA-shared span of a devcomm between two

@@ -492,15 +492,18 @@ ncclResult_t symTeamObtainMcLe(struct ncclComm*, struct ncclDevrTeam*, struct nc
 // Barrier requirement builders (host variants).
 // ---------------------------------------------------------------------------
 extern "C" ncclResult_t ncclLsaBarrierCreateRequirement(ncclTeam_t, int, ncclLsaBarrierHandle_t*,
-                                                        ncclDevResourceRequirements_t*) {
+                                                        ncclDevResourceRequirements_t* req) {
+  if (req != nullptr) *req = {};
   return ncclSuccess;
 }
 extern "C" ncclResult_t ncclGinBarrierCreateRequirement(ncclComm_t, ncclTeam_t, int, ncclGinBarrierHandle_t*,
-                                                        ncclDevResourceRequirements_t*) {
+                                                        ncclDevResourceRequirements_t* req) {
+  if (req != nullptr) *req = {};
   return ncclSuccess;
 }
 extern "C" ncclResult_t ncclCftBarrierCreateRequirement(ncclTeam_t, int, ncclCftBarrierHandle_t*,
-                                                        ncclDevResourceRequirements_t*) {
+                                                        ncclDevResourceRequirements_t* req) {
+  if (req != nullptr) *req = {};
   return ncclSuccess;
 }