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

Trim NICs when all GPUs are connected by XGMI (#430)

Browse Source 
* Trim NICs when all GPUs are connected by XGMI

* Only enable clique with maximum of 2 hops
This commit is contained in:
Wenkai Du
2021-10-05 18:27:43 -07:00
committed by GitHub
parent 51a1cf428e
commit 29c729d8b6
6 changed files with 159 additions and 78 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/graph/paths.cc
+16 -2
View File
@@ -510,8 +510,22 @@ ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm*
}
if (ret) {
system->type |= RCCL_TOPO_GDR_ALL;
remove = 0;
INFO(NCCL_GRAPH, "GDR is available on all GPUs");
bool allXgmi = true;
// don't trim NICs unless all GPUs are connected by XGMI
for (int i = 0; i < system->nodes[GPU].count && allXgmi; i++) {
int cudaDev1 = system->nodes[GPU].nodes[i].gpu.dev;
for (int j = 0; j < system->nodes[GPU].count && allXgmi; j++) {
if (i == j) continue;
int cudaDev2 = system->nodes[GPU].nodes[j].gpu.dev;
bool isXGMI;
NCCLCHECK(ncclTopoGetLinkType(comm->topo, cudaDev1, cudaDev2, &isXGMI));
allXgmi &= isXGMI;
}
}
if (!allXgmi) {
remove = 0;
INFO(NCCL_GRAPH, "GDR is available on all GPUs");
}
}
}
comm->localRanks = system->nodes[GPU].count;
src/graph/search.cc
+34 -9
View File
@@ -1046,7 +1046,8 @@ ncclResult_t ncclTopoGetIntraNetDev(struct ncclTopoSystem* system, int rank, str
return ncclSuccess;
}
ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, int cudaDev2, bool* isXGMI, bool direct_only) {
ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, int cudaDev2, bool* isXGMI, int maxInter, int nInter, int *inter) {
int interGpus[MAX_XGMI_INTER_GPUS+1];
int ngpus = system->nodes[GPU].count;
*isXGMI = false;
// check for direct XGMI connection
@@ -1065,15 +1066,39 @@ ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, in
}
}
}
if (direct_only) return ncclSuccess;
// check if there is intermediate GPU that is connected to both
if (maxInter == 0) return ncclSuccess;
// check if there are intermediate GPUs that are connected to both
bool res1, res2, res3;
int j;
for (j=0; j<nInter; j++) {
bool res1;
ncclTopoGetLinkType(system, inter[j], inter[j+1], &res1, 0);
if (!res1) break;
}
if (j<nInter) return ncclSuccess;
if (nInter > 0 && inter != nullptr) {
ncclTopoGetLinkType(system, inter[nInter], cudaDev2, &res2, 0);
if (res2) {
*isXGMI = true;
return ncclSuccess;
}
memcpy(interGpus+1, inter+1, sizeof(int)*nInter);
}
interGpus[0] = cudaDev1;
// add one more intermediate GPU recursively util reaching max depth
nInter++;
if (nInter+2 > ngpus || nInter > MAX_XGMI_INTER_GPUS || nInter > maxInter) return ncclSuccess;
for (int i=0; i<ngpus; i++) {
if (system->nodes[GPU].nodes[i].gpu.dev == cudaDev1 || system->nodes[GPU].nodes[i].gpu.dev == cudaDev2)
continue;
bool res1, res2;
ncclTopoGetLinkType(system, system->nodes[GPU].nodes[i].gpu.dev, cudaDev1, &res1, true);
ncclTopoGetLinkType(system, system->nodes[GPU].nodes[i].gpu.dev, cudaDev2, &res2, true);
if (res1 && res2) {
int dev = system->nodes[GPU].nodes[i].gpu.dev;
// skip duplicated GPU
if (dev == cudaDev2) continue;
for (j=0; j<nInter; j++)
if (dev == interGpus[j]) break;
if (j<nInter) continue;
// check connectivity with intermediate GPUs
interGpus[nInter] = dev;
ncclTopoGetLinkType(system, cudaDev1, cudaDev2, &res3, maxInter, nInter, interGpus);
if (res3) {
*isXGMI = true;
return ncclSuccess;
}