Skip to content

fix(rccl): don't set hasMloPart for non-partitioned HIP GPUs #10817

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

fix(rccl): don't set hasMloPart for non-partitioned HIP GPUs #10817

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
20 changes: 13 additions & 7 deletions projects/rccl/src/init.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1097,12 +1097,18 @@ static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, u
deviceClass[0] = '\0';
if (int64ToBusId(info->busId, busIdStr) == ncclSuccess) {
(void)ncclOsGetPciDeviceClassByBusId(busIdStr, deviceClass, sizeof(deviceClass));
int isGpu = strncmp(deviceClass, PCI_ACCELERATOR_CLASS, strlen(PCI_ACCELERATOR_CLASS)) == 0 ||
strncmp(deviceClass, "0x03", 4) == 0;
if (fn == 0) {
if (isGpu) info->mloPart = 0;
} else if (!isGpu) {
info->mloPart = fn;
// Only set mloPart for fake HIP functions (fn > 0 that aren't GPUs in
// sysfs). fn=0 GPUs keep mloPart as NCCL_TOPO_UNDEF — we can't
// distinguish non-partitioned GPUs from CPX partition 0, and setting
// mloPart=0 would disable GIN and NET GDR on non-partitioned systems.
// If lookup failed (deviceClass empty), leave mloPart undefined to
// avoid misclassifying a real GPU at fn>0 as a partition.
if (fn > 0 && deviceClass[0] != '\0') {
int isGpu = strncmp(deviceClass, PCI_ACCELERATOR_CLASS, strlen(PCI_ACCELERATOR_CLASS)) == 0 ||
strncmp(deviceClass, "0x03", 4) == 0;
if (!isGpu) {
info->mloPart = fn;
}
}
}
}
Expand Down Expand Up @@ -1506,7 +1512,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
}
if (comm->peerInfo[i].hostHash != comm->peerInfo[rank].hostHash) nNodes++;
if (!comm->peerInfo[i].cuMemSupport) comm->cuMemSupport = 0;
if (comm->peerInfo[i].mloPart != -1) comm->hasMloPart = true;
if (comm->peerInfo[i].mloPart != NCCL_TOPO_UNDEF) comm->hasMloPart = true;
for (int j = 0; j < i; j++) {
// NVML device is agnostic to MloPart being used. With MloPart, each partition has a different GPU UUID.
comm->hasMultiRankNvml = (comm->peerInfo[i].hostHash == comm->peerInfo[j].hostHash) &&
Expand Down
11 changes: 11 additions & 0 deletions projects/rccl/test/host/init-test.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1938,6 +1938,17 @@ TEST_F(InitMicrotest, InitTransportsRank_NoPeerWithMloPart_LeavesHasMloPartUnset
EXPECT_FALSE(c.get()->hasMloPart);
}

// mloPart=1 is the lowest partition index for fake HIP functions (fn 1-7).
// Verify the boundary: mloPart=1 must trigger hasMloPart.
TEST_F(InitMicrotest, InitTransportsRank_MloPart1_SetsHasMloPart) {
TransportsRankComm c(/*nRanks=*/4, /*rank=*/0);
std::vector<PeerSpec> specs(4);
specs[3].mloPart = 1;
InstallPeerInfoAllGather(c, specs);
EXPECT_EQ(ncclRemoteError, initTransportsRank(c.get(), nullptr, c.timers()));
EXPECT_TRUE(c.get()->hasMloPart);
}

// NOT ASSERTABLE FROM THIS RUNG, deliberately: the four `global*Support` accumulators at :1491-1494 are
// function-locals first read at :2347-2363, ~700 lines past the terminator. They execute (so they count
// as covered) but nothing here can observe them, and deleting any of the four leaves the suite green.
Expand Down
Loading