From fb58c7312af496adeca12ba6b4a5fe98ac3cb733 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 | 32 ++++++++++++++++++- .../host/fakes/dev_runtime_micro_fakes.cc | 7 +++- .../test/host/fakes/dev_runtime_micro_fakes.h | 1 + 4 files changed, 41 insertions(+), 3 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index b2e986363d4..360befc9337 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -2060,7 +2060,6 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom } ginActivated = !devr->ginEnabled; - devr->ginEnabled = true; } if (reqs->worldGinBarrierCount > 0 && requestedConnectionType == NCCL_GIN_CONNECTION_RAIL) { @@ -2078,6 +2077,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 1c3523cfe3d..d8171f1275f 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -6711,7 +6711,8 @@ TEST(DevrGinContextsTotal, ZeroConnectionBackend_TreatedAsOne) { // resource-window sizing that follows it, up to the cuMemCreate call. // // The rest of the devcomm build -- symmetric-memory window creation, GIN -// window registration, barriers -- is not covered here. +// window registration, barriers -- is not covered here beyond a failed GIN +// connect. class DevrCommCreateInternalTest : public NcclVersionCompatTest { protected: @@ -6828,6 +6829,35 @@ TEST_F(DevrCommCreateInternalTest, GinSignalRegions_UseLargestRoundedContextCoun EXPECT_EQ(allocationSize, 96u); } +// 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 8240ac137ef..d6ba49a9eea 100644 --- a/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc +++ b/projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc @@ -305,7 +305,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; } @@ -563,6 +567,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 dfc7a1c84e4..cdba07c6918 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 From d77027c9e0b04f3b6853580ce616deaa4ac68f12 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 360befc9337..a7e9e4b42c3 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -2078,7 +2078,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 632498ed23090cdfc1167066f81e5e9320a1351e 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 a7e9e4b42c3..48ed6c0a1c9 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -2077,7 +2077,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 d8171f1275f..a2e9eab275a 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -6838,7 +6838,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 f6f5cb86ac1154c1b71e5430a14393f9178ed8d6 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Mon, 26 Oct 2026 15:07:10 -0400 Subject: [PATCH 4/4] fix(rccl): mark the post-connect GIN latch as an RCCL divergence Upstream sets ginEnabled before ncclGinConnectOnce. Tag the comment RCCL: so a future NCCL sync keeps the latch after connect. AICOMRCCL-2879 Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index 48ed6c0a1c9..0fd610a1769 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -2077,8 +2077,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 stale or 0, and - // a retry must reconnect rather than reach ncclGinDevCommSetup. + // RCCL: upstream latches ginEnabled before connect; keep it here on a sync. + // 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.