Skip to content

fix(rccl): Don't gate hierarchical CE on hierarchical AllGather sub-communicators #12215

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
thomas-huber wants to merge 1 commit into from
Closed

fix(rccl): Don't gate hierarchical CE on hierarchical AllGather sub-communicators #12215

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
6 changes: 0 additions & 6 deletions projects/rccl/src/ce_coll.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1258,12 +1258,6 @@ bool ncclHierCeAvailable(struct ncclComm* comm, ncclFunc_t coll, int /*ncclDevRe
TRACE(NCCL_TUNING, "Skipping hierarchical CE collective: not multi-node");
return false;
}
// Sub-comms are only built at nNodes >= 8; rcclHierarchicalAlgoInfo dereferences
// them unconditionally, so bail out here if they are not initialized.
if (!comm->hierarchicalCommsInitialized) {
TRACE(NCCL_TUNING, "Skipping hierarchical CE collective: hierarchical sub-comms not initialized");
return false;
}
// If LSA already spans the whole comm, use CE path instead
if (ncclDevrIsOneLsaTeam(comm)) {
TRACE(NCCL_TUNING, "Skipping hierarchical CE collective: LSA spans the comm; use CE path instead");
Expand Down
17 changes: 4 additions & 13 deletions projects/rccl/src/rccl_wrap.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1676,17 +1676,13 @@ ncclResult_t rcclSelectAllGather(struct ncclComm* comm, const void* sendbuff, vo
// Branch #3.5: Hierarchical CE (multi-node, both buffers registered).
// ceCollTaskAppend routes to ncclHierCeAllGather via ncclHierCeDispatch(comm),
// so RCCL_CE_REGISTERED is correct here — same as rcclSelectAlltoAll Branch #5.
// It uses none of the hierarchical AllGather sub-communicators, which need not
// exist on this comm, so it is reported like the other CE branches.
const bool hierCeAvailable =
!ceCapturing && ncclHierCeAvailable(comm, ncclFuncAllGather, (int)ncclSum, datatype, winRegType, sendWin, recvWin);
if (hierCeAvailable && !hasSysmemSegment &&
(comm->config.CTAPolicy & NCCL_CTA_POLICY_ZERO)) {
decision->algo = RCCL_CE_REGISTERED;
if (query) {
int a, p, ch;
NCCLCHECK(rcclHierarchicalAlgoInfo(comm, ncclFuncAllGather, sendcount, datatype, &a, &p, &ch));
decision->protocol = p;
decision->nMaxChannels = ch;
}
return ncclSuccess;
}
}
Expand Down Expand Up @@ -2006,16 +2002,11 @@ ncclResult_t rcclSelectAlltoAll(struct ncclComm* comm, const void* sendbuff, voi
(int)ncclCeAvailable(comm, ncclFuncAlltoAll, ncclDevSum, datatype, a2aWinRegType, a2aSendWin, a2aRecvWin));

// (5) Hierarchical CE: multi-node, non-LSA-spanning.
// Require CTA_POLICY_ZERO and no sysmem segment, matching the AllGather twin.
// Require CTA_POLICY_ZERO and no sysmem segment, and report it like the other
// CE branches, matching the AllGather twin.
if ((comm->config.CTAPolicy & NCCL_CTA_POLICY_ZERO) && !a2aHasSysmem &&
ncclHierCeAvailable(comm, ncclFuncAlltoAll, ncclDevSum, datatype, a2aWinRegType, a2aSendWin, a2aRecvWin)) {
decision->algo = RCCL_CE_REGISTERED; // reports as CE; hier dispatch in taskAppend
if (query) {
int a, p, ch;
NCCLCHECK(rcclHierarchicalAlgoInfo(comm, ncclFuncAlltoAll, count, datatype, &a, &p, &ch));
decision->protocol = p;
decision->nMaxChannels = ch;
}
return ncclSuccess;
}

Expand Down
67 changes: 67 additions & 0 deletions projects/rccl/test/host/wrap-test.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4480,6 +4480,73 @@ TEST(WrapMicrotestIsolated, SelectAllGather_CeRegisteredViaSymmetricWindowsChose
});
}

// Hierarchical CE (branch #3.5) needs none of the hierarchical AllGather
// sub-communicators, which are only built on 8+ nodes. A 2-node comm has none,
// and both the report and the live decision must still name CE without touching
// them. The control proves the decision came from the hierarchical CE branch.
TEST(WrapMicrotestIsolated, SelectAllGather_HierCeWithoutHierarchicalSubComms) {
RUN_ISOLATED_TEST(
"Wrap_SelectAllGather_HierCeWithoutHierarchicalSubComms",
[]() {
g_hierCeAvailable = true;
ncclComm* comm = MakeCommWithArch("gfx942");
comm->nNodes = 2;
comm->nRanks = 16;
comm->config.CTAPolicy = NCCL_CTA_POLICY_ZERO;
ASSERT_EQ(nullptr, comm->hierarchicalInterComm);
ASSERT_EQ(nullptr, comm->hierarchicalIntraComm);
for (bool query : {true, false}) {
rcclCollDecision decision{};
EXPECT_EQ(ncclSuccess, rcclSelectAllGather(comm, nullptr, nullptr, /*sendcount=*/1024, ncclFloat32,
/*stream=*/nullptr, query, /*graphCapturingHint=*/false,
&decision));
EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_CE_REGISTERED, decision.algo) << "query=" << query;
EXPECT_EQ(NCCL_PROTO_SIMPLE, decision.protocol) << "query=" << query;
EXPECT_EQ(0, decision.nMaxChannels) << "query=" << query;
}

@amd-argus amd-argus Bot Sep 24, 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

The new tests set CTAPolicy to NCCL_CTA_POLICY_ZERO at line 4495 and never clear it, so nothing pins the comm->config.CTAPolicy & NCCL_CTA_POLICY_ZERO conjunct that branch #3.5 carries, and a change dropping it would still pass. The CE-registered test just above adds exactly this half at line 4472 and gives the reason. Could we add it here and in the AlltoAll twin at line 4527?

Suggested change
}
// ...and the CTA-policy conjunct itself, as branch #3's test does at line 4472.
comm->config.CTAPolicy = NCCL_CTA_POLICY_DEFAULT;
rcclCollDecision noPolicy{};
EXPECT_EQ(ncclSuccess, rcclSelectAllGather(comm, nullptr, nullptr, /*sendcount=*/1024, ncclFloat32,
/*stream=*/nullptr, /*query=*/true, /*graphCapturingHint=*/false,
&noPolicy));
EXPECT_NE((int)rcclAddonAlgos_t::RCCL_CE_REGISTERED, noPolicy.algo)
<< "hierarchical CE must require CTA_POLICY_ZERO";
comm->config.CTAPolicy = NCCL_CTA_POLICY_ZERO;

g_hierCeAvailable = false;
rcclCollDecision control{};
EXPECT_EQ(ncclSuccess, rcclSelectAllGather(comm, nullptr, nullptr, /*sendcount=*/1024, ncclFloat32,
/*stream=*/nullptr, /*query=*/true, /*graphCapturingHint=*/false,
&control));
EXPECT_NE((int)rcclAddonAlgos_t::RCCL_CE_REGISTERED, control.algo);
DeleteCommWithArch(comm);
});
}

// The AlltoAll twin of the test above, for rcclSelectAlltoAll branch (5).
TEST(WrapMicrotestIsolated, SelectAlltoAll_HierCeWithoutHierarchicalSubComms) {
RUN_ISOLATED_TEST(
"Wrap_SelectAlltoAll_HierCeWithoutHierarchicalSubComms",
[]() {
g_hierCeAvailable = true;
ncclComm* comm = MakeCommWithArch("gfx942");
comm->nNodes = 2;
comm->nRanks = 16;
comm->config.CTAPolicy = NCCL_CTA_POLICY_ZERO;
ASSERT_EQ(nullptr, comm->hierarchicalInterComm);
ASSERT_EQ(nullptr, comm->hierarchicalIntraComm);
for (bool query : {true, false}) {
rcclCollDecision decision{};
EXPECT_EQ(ncclSuccess, rcclSelectAlltoAll(comm, nullptr, nullptr, /*count=*/1024, ncclFloat32,
/*stream=*/nullptr, query, /*graphCapturingHint=*/false,
&decision));
EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_CE_REGISTERED, decision.algo) << "query=" << query;
EXPECT_EQ(NCCL_PROTO_SIMPLE, decision.protocol) << "query=" << query;
EXPECT_EQ(0, decision.nMaxChannels) << "query=" << query;
}

g_hierCeAvailable = false;
rcclCollDecision control{};
EXPECT_EQ(ncclSuccess, rcclSelectAlltoAll(comm, nullptr, nullptr, /*count=*/1024, ncclFloat32,
/*stream=*/nullptr, /*query=*/true, /*graphCapturingHint=*/false,
&control));
EXPECT_NE((int)rcclAddonAlgos_t::RCCL_CE_REGISTERED, control.algo);
DeleteCommWithArch(comm);
});
}

// Complementary proof for AllGather's own CE-registered check (a separate
// hasSysmemSegment computation from rcclSelectAllReduce's).
TEST(WrapMicrotestIsolated, SelectAllGather_SysmemSegmentBlocksCeRegistered) {
Expand Down
Loading