From bb5968ddb636e3c8416a7968ee0e1096d026ef68 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sat, 10 Oct 2026 01:44:02 -0400 Subject: [PATCH 1/4] fix(rccl): latch GIN only after a successful connect ncclDevrCommCreateInternal set devr->ginEnabled before ncclGinConnectOnce and never cleared it. A failed connect left ginCommCount at 0, so a retry skipped the reconnect and divided by zero in ginDevCommSetupWithBackend. Set the latch only once connect succeeds, so the next create reconnects. AICOMRCCL-2879 Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 4 ++- projects/rccl/test/host/dev-runtime-test.cc | 31 ++++++++++++++++++- .../host/fakes/dev_runtime_micro_fakes.cc | 7 ++++- .../test/host/fakes/dev_runtime_micro_fakes.h | 1 + 4 files changed, 40 insertions(+), 3 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index f5376919b787..494c96c50d70 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -1910,7 +1910,6 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom } ginActivated = !devr->ginEnabled; - devr->ginEnabled = true; } if (reqs->worldGinBarrierCount > 0 && requestedConnectionType == NCCL_GIN_CONNECTION_RAIL) { @@ -1928,6 +1927,9 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom if (ginActivated) { NCCLCHECKGOTO(ncclGinConnectOnce(comm), ret, fail); + // Latch only on success: a failed connect leaves ginCommCount at 0, and a + // retry must reconnect rather than reach ncclGinDevCommSetup. + devr->ginEnabled = true; // Register all preexisting memories with GIN. Update the windows later when // we have a stream. for (struct ncclDevrMemory* mem = devr->memHead; mem != nullptr; mem = mem->next) { diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 991b180dea8f..5576533af0c1 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -5508,7 +5508,7 @@ TEST_F(DevCommCreateTest, RequirementsFilterFails_ReturnsErrorWithoutQueueing) { // is touched and rejects combinations the communicator cannot serve. // // The body past that gate builds the whole devcomm -- GIN activation, resource -// windows, barriers -- and is not covered here. +// windows, barriers -- and is not covered here beyond a failed GIN connect. class DevrCommCreateInternalTest : public NcclVersionCompatTest { protected: @@ -5581,6 +5581,35 @@ TEST_F(DevrCommCreateInternalTest, GinForceEnable_BehavesAsFullConnection) { EXPECT_EQ(Create(), ncclInvalidArgument); } +// Branch: GIN connect fails. GIN must stay unlatched so the next create +// reconnects; a latched retry would skip connect and reach ncclGinDevCommSetup +// with ginCommCount still 0 (AICOMRCCL-2879). +TEST_F(DevrCommCreateInternalTest, GinConnectFails_RetryConnectsAgain) { + reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL; + comm->globalGinSupport = NCCL_GIN_CONNECTION_FULL; + ScopedHook connect(g_devrGinConnectOnce, [](ncclComm*) { return ncclSystemError; }); + + EXPECT_EQ(Create(), ncclSystemError); + EXPECT_FALSE(comm->devrState.ginEnabled); + + EXPECT_EQ(Create(), ncclSystemError); + EXPECT_EQ(connect.calls, 2); + EXPECT_FALSE(comm->devrState.ginEnabled); +} + +// Branch: a GIN request rejected after the support checks but before connect +// leaves GIN unlatched, so a later valid create still connects. +TEST_F(DevrCommCreateInternalTest, WorldGinBarrierOnRail_RejectedWithoutLatchingGin) { + reqs.ginConnectionType = NCCL_GIN_CONNECTION_RAIL; + reqs.worldGinBarrierCount = 1; + comm->globalGinSupport = NCCL_GIN_CONNECTION_RAIL; + ScopedHook connect(g_devrGinConnectOnce, [](ncclComm*) { return ncclSuccess; }); + + EXPECT_EQ(Create(), ncclInvalidArgument); + EXPECT_EQ(connect.calls, 0); + EXPECT_FALSE(comm->devrState.ginEnabled); +} + // --------------------------------------------------------------------------- // 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 0cb16f662c1b..eaa9f0106218 100644 --- a/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc +++ b/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc @@ -312,7 +312,11 @@ ncclResult_t ncclGetRailedGinType(struct ncclComm*, ncclGinType_t* ginType) { if (ginType) *ginType = NCCL_GIN_TYPE_NONE; return ncclSuccess; } -ncclResult_t ncclGinConnectOnce(struct ncclComm*) { return ncclSuccess; } +// Seam: ncclDevrCommCreateInternal latches GIN only after this succeeds. +static ncclResult_t DefaultGinConnectOnce(struct ncclComm*) { return ncclSuccess; } +std::function g_devrGinConnectOnce = DefaultGinConnectOnce; + +ncclResult_t ncclGinConnectOnce(struct ncclComm* comm) { return g_devrGinConnectOnce(comm); } ncclResult_t ncclGinDevCommSetup(struct ncclComm*, struct ncclDevCommRequirements const*, struct ncclDevComm*) { return ncclSuccess; } @@ -546,6 +550,7 @@ void ResetDevRuntimeMicroFakes() { g_devrBootstrapAllGather = DefaultAllGather; g_devrBootstrapSend = DefaultBootstrapSend; g_devrBootstrapRecv = DefaultBootstrapRecv; + g_devrGinConnectOnce = DefaultGinConnectOnce; g_devrGinRegister = DefaultGinRegister; g_devrGinDeregister = DefaultGinDeregister; g_devrRmaProxyConnectOnce = DefaultRmaProxyConnectOnce; diff --git a/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.h b/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.h index 2674fa211227..200027f4b5d2 100644 --- a/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.h +++ b/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.h @@ -57,6 +57,7 @@ extern std::function g_devrBoot extern std::function g_devrBootstrapSend; extern std::function g_devrBootstrapRecv; +extern std::function g_devrGinConnectOnce; extern std::function g_devrGinRegister; From 1da36b4d24634edbff1f84ac40d01719ce0ee8aa Mon Sep 17 00:00:00 2001 From: Aurelien Bouteiller Date: Fri, 23 Oct 2026 17:53:31 -0400 Subject: [PATCH 2/4] Update projects/rccl/src/dev_runtime.cc Co-authored-by: amd-argus[bot] <321830313+amd-argus[bot]@users.noreply.github.com> --- projects/rccl/src/dev_runtime.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index 494c96c50d70..04f6b73e2291 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -1928,7 +1928,8 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom if (ginActivated) { NCCLCHECKGOTO(ncclGinConnectOnce(comm), ret, fail); // Latch only on success: a failed connect leaves ginCommCount at 0, and a - // retry must reconnect rather than reach ncclGinDevCommSetup. + // Latch only on success: a failed connect leaves ginCommCount stale or 0, and + // a retry must reconnect rather than reach ncclGinDevCommSetup. devr->ginEnabled = true; // Register all preexisting memories with GIN. Update the windows later when // we have a stream. From 552719213d1e363ce56edb719c414a2952226d84 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 25 Oct 2026 23:30:29 -0400 Subject: [PATCH 3/4] fix(rccl): drop the duplicated GIN latch comment and fail the retry test fast The suggestion apply left both comment lines in place, and EXPECT_FALSE let a latched ginEnabled fall through into a host-test divide by zero. Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 1 - projects/rccl/test/host/dev-runtime-test.cc | 2 +- 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index 04f6b73e2291..559a6a765c47 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -1927,7 +1927,6 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom if (ginActivated) { NCCLCHECKGOTO(ncclGinConnectOnce(comm), ret, fail); - // Latch only on success: a failed connect leaves ginCommCount at 0, and a // Latch only on success: a failed connect leaves ginCommCount stale or 0, and // a retry must reconnect rather than reach ncclGinDevCommSetup. devr->ginEnabled = true; diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 5576533af0c1..4437f267bdd8 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -5590,7 +5590,7 @@ TEST_F(DevrCommCreateInternalTest, GinConnectFails_RetryConnectsAgain) { ScopedHook connect(g_devrGinConnectOnce, [](ncclComm*) { return ncclSystemError; }); EXPECT_EQ(Create(), ncclSystemError); - EXPECT_FALSE(comm->devrState.ginEnabled); + ASSERT_FALSE(comm->devrState.ginEnabled); EXPECT_EQ(Create(), ncclSystemError); EXPECT_EQ(connect.calls, 2); From ad48d9dde0523cdb17d1113585f26f4fd03e5ad8 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Mon, 26 Oct 2026 00:28:23 -0400 Subject: [PATCH 4/4] fix(rccl): roll back a GIN activation that fails after connect ncclDevrCommCreateInternal latches ginEnabled once ncclGinConnectOnce succeeds, then registers existing memories with GIN and updates existing windows. A failure in between left GIN latched with some memories unregistered and some windows missing their GIN handles. The next create saw ginEnabled, skipped activation, and built a GIN devcomm over them. Until the window update finishes, a failure now deregisters every memory this activation registered, frees the segment windows it attached, and clears the latch, so the next create activates GIN again from scratch. AICOMRCCL-2879 Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 42 +++++++++++++++ projects/rccl/test/host/dev-runtime-test.cc | 60 ++++++++++++++++++++- 2 files changed, 101 insertions(+), 1 deletion(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index 559a6a765c47..5cc8af79a57e 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -827,6 +827,42 @@ static void symMemoryUnregister(struct ncclComm* comm, struct ncclDevrMemory* me } } +// Undo an incomplete GIN activation in ncclDevrCommCreateInternal. Memories and +// windows that existed before it carried no GIN state, so everything GIN holds +// now came from this activation. +static void symGinActivationRollbackWindows(struct ncclDevrState* devr, cudaStream_t stream) { + for (int i = 0; i < devr->winSortedCount; i++) { + struct ncclDevrWindow* win = devr->winSorted[i].win; + struct ncclWindow_vidmem* winHost; + if (ncclShadowPoolToHost(&devr->shadows, win->vidmem, &winHost) != ncclSuccess) continue; + if (winHost->ginMultiSegmentWins == nullptr) continue; + (void)ncclShadowPoolFree(&devr->shadows, winHost->ginMultiSegmentWins, stream); + winHost->ginMultiSegmentWins = nullptr; + winHost->numSegments = 0; + for (int c = 0; c < NCCL_GIN_MAX_CONNECTIONS; c++) winHost->ginWinsDefaultBackend[c] = nullptr; + CUDACHECKIGNORE( + cudaMemcpyAsync(win->vidmem, winHost, sizeof(struct ncclWindow_vidmem), cudaMemcpyHostToDevice, stream)); + } +} + +static void symGinActivationRollbackMemories(struct ncclComm* comm) { + struct ncclDevrState* devr = &comm->devrState; + for (struct ncclDevrMemory* mem = devr->memHead; mem != nullptr; mem = mem->next) { + if (mem->ginSegmentInfos == nullptr) continue; + for (int segment = 0; segment < mem->numGinSegments; segment++) { + (void)ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); + } + free(mem->ginSegmentInfos); + mem->ginSegmentInfos = nullptr; + mem->numGinSegments = 1; + for (int i = 0; i < NCCL_GIN_MAX_CONNECTIONS * NCCL_GIN_MAX_ACTIVE_BACKENDS; i++) { + mem->ginHostWins[i] = nullptr; + mem->ginDevWins[i] = nullptr; + } + } + devr->ginEnabled = false; +} + static void symMemoryUnmapLsaRank(struct ncclDevrState* devr, struct ncclDevrMemory* mem, int rank) { uintptr_t base = reinterpret_cast(devr->lsaFlatBase); uintptr_t addr = base + rank * devr->bigSize + mem->bigOffset; @@ -1847,6 +1883,7 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom struct ncclTeam world = ncclTeamWorld(comm); struct ncclTeam lsa = ncclTeamInnerFactor(world, devr->lsaSize); bool ginActivated = false; + bool ginActivationDone = false; bool cftUcActivated = false; bool hasCountedCftMemory = symHasCountedCftMemory(devr); struct ncclDevrTeam* tmLsa; @@ -2133,6 +2170,7 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom stream), ret, fail_stream); } + ginActivationDone = true; } NCCLCHECKGOTO(symWindowTableInitOnce(comm, stream), ret, fail_stream); // ensure devr->windowTable exists @@ -2209,8 +2247,12 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom } symMemoryDropRef(comm, mem); fail_stream: + if (ginActivated && !ginActivationDone) symGinActivationRollbackWindows(devr, stream); CUDACHECKIGNORE(cudaStreamDestroy(stream)); fail: + // A retry sees ginEnabled and skips activation, so leaving GIN latched here + // would build GIN devcomms over unregistered memories and stale windows. + if (ginActivated && !ginActivationDone) symGinActivationRollbackMemories(comm); CUDACHECKIGNORE(cudaThreadExchangeStreamCaptureMode(&captureMode)); return ret; } diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 4437f267bdd8..b5a0535029ed 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -5508,7 +5508,7 @@ TEST_F(DevCommCreateTest, RequirementsFilterFails_ReturnsErrorWithoutQueueing) { // is touched and rejects combinations the communicator cannot serve. // // The body past that gate builds the whole devcomm -- GIN activation, resource -// windows, barriers -- and is not covered here beyond a failed GIN connect. +// windows, barriers -- and is not covered here beyond a failed GIN activation. class DevrCommCreateInternalTest : public NcclVersionCompatTest { protected: @@ -5597,6 +5597,64 @@ TEST_F(DevrCommCreateInternalTest, GinConnectFails_RetryConnectsAgain) { EXPECT_FALSE(comm->devrState.ginEnabled); } +// Branch: connect succeeds and GIN latches, then registering an existing memory +// fails. The memory registered before it is deregistered again and GIN +// unlatches, so the next create runs the whole activation again instead of +// building a GIN devcomm over unregistered memory. +TEST_F(DevrCommCreateInternalTest, GinRegisterFailsAfterConnect_RollsBackActivation) { + reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL; + comm->globalGinSupport = NCCL_GIN_CONNECTION_FULL; + + size_t segmentSizes[2] = {8192, 8192}; + hipMemGenericAllocationHandle_t memHandles[2] = {reinterpret_cast(0x55), + reinterpret_cast(0x56)}; + ncclDevrMemory mems[2] = {}; + for (int m = 0; m < 2; m++) { + mems[m].primaryAddr = reinterpret_cast(0x100000 * (m + 1)); + mems[m].size = segmentSizes[m]; + mems[m].numSegments = 1; + mems[m].maxGlobalNumSegments = 1; + mems[m].segmentSizes = &segmentSizes[m]; + mems[m].memHandles = &memHandles[m]; + } + mems[0].next = &mems[1]; + comm->devrState.memHead = &mems[0]; + + void* const goodHostWin = reinterpret_cast(0xa1); + int goodRegistrations = 0; + int goodDeregistrations = 0; + ScopedHook connect(g_devrGinConnectOnce, [](ncclComm*) { return ncclSuccess; }); + ScopedHook reg(g_devrGinRegister, [&](ncclComm*, void* addr, size_t, void* hostWins[], ncclGinWindow_t[], int, bool, + int) { + // Registration only runs once connect has latched GIN. + EXPECT_TRUE(comm->devrState.ginEnabled); + if (addr != mems[0].primaryAddr) return ncclSystemError; + hostWins[0] = goodHostWin; + goodRegistrations++; + return ncclSuccess; + }); + ScopedHook dereg(g_devrGinDeregister, [&](ncclComm*, void* hostWins[]) { + if (hostWins[0] == goodHostWin) goodDeregistrations++; + return ncclSuccess; + }); + + EXPECT_EQ(Create(), ncclSystemError); + ASSERT_FALSE(comm->devrState.ginEnabled); + EXPECT_EQ(goodRegistrations, 1); + EXPECT_EQ(goodDeregistrations, 1); + EXPECT_EQ(mems[0].ginSegmentInfos, nullptr); + EXPECT_EQ(mems[1].ginSegmentInfos, nullptr); + + EXPECT_EQ(Create(), ncclSystemError); + EXPECT_EQ(connect.calls, 2); + EXPECT_EQ(goodRegistrations, 2); + EXPECT_EQ(goodDeregistrations, 2); + EXPECT_FALSE(comm->devrState.ginEnabled); + + comm->devrState.memHead = nullptr; + for (ncclDevrMemory& mem : mems) free(mem.ginSegmentInfos); +} + // Branch: a GIN request rejected after the support checks but before connect // leaves GIN unlatched, so a later valid create still connects. TEST_F(DevrCommCreateInternalTest, WorldGinBarrierOnRail_RejectedWithoutLatchingGin) {