Skip to content

fix(rccl): latch GIN only after a successful connect #14321

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
dlamd1dai merged 4 commits into from
Oct 27, 2026
Merged

fix(rccl): latch GIN only after a successful connect #14321

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
fb58c73
fix(rccl): latch GIN only after a successful connect
dlamd1dai Oct 10, 2026
d77027c
Update projects/rccl/src/dev_runtime.cc
abouteiller Oct 23, 2026
632498e
fix(rccl): drop the duplicated GIN latch comment and fail the retry t…
dlamd1dai Oct 26, 2026
f6f5cb8
fix(rccl): mark the post-connect GIN latch as an RCCL divergence
dlamd1dai Oct 26, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion projects/rccl/src/dev_runtime.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand All @@ -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.
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
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) {
Expand Down
32 changes: 31 additions & 1 deletion projects/rccl/test/host/dev-runtime-test.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down Expand Up @@ -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
Expand Down
7 changes: 6 additions & 1 deletion projects/rccl/test/host/fakes/dev_runtime_micro_fakes.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<ncclResult_t(struct ncclComm*)> 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;
}
Expand Down Expand Up @@ -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;
Expand Down
1 change: 1 addition & 0 deletions projects/rccl/test/host/fakes/dev_runtime_micro_fakes.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,6 +57,7 @@ extern std::function<ncclResult_t(void*, int*, int, int, void*, int)> g_devrBoot
extern std::function<ncclResult_t(void*, int, int, void*, int)> g_devrBootstrapSend;
extern std::function<ncclResult_t(void*, int, int, void*, int)> g_devrBootstrapRecv;

extern std::function<ncclResult_t(struct ncclComm*)> g_devrGinConnectOnce;
extern std::function<ncclResult_t(struct ncclComm*, void*, size_t, void*[NCCL_GIN_MAX_CONNECTIONS],
ncclGinWindow_t[NCCL_GIN_MAX_CONNECTIONS],
uint32_t[NCCL_GIN_MAX_ACTIVE_BACKENDS], int, bool, int)>
Expand Down
Loading