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

Allow setup ring through NCCL_RINGS to facilitate testing

Browse Source 
[ROCm/rccl commit: 5b03132ace]
This commit is contained in:
Wenkai Du
2020-08-04 21:07:00 +00:00
parent 22a6211eaf
commit 5f96f13e75
2 changed files with 7 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/graph/connect.cc
+1
View File
@@ -291,6 +291,7 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, struct nccl
int nc = nChannels*2;
if (comm->topo->nodes[NET].count == 0 && comm->topo->type == RCCL_TOPO_CR8G) nc = nChannels*3;
if (comm->topo->nodes[NET].count && comm->topo->type == RCCL_TOPO_4P2H_ROME) nc = 4*comm->topo->nodes[NET].count;
int end = std::min((int)ncclMaxNchannels(), std::max(nc, ncclMinNchannels()));
// Duplication should be complete now
projects/rccl/src/graph/search.cc
+6 -2
View File
@@ -931,6 +931,8 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph
if (graph->nChannels > 0) return ncclSuccess;
}
str = getenv("NCCL_RINGS");
if (str) system->type = RCCL_TOPO_4P2H_ROME;
if (!str) NCCLCHECK(parseChordalRing(system, &str));
if (!str) NCCLCHECK(parseRome4P2H(system, &str));
if (str) {
@@ -950,8 +952,10 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph
if (system->nodes[NET].count) {
// do not change ring order for multi node 4P2H on Rome
if (system->type == RCCL_TOPO_4P2H_ROME) {
for (int n = 0; n < graph->nChannels; n++)
graph->inter[n*2] = graph->inter[n*2+1] = n%system->nodes[NET].count;
for (int n = 0; n < graph->nChannels; n++) {
graph->inter[n*2] = n%system->nodes[NET].count;
graph->inter[n*2+1] = (n+1)%system->nodes[NET].count;
}
} else {
int *intra, *used;
graph->nChannels = system->nodes[NET].count;