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

Limit P2P channels on Rome

Browse Source 
[ROCm/rccl commit: 42955f5f4f]
This commit is contained in:
Wenkai Du
2020-09-17 16:01:45 -07:00
parent f6b08ca547
commit f0a303664e
5 changed files with 22 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/graph/search.cc
+6 -3
View File
@@ -892,14 +892,18 @@ static ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, char **str) {
*str = 0;
int ngpus = system->nodes[GPU].count;
int ncpus = system->nodes[CPU].count;
// 8 GPUs and 4 numa nodes only
if (ngpus != 8 || (ncpus != 4 && ncpus != 8))
// 8 GPUs only
if (ngpus != 8)
return ncclSuccess;
// only valid on Rome
int arch, vendor, model;
NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model));
if (arch != NCCL_TOPO_CPU_ARCH_X86 || vendor != NCCL_TOPO_CPU_VENDOR_AMD || model != NCCL_TOPO_CPU_TYPE_ROME)
return ncclSuccess;
system->type = RCCL_TOPO_4P2H_ROME;
// 4 or 8 numa nodes only
if (ncpus != 4 && ncpus != 8)
return ncclSuccess;
// number of GPUs and NICs on each numa node is used as first screening pattern
char pattern[256];
for (i = 0; i < ncpus; i++) {
@@ -998,7 +1002,6 @@ static ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, char **str) {
}
ringRemap[i] = 0;
*str = ringRemap;
system->type = RCCL_TOPO_4P2H_ROME;
INFO(NCCL_GRAPH, "Use 4P2H on Rome: %s", ringRemap);
return ncclSuccess;
}