Skip to content

fix(rccl): roll back a GIN activation that fails after connect #14231

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

Closed
dlamd1dai wants to merge 4 commits into from
Closed

fix(rccl): roll back a GIN activation that fails after connect #14231

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
bb5968d
fix(rccl): latch GIN only after a successful connect
dlamd1dai Oct 10, 2026
1da36b4
Update projects/rccl/src/dev_runtime.cc
abouteiller Oct 23, 2026
5527192
fix(rccl): drop the duplicated GIN latch comment and fail the retry t…
dlamd1dai Oct 26, 2026
ad48d9d
fix(rccl): roll back a GIN activation that fails after connect
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
46 changes: 45 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 @@ -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.

@amd-argus amd-argus Bot Oct 26, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Low

Could this header fit the one-line budget? The safety argument still fits:

Suggested change
// now came from this activation.
// Undo a partial GIN activation; memories and windows held no GIN state before it.

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<uintptr_t>(devr->lsaFlatBase);
uintptr_t addr = base + rank * devr->bigSize + mem->bigOffset;
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -1910,7 +1947,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 @@ -1928,6 +1964,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.
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 Expand Up @@ -2131,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
Expand Down Expand Up @@ -2207,8 +2247,12 @@ ncclResult_t ncclDevrCommCreateInternal(struct ncclComm* comm, struct ncclDevCom
}
symMemoryDropRef(comm, mem);
fail_stream:
if (ginActivated && !ginActivationDone) symGinActivationRollbackWindows(devr, stream);

@amd-argus amd-argus Bot Oct 26, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Medium

symGinActivationRollbackWindows has no test, so the suite passes with this call removed. The fixture stops early because the requirement-builder fakes at dev_runtime_micro_fakes.cc:498 leave the stack ncclDevResourceRequirements uninitialised, and line 2109 writes through outBufferHandle. Could those fakes zero their output? A test could then fail g_devrAllocAndPopulateSegmentWindows on a second GIN window. The first window's hook must return a shadow-pool object, since the default nullptr (dev_runtime_micro_fakes.cc:398) is skipped at line 838.

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;
}
Expand Down
89 changes: 88 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 @@ -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 activation.

class DevrCommCreateInternalTest : public NcclVersionCompatTest {
protected:
Expand Down Expand Up @@ -5581,6 +5581,93 @@ 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);
ASSERT_FALSE(comm->devrState.ginEnabled);

EXPECT_EQ(Create(), ncclSystemError);
EXPECT_EQ(connect.calls, 2);
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<hipMemGenericAllocationHandle_t>(0x55),
reinterpret_cast<hipMemGenericAllocationHandle_t>(0x56)};
ncclDevrMemory mems[2] = {};
for (int m = 0; m < 2; m++) {
mems[m].primaryAddr = reinterpret_cast<void*>(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<void*>(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) {
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 @@ -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<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 @@ -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;
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], int, bool, int)>
g_devrGinRegister;
Expand Down
Loading