diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index b2e986363d47..0fd610a1769f 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,10 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom if (ginActivated) { NCCLCHECKGOTO(ncclGinConnectOnce(comm), ret, fail); + // 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. 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 1c3523cfe3d5..a2e9eab275ac 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); + ASSERT_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 8240ac137ef3..d6ba49a9eea4 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 dfc7a1c84e4b..cdba07c69180 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