AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Consistent channel shuffling for MI300X multi-node (#1255)

Browse Source 
* Revert "[GRAPH] Use channel shuffling only for IB systems (#1228)"

This reverts commit 8a7dd0e590.

* Revert "Revert "Changing channel stride for MI300X multinode (#1196)" (#1224)"

This reverts commit bb6fab3d8e.

[ROCm/rccl commit: a1ef217b32]
This commit is contained in:
Nilesh M Negi
2024-07-18 10:18:09 -05:00
committed by GitHub
parent 13134c6c64
commit 73e17b3e70
2 changed files with 5 additions and 166 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/graph/connect.cc
+4 -22
View File
@@ -624,29 +624,11 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa
NCCLCHECK(connectRings(comm, ringRecv, ringSend, ringPrev, ringNext));
NCCLCHECK(connectTrees(comm, treeToParent, treeToChild0, treeToChild1, treePatterns));
// 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;
}
// Only use full MAXCHANNELS for gfx94x
int maxChannels = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ? MAXCHANNELS : 2*CHANNEL_LIMIT;
if (graphs[NCCL_ALGO_RING]->nIntraChannels > 0 || comm->nNodes > 1) {
maxChannels = std::min(multiNodeNchannels, maxChannels);
maxChannels = std::min(64, maxChannels);
}
// Duplicate ringPrev/ringNext for ncclBuildRing
@@ -692,7 +674,7 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa
int minNchannels = ncclMinNchannels();
if (comm->nNodes > 1) {
minNchannels = std::min(multiNodeNchannels, minNchannels);
minNchannels = std::min(64, minNchannels);
}
if (comm->nRanks < 8 && 64 < minNchannels) {
minNchannels = 2;