From fbc145e2425059827524529019ef8cdc96ae9f94 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Fri, 9 Oct 2026 23:45:11 -0400 Subject: [PATCH 1/5] fix(devr): size GIN signal shadows for the rounded context count The device indexes ginSignalShadows by the devComm's ginContextCount, which ncclGinDevCommSetup rounds up to the chosen backend's connection count, but the shadow region was sized for the requested count. Size it, and the Anvil signal span, for the largest rounded count over every active backend. Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 1 + projects/rccl/src/dev_runtime.cc | 27 +++++++++------- projects/rccl/test/host/dev-runtime-test.cc | 35 +++++++++++++++++++++ 3 files changed, 52 insertions(+), 11 deletions(-) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 1c69375efc02..9d375d416533 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -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. diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index 6de9f55b54e8..e17321a7829d 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -1927,6 +1927,19 @@ void ncclDevCommDump(struct ncclDevComm* devComm) { printf(" GIN World Barrier signal0 %d\n", devComm->worldGinBarrier.signal0); } +// The devComm's ginContextCount is nGinContexts rounded up to the ginCommCount of +// 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 nGinContexts) { + int total = nGinContexts; + for (int b = 0; b < ginState->numActiveBackends; b++) { + int ginCommCount = std::max(ginState->backends[b].ginCommCount, 1); + total = std::max(total, ROUNDUP(nGinContexts, ginCommCount)); + } + return total; +} + ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCommRequirements* reqs, struct ncclDevComm* outDevComm, bool isInternal, uint32_t deviceCodeVersion) { ncclResult_t ret = ncclSuccess; @@ -2164,21 +2177,13 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom rr = rr->next; } bufSizeTotal = alignUp(bufSizeTotal, 128); + nGinContextsTotal = + devr->ginEnabled ? devrGinContextsTotal(&comm->sharedRes->ginState, nGinContexts) : nGinContexts; 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); } diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 449430e107f0..7317ca163194 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -6168,6 +6168,41 @@ TEST_F(DevCommCreateTest, RequirementsFilterFails_ReturnsErrorWithoutQueueing) { } +// --------------------------------------------------------------------------- +// devrGinContextsTotal sizes the per-context resource regions for the largest +// ginContextCount any active backend could round nGinContexts 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); +} + +// 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 From 92006c8d6a33591c837cf3060b65f0adcbdc333e Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sat, 10 Oct 2026 01:03:23 -0400 Subject: [PATCH 2/5] test(devr): pin max-over-backends order in gin context sizing MixedBackends_TakesLargestRoundedCount still passed if the helper kept only the last ROUNDUP, because the furthest-rounding backend was last. Assert the mirrored order so the max accumulator is what the suite locks. Co-authored-by: Cursor --- projects/rccl/test/host/dev-runtime-test.cc | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 7317ca163194..367342e9ad4d 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -6192,6 +6192,12 @@ TEST(DevrGinContextsTotal, MixedBackends_TakesLargestRoundedCount) { 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. From 49e0beed6cda6aa4034f5703908cc9b6bdf7ee8e Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Mon, 12 Oct 2026 02:39:31 -0400 Subject: [PATCH 3/5] test(devr): cover rounded GIN resource sizing Drive ncclDevrCommCreateInternal through the resource-window cuMemCreate call and assert that both signal regions use the largest backend-rounded context count. This pins the production integration at dev_runtime.cc:2092-2098, which helper-only tests did not cover. Co-authored-by: Cursor --- projects/rccl/test/host/dev-runtime-test.cc | 41 +++++++++++++++++++ .../host/fakes/dev_runtime_micro_fakes.cc | 9 ++-- 2 files changed, 47 insertions(+), 3 deletions(-) diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 367342e9ad4d..715d06dd1ea1 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -6220,6 +6220,7 @@ TEST(DevrGinContextsTotal, ZeroConnectionBackend_TreatedAsOne) { class DevrCommCreateInternalTest : public NcclVersionCompatTest { protected: std::unique_ptr commStorage; + std::unique_ptr sharedResStorage; ncclComm* comm = nullptr; ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER; ncclDevComm outDevComm{}; @@ -6236,9 +6237,23 @@ class DevrCommCreateInternalTest : public NcclVersionCompatTest { ncclDevCommCompat_v22902.maxVersion = NCCL_VERSION_CODE; commStorage = std::make_unique(); // value-initialised: POD members zeroed + sharedResStorage = std::make_unique(); 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() { @@ -6288,6 +6303,32 @@ 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; + 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 diff --git a/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc b/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc index 3ece503ad9bb..b79afaff0914 100644 --- a/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc +++ b/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc @@ -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; } From 6578783e3758790185f719e70fe57ec75bdda696 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Mon, 12 Oct 2026 14:43:23 -0400 Subject: [PATCH 4/5] test(devr): pin ginContextCount and refresh create-suite comments Co-authored-by: Cursor --- projects/rccl/test/host/dev-runtime-test.cc | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 715d06dd1ea1..1b1a1ec62e98 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -6212,10 +6212,11 @@ TEST(DevrGinContextsTotal, ZeroConnectionBackend_TreatedAsOne) { // --------------------------------------------------------------------------- // 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: @@ -6293,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 @@ -6309,6 +6311,8 @@ TEST_F(DevrCommCreateInternalTest, GinForceEnable_BehavesAsFullConnection) { // 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; From 913e54a94c643958600bc5fccb144e2b37fb3a94 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Fri, 23 Oct 2026 14:34:21 -0400 Subject: [PATCH 5/5] refactor(gin): share the context round-up between devComm setup and devr sizing Add ncclGinContextCount() in gin_host.h and use it both where ginDevCommSetupWithBackend sets devComm->ginContextCount and where devrGinContextsTotal sizes the per-context resource regions, so the two cannot diverge. Rename nGinContexts to nGinContextsRequested in ncclDevrCommCreateInternal: it only seeds the round-up. Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 19 ++++++++++--------- projects/rccl/src/gin/gin_host.cc | 2 +- projects/rccl/src/include/gin/gin_host.h | 8 ++++++++ projects/rccl/test/host/dev-runtime-test.cc | 2 +- 4 files changed, 20 insertions(+), 11 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index e17321a7829d..4b3dea639216 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -1927,15 +1927,15 @@ void ncclDevCommDump(struct ncclDevComm* devComm) { printf(" GIN World Barrier signal0 %d\n", devComm->worldGinBarrier.signal0); } -// The devComm's ginContextCount is nGinContexts rounded up to the ginCommCount of -// 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 nGinContexts) { - int total = nGinContexts; +// 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, ROUNDUP(nGinContexts, ginCommCount)); + total = std::max(total, ncclGinContextCount(nGinContextsRequested, ginCommCount)); } return total; } @@ -1952,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; @@ -2178,7 +2178,8 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom } bufSizeTotal = alignUp(bufSizeTotal, 128); nGinContextsTotal = - devr->ginEnabled ? devrGinContextsTotal(&comm->sharedRes->ginState, nGinContexts) : nGinContexts; + devr->ginEnabled ? devrGinContextsTotal(&comm->sharedRes->ginState, nGinContextsRequested) + : nGinContextsRequested; ginSignalShadowsOffset = bufSizeTotal; bufSizeTotal += (size_t)nGinContextsTotal * ginSignalTotal * sizeof(uint64_t); // include signal shadows ginAnvilNetSignalsOffset = bufSizeTotal; diff --git a/projects/rccl/src/gin/gin_host.cc b/projects/rccl/src/gin/gin_host.cc index 96ea672d6c59..a336b2658099 100644 --- a/projects/rccl/src/gin/gin_host.cc +++ b/projects/rccl/src/gin/gin_host.cc @@ -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, diff --git a/projects/rccl/src/include/gin/gin_host.h b/projects/rccl/src/include/gin/gin_host.h index bccc2a107c45..bbc8a4e8278e 100644 --- a/projects/rccl/src/include/gin/gin_host.h +++ b/projects/rccl/src/include/gin/gin_host.h @@ -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 diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 1b1a1ec62e98..9bd849a858ad 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -6170,7 +6170,7 @@ TEST_F(DevCommCreateTest, RequirementsFilterFails_ReturnsErrorWithoutQueueing) { // --------------------------------------------------------------------------- // devrGinContextsTotal sizes the per-context resource regions for the largest -// ginContextCount any active backend could round nGinContexts up to. +// ginContextCount any active backend could round nGinContextsRequested up to. TEST(DevrGinContextsTotal, NoActiveBackends_ReturnsRequested) { struct ncclGinState ginState{};