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

Rework Rome detection and add multiple network ports models (#274)

Browse Source 
* Rework Rome detection and add multiple network ports models

* Remove unused opCount in p2p transport

[ROCm/rccl commit: ae008fd2db]
This commit is contained in:
Wenkai Du
2020-10-07 13:37:36 -07:00
committed by GitHub
parent e12db6f2ab
commit 41260bb948
11 changed files with 616 additions and 352 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/init.cc
+16 -1
View File
@@ -849,6 +849,21 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
}
INFO(NCCL_INIT, "RCCL AllToAll(v)/Scatter/Gather kernels %s", comm->alltoallDisable ? "disabled" : "enabled");
// count NETs used by ring
int nNets = 0;
int nets[MAXCHANNELS*2];
for (int i = 0; i < ringGraph.nChannels; i++) {
for (int j = 0; j < 2; j++) {
int k;
for (k = 0; k < nNets; k++)
if (nets[k] == ringGraph.inter[i*2+j]) break;
if (k >= nNets) {
nets[nNets] = ringGraph.inter[i*2+j];
nNets++;
}
}
}
if (comm->nChannels < nChannelsOrig) {
// We started duplicating channels during Preset(), so we need to move the
// duplicated channels since we have removed some.
@@ -858,7 +873,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
int *rings;
NCCLCHECK(ncclCalloc(&rings, nranks*MAXCHANNELS));
NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, allTopoRanks, rings, gcn));
NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, allTopoRanks, rings, gcn, nNets));
if (comm->nNodes > 1 &&
ncclParamCollNetEnable() == 1 &&
collNetSupport() && collNetGraph.nChannels) {