Skip to content

[GRAPH] Use channel shuffling only for IB systems #1228

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
nileshnegi merged 3 commits into from
Jul 2, 2024
Merged

[GRAPH] Use channel shuffling only for IB systems #1228

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
a3702fb
[GRAPH] Use channel shuffling only for IB systems
nileshnegi Jun 26, 2024
120f93f
[GRAPH] Define channels=48 for gfx94 RoCE systems
nileshnegi Jun 28, 2024
776a9e2
[GRAPH] Increase channels for RoCE gfx94 systems
nileshnegi Jul 1, 2024
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
26 changes: 22 additions & 4 deletions src/graph/connect.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -624,11 +624,29 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa
NCCLCHECK(connectRings(comm, ringRecv, ringSend, ringPrev, ringNext));
NCCLCHECK(connectTrees(comm, treeToParent, treeToChild0, treeToChild1, treePatterns));

// Only use full MAXCHANNELS for gfx94x
int maxChannels = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ? MAXCHANNELS : 2*CHANNEL_LIMIT;
// Define channels for non-gfx94 GPU architectures
int maxChannels = 2*CHANNEL_LIMIT;
int multiNodeNchannels = maxChannels;

// Define channels for gfx94 GPU architectures
if (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94")) {
// Only use full MAXCHANNELS for gfx94x
maxChannels = MAXCHANNELS;

// Define channels=64 for gfx94 multi-node systems
multiNodeNchannels = 64;

// Check if NCCL_IB_GID_INDEX=3 -- needed for RoCE systems
const char* ncclIbGidIndex = ncclGetEnv("NCCL_IB_GID_INDEX");
int gid_index = 0;
if (ncclIbGidIndex) gid_index = atoi(ncclIbGidIndex);

// Limit channels=48 for RoCE gfx94 multi-node systems
multiNodeNchannels = gid_index == 3 ? 48 : multiNodeNchannels;
}

if (graphs[NCCL_ALGO_RING]->nIntraChannels > 0 || comm->nNodes > 1) {
maxChannels = std::min(64, maxChannels);
maxChannels = std::min(multiNodeNchannels, maxChannels);
}

// Duplicate ringPrev/ringNext for ncclBuildRing
Expand Down Expand Up @@ -674,7 +692,7 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa

int minNchannels = ncclMinNchannels();
if (comm->nNodes > 1) {
minNchannels = std::min(64, minNchannels);
minNchannels = std::min(multiNodeNchannels, minNchannels);
}
if (comm->nRanks < 8 && 64 < minNchannels) {
minNchannels = 2;
Expand Down
Loading