From d5f90e19b54a85d65124e254a7221686688f362a Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Tue, 21 Jul 2020 14:10:36 -0700 Subject: [PATCH] Add 8P6L multi-node models (#239) --- src/graph/connect.cc | 6 +- src/graph/paths.cc | 6 + src/graph/search.cc | 50 +++++--- src/graph/topo.cc | 6 + src/graph/topo.h | 3 + src/include/graph.h | 4 +- tools/scripts/topo_val.sh | 2 +- tools/topo_expl/models/topo_8p6l.xml | 20 ++-- tools/topo_expl/models/topo_8p6l_1nic.xml | 107 ++++++++++++++++++ tools/topo_expl/models/topo_8p6l_2nic.xml | 112 ++++++++++++++++++ tools/topo_expl/models/topo_8p6l_3nic.xml | 117 +++++++++++++++++++ tools/topo_expl/models/topo_8p6l_4nic.xml | 122 ++++++++++++++++++++ tools/topo_expl/models/topo_8p6l_5nic.xml | 127 +++++++++++++++++++++ tools/topo_expl/models/topo_8p6l_6nic.xml | 132 ++++++++++++++++++++++ tools/topo_expl/topo_expl.cpp | 48 ++++++++ 15 files changed, 830 insertions(+), 32 deletions(-) create mode 100644 tools/topo_expl/models/topo_8p6l_1nic.xml create mode 100644 tools/topo_expl/models/topo_8p6l_2nic.xml create mode 100644 tools/topo_expl/models/topo_8p6l_3nic.xml create mode 100644 tools/topo_expl/models/topo_8p6l_4nic.xml create mode 100644 tools/topo_expl/models/topo_8p6l_5nic.xml create mode 100644 tools/topo_expl/models/topo_8p6l_6nic.xml diff --git a/src/graph/connect.cc b/src/graph/connect.cc index 084a98bc22..44955dab48 100644 --- a/src/graph/connect.cc +++ b/src/graph/connect.cc @@ -9,6 +9,7 @@ #include "graph.h" #include "trees.h" #include "rings.h" +#include "topo.h" /******************************************************************/ /********************* Internode connection ***********************/ @@ -288,9 +289,8 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, struct nccl memcpy(ringPrev+nChannels*nranks, ringPrev, nChannels*nranks*sizeof(int)); memcpy(ringNext+nChannels*nranks, ringNext, nChannels*nranks*sizeof(int)); - char *str = NULL; - NCCLCHECK(parseChordalRing(comm->topo, &str)); - int end = std::min((int)ncclMaxNchannels(), (str ? nChannels*3 : ncclMinNchannels())); + int end = std::min((int)ncclMaxNchannels(), + ((comm->topo->nodes[NET].count == 0 && comm->topo->type == RCCL_TOPO_CR8G) ? nChannels*3 : ncclMinNchannels())); // Duplication should be complete now nChannels = comm->nChannels = std::min(MAXCHANNELS,nChannels*2); diff --git a/src/graph/paths.cc b/src/graph/paths.cc index e2e2d396b6..10c729647a 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -329,6 +329,12 @@ ncclResult_t ncclTopoCheckGdr(struct ncclTopoSystem* system, int64_t busId, int // Check if we are close enough that it makes sense to enable GDR int netGdrLevel = PATH_PXB; +#ifdef TOPO_EXPL + 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) + netGdrLevel = PATH_PHB; +#endif NCCLCHECK(ncclGetLevel(&ncclTopoUserGdrLevel, NULL, "NCCL_NET_GDR_LEVEL")); if (ncclTopoUserGdrLevel != -2) netGdrLevel = ncclTopoUserGdrLevel; int distance = gpu->paths[NET][n].type; diff --git a/src/graph/search.cc b/src/graph/search.cc index 4b35d7cf21..18e9ee793e 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -25,6 +25,7 @@ static float getMaxWidth(struct ncclTopoSystem* system, struct ncclTopoNode* gpu } ncclResult_t ncclTopoSearchInit(struct ncclTopoSystem* system) { system->maxWidth = 0.0; + system->type = 0; int inter = system->nodes[NET].count; if (inter == 0 && system->nodes[GPU].count == 1) { system->maxWidth = LOC_WIDTH; @@ -714,30 +715,30 @@ end: return ncclSuccess; } -ncclResult_t parseChordalRing(struct ncclTopoSystem* system, char **str) { - static const char *ringBase = "0 6 7 4 5 3 2 1|0 5 6 3 7 1 4 2|0 4 6 2 7 5 1 3|0 1 2 3 5 4 7 6|0 2 4 1 7 3 6 5|0 3 1 5 7 2 6 4"; +static ncclResult_t parseChordalRing(struct ncclTopoSystem* system, char **str) { + static const char *ringBase = "0 1 2 3 5 4 7 6|0 2 4 1 7 3 6 5|0 3 1 5 7 2 6 4|0 6 7 4 5 3 2 1|0 5 6 3 7 1 4 2|0 4 6 2 7 5 1 3"; static char ringRemap[256]; int id[8], dist[8]; int i; + *str = 0; int ngpus = system->nodes[GPU].count; - // single node CR8G only - if (ngpus != 8 || system->nodes[NET].count != 0) + if (ngpus != 8) return ncclSuccess; // validate chordal ring and calculate distance for (i=0; inodes[GPU].nodes+i; if (node->paths[GPU] == NULL) continue; - int sum = ngpus*(ngpus-1)/2 - node->gpu.rank; + int sum = ngpus*(ngpus-1)/2 - node->gpu.dev; int count = 0; for (int n = 0; nlinks; link->remNode; link++) { - if (link->remNode->gpu.rank == n) break; + if (link->remNode->gpu.dev == n) break; } if (!link->remNode) continue; if (link->type != LINK_NVL) continue; - sum -= system->nodes[GPU].nodes[n].gpu.rank; + sum -= system->nodes[GPU].nodes[n].gpu.dev; count ++; } if(count != ngpus-2 || sum < 0 || sum > ngpus-1) { @@ -765,6 +766,7 @@ ncclResult_t parseChordalRing(struct ncclTopoSystem* system, char **str) { } ringRemap[i] = 0; *str = ringRemap; + system->type = RCCL_TOPO_CR8G; INFO(NCCL_GRAPH, "Use chordal ring: %s", ringRemap); return ncclSuccess; } @@ -806,18 +808,34 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph // Translate gpu numbers into ranks graph->intra[i] = system->nodes[GPU].nodes[graph->intra[i]].gpu.rank; } - // TODO : let user specify NICs - graph->inter[0] = graph->inter[1] = 0; graph->speedIntra = graph->speedInter = system->maxWidth; - if (graph->pattern == NCCL_TOPO_PATTERN_RING) { - // Reverse the loop - for (int c=0; cnChannels; c++) { - for (int i=0; i<=ngpus/2; i++) { - int tmp = graph->intra[ngpus*c+i]; - graph->intra[ngpus*c+i] = graph->intra[ngpus*c+(ngpus-i)%ngpus]; - graph->intra[ngpus*c+ngpus-i] = tmp; + if (system->nodes[NET].count) { + int *intra, *used; + graph->nChannels = system->nodes[NET].count; + NCCLCHECK(ncclCalloc(&intra, ngpus)); + NCCLCHECK(ncclCalloc(&used,system->nodes[NET].count)); + for (int n = 0; n < system->nodes[NET].count; n++) { + graph->inter[n*2] = graph->inter[n*2+1] = n; + struct ncclTopoNode* net = system->nodes[NET].nodes+n; + struct ncclTopoLinkList* paths = net->paths[GPU]; + // find the first unsed GPU that is closest to NIC + int f, m; + for (f = 0; f < ngpus; f++) { + int j = 0; for (j = 0; j < n; j++) if(used[j] == system->nodes[GPU].nodes[f].gpu.rank) break; + if(j >= n) break; } + for (int i = 0; i < ngpus; i++) { + int j = 0; for (j = 0; j < n; j++) if(used[j] == system->nodes[GPU].nodes[i].gpu.rank) break; + if (j < n) continue; + if (paths[i].count < paths[f].count) f = i; + } + for (m = 0; mintra[n*ngpus+m] == system->nodes[GPU].nodes[f].gpu.rank) break; + used[n] = graph->intra[n*ngpus+m]; + for (int i = 0; i < ngpus; i++) intra[i] = graph->intra[n*ngpus+((i+m)%ngpus)]; + for (int i = 0; i < ngpus; i++) graph->intra[n*ngpus+i] = intra[i]; } + free(used); + free(intra); } if (graph->nChannels) return ncclSuccess; } diff --git a/src/graph/topo.cc b/src/graph/topo.cc index ea53d90324..514a897900 100644 --- a/src/graph/topo.cc +++ b/src/graph/topo.cc @@ -400,6 +400,12 @@ ncclResult_t ncclTopoAddCpu(struct ncclXmlNode* xmlCpu, struct ncclTopoSystem* s NCCLCHECK(xmlGetAttrInt(xmlCpu, "modelid", &modelId)); cpu->cpu.model = (familyId == 6 && modelId >= 0x55) ? NCCL_TOPO_CPU_TYPE_SKL : NCCL_TOPO_CPU_INTEL_BDW; } + if (cpu->cpu.vendor == NCCL_TOPO_CPU_VENDOR_AMD) { + int familyId, modelId; + NCCLCHECK(xmlGetAttrInt(xmlCpu, "familyid", &familyId)); + NCCLCHECK(xmlGetAttrInt(xmlCpu, "modelid", &modelId)); + cpu->cpu.model = (familyId == 143 && modelId >= 49) ? NCCL_TOPO_CPU_TYPE_ROME : NCCL_TOPO_CPU_TYPE_ZEN; + } } for (int s=0; snSubs; s++) { struct ncclXmlNode* node = xmlCpu->subs[s]; diff --git a/src/graph/topo.h b/src/graph/topo.h index cd8bf801f9..6443ce36e1 100644 --- a/src/graph/topo.h +++ b/src/graph/topo.h @@ -77,6 +77,8 @@ struct ncclTopoLinkList { #define NCCL_TOPO_UNDEF (-1) +#define RCCL_TOPO_CR8G 1 + struct ncclTopoNode { int type; int64_t id; @@ -119,6 +121,7 @@ struct ncclTopoNodeSet { struct ncclTopoSystem { struct ncclTopoNodeSet nodes[NCCL_TOPO_NODE_TYPES]; float maxWidth; + int type; }; ncclResult_t ncclTopoGetNode(struct ncclTopoSystem* system, struct ncclTopoNode** node, int type, uint64_t id); diff --git a/src/include/graph.h b/src/include/graph.h index f4d623584b..b326c71b63 100644 --- a/src/include/graph.h +++ b/src/include/graph.h @@ -42,6 +42,8 @@ ncclResult_t ncclTopoSetAffinity(struct ncclTopoSystem* system, int rank); #define NCCL_TOPO_CPU_VENDOR_AMD 2 #define NCCL_TOPO_CPU_TYPE_BDW 1 #define NCCL_TOPO_CPU_TYPE_SKL 2 +#define NCCL_TOPO_CPU_TYPE_ZEN 3 +#define NCCL_TOPO_CPU_TYPE_ROME 4 ncclResult_t ncclTopoCpuType(struct ncclTopoSystem* system, int* arch, int* vendor, int* model); #define NCCL_TOPO_MAX_NODES 256 @@ -101,6 +103,4 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom #include "info.h" ncclResult_t ncclTopoGetAlgoTime(struct ncclInfo* info, int algorithm, int protocol, float* time); -ncclResult_t parseChordalRing(struct ncclTopoSystem* system, char **str); - #endif diff --git a/tools/scripts/topo_val.sh b/tools/scripts/topo_val.sh index 02273a3d1d..9a6b475a42 100755 --- a/tools/scripts/topo_val.sh +++ b/tools/scripts/topo_val.sh @@ -21,7 +21,7 @@ DIR="$(cd -P "$(dirname "${BASH_SOURCE[0]}")" && pwd)" -for i in {0..13} +for i in {0..20} do $DIR/../topo_expl/topo_expl -m $i > "topo_m$i.log" $DIR/../TopoVisual/topo_visual.sh -i "topo_m$i.log" diff --git a/tools/topo_expl/models/topo_8p6l.xml b/tools/topo_expl/models/topo_8p6l.xml index e54ac50a1b..6afbc5ad42 100644 --- a/tools/topo_expl/models/topo_8p6l.xml +++ b/tools/topo_expl/models/topo_8p6l.xml @@ -2,7 +2,7 @@ - + @@ -14,7 +14,7 @@ - + @@ -26,7 +26,7 @@ - + @@ -38,7 +38,7 @@ - + @@ -52,7 +52,7 @@ - + @@ -64,7 +64,7 @@ - + @@ -76,7 +76,7 @@ - + @@ -88,7 +88,7 @@ - + @@ -98,9 +98,9 @@ - + - + diff --git a/tools/topo_expl/models/topo_8p6l_1nic.xml b/tools/topo_expl/models/topo_8p6l_1nic.xml new file mode 100644 index 0000000000..92a08256f0 --- /dev/null +++ b/tools/topo_expl/models/topo_8p6l_1nic.xml @@ -0,0 +1,107 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/models/topo_8p6l_2nic.xml b/tools/topo_expl/models/topo_8p6l_2nic.xml new file mode 100644 index 0000000000..f4e66bc110 --- /dev/null +++ b/tools/topo_expl/models/topo_8p6l_2nic.xml @@ -0,0 +1,112 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/models/topo_8p6l_3nic.xml b/tools/topo_expl/models/topo_8p6l_3nic.xml new file mode 100644 index 0000000000..2cc7800c22 --- /dev/null +++ b/tools/topo_expl/models/topo_8p6l_3nic.xml @@ -0,0 +1,117 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/models/topo_8p6l_4nic.xml b/tools/topo_expl/models/topo_8p6l_4nic.xml new file mode 100644 index 0000000000..554d2ea3e2 --- /dev/null +++ b/tools/topo_expl/models/topo_8p6l_4nic.xml @@ -0,0 +1,122 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/models/topo_8p6l_5nic.xml b/tools/topo_expl/models/topo_8p6l_5nic.xml new file mode 100644 index 0000000000..c9ffda03a0 --- /dev/null +++ b/tools/topo_expl/models/topo_8p6l_5nic.xml @@ -0,0 +1,127 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/models/topo_8p6l_6nic.xml b/tools/topo_expl/models/topo_8p6l_6nic.xml new file mode 100644 index 0000000000..338c6876ce --- /dev/null +++ b/tools/topo_expl/models/topo_8p6l_6nic.xml @@ -0,0 +1,132 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp index 1f87fa0646..92e6b9fe9e 100644 --- a/tools/topo_expl/topo_expl.cpp +++ b/tools/topo_expl/topo_expl.cpp @@ -77,6 +77,12 @@ const char *model_descriptions[] = { "4 nodes with 8 VEGA20 GPUs XGMI 4P2H 1 NIC 2nd Hive", "4 nodes with 8 VEGA20 GPUs XGMI 4P2H 2 NIC", "single node 8 VEGA20 Rome", + "4 nodes gfx908 8P6L 1 NIC", + "4 nodes gfx908 8P6L 2 NICs", + "4 nodes gfx908 8P6L 3 NICs", + "4 nodes gfx908 8P6L 4 NICs", + "4 nodes gfx908 8P6L 5 NICs", + "4 nodes gfx908 8P6L 6 NICs", NULL, }; @@ -171,6 +177,48 @@ int main(int argc,char* argv[]) node = new NodeModel("topo_8p_rome.xml"); network.AddNode(node); break; + case 14: + for (int i=0; i<4; i++) { + node = new NodeModel("topo_8p6l.xml"); + network.AddNode(node); + } + break; + case 15: + for (int i=0; i<4; i++) { + node = new NodeModel("topo_8p6l_1nic.xml"); + network.AddNode(node); + } + break; + case 16: + for (int i=0; i<4; i++) { + node = new NodeModel("topo_8p6l_2nic.xml"); + network.AddNode(node); + } + break; + case 17: + for (int i=0; i<4; i++) { + node = new NodeModel("topo_8p6l_3nic.xml"); + network.AddNode(node); + } + break; + case 18: + for (int i=0; i<4; i++) { + node = new NodeModel("topo_8p6l_4nic.xml"); + network.AddNode(node); + } + break; + case 19: + for (int i=0; i<4; i++) { + node = new NodeModel("topo_8p6l_5nic.xml"); + network.AddNode(node); + } + break; + case 20: + for (int i=0; i<4; i++) { + node = new NodeModel("topo_8p6l_6nic.xml"); + network.AddNode(node); + } + break; default: printf("Invalid model_id %d\n", model_id); exit(0);