From d24df908826b1e0813b555e8caf5a082f1416f5f Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Mon, 22 Feb 2021 16:46:45 -0500 Subject: [PATCH 1/4] Match NBIO only when GPUs and NICs are directly connected to CPU [ROCm/rccl commit: ec8d89b1ddd8a3b257c8bbee57b8706e9ff797ee] --- projects/rccl/src/graph/paths.cc | 15 ++++++++-- projects/rccl/src/graph/rome_models.h | 2 +- projects/rccl/src/graph/search.cc | 42 +++++++++++++++++---------- 3 files changed, 40 insertions(+), 19 deletions(-) diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc index 13a6434dd4..8475bf1ef0 100644 --- a/projects/rccl/src/graph/paths.cc +++ b/projects/rccl/src/graph/paths.cc @@ -342,8 +342,19 @@ ncclResult_t ncclTopoCheckGdr(struct ncclTopoSystem* system, int64_t busId, int else { int arch, vendor, model; NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model)); - if((system->nodes[GPU].nodes[g].id & 0xf0000) == (system->nodes[NET].nodes[n].net.busId & 0xf0000)) - netGdrLevel = PATH_PHB; + if (arch == NCCL_TOPO_CPU_ARCH_X86 && vendor == NCCL_TOPO_CPU_VENDOR_AMD && model == NCCL_TOPO_CPU_TYPE_ROME) { + int i, d1 = -1, d2 = -1; + for (i = 0; i < system->nodes[CPU].count; i++) + if (system->nodes[GPU].nodes[g].paths[CPU][i].count == 2) break; + if (i nodes[CPU].count) d1 = system->nodes[CPU].nodes[i].id; + for (i = 0; i < system->nodes[CPU].count; i++) + if (system->nodes[NET].nodes[n].paths[CPU][i].count == 2) break; + if (i nodes[CPU].count) d2 = system->nodes[CPU].nodes[i].id; + if (d1 != -1 && d2 != -1 && d1 == d2 && + (system->nodes[GPU].nodes[g].id & 0xf0000) == (system->nodes[NET].nodes[n].net.busId & 0xf0000)) { + netGdrLevel = PATH_PHB; + } + } } int distance = gpu->paths[NET][n].type; diff --git a/projects/rccl/src/graph/rome_models.h b/projects/rccl/src/graph/rome_models.h index 21b089c0b4..556e01977e 100755 --- a/projects/rccl/src/graph/rome_models.h +++ b/projects/rccl/src/graph/rome_models.h @@ -254,7 +254,7 @@ static struct rcclRomeModel rome_model_49 = { .nicNuma = { 0, 1, 2, 3, }, .connMatrix = { 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, }, .pattern = "21212121", - .ringBase = "N0 0 1 2 3 4 5 6 7 N3|N3 7 6 5 4 3 2 1 0 N0", + .ringBase = "N0 0 1 2 3 4 5 6 7 N3|N3 7 6 5 4 3 2 1 0 N0|N1 2 3 0 1 6 7 4 5 N2|N2 5 4 7 6 1 0 3 2 N1", }; static struct rcclRomeModel romeTopoModels[] = { diff --git a/projects/rccl/src/graph/search.cc b/projects/rccl/src/graph/search.cc index 65eb5040f3..cb34d515b6 100644 --- a/projects/rccl/src/graph/search.cc +++ b/projects/rccl/src/graph/search.cc @@ -991,7 +991,7 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo return ncclSuccess; } -static bool permuteGpuIds(int *g, int n, int last, struct rcclRomeModel* ref, struct rcclRomeModel* topo, int* time) { +static bool permuteGpuIds(int *g, int n, int last, struct rcclRomeModel* ref, struct rcclRomeModel* topo, int* time, bool nbio) { (*time) ++; if (n == last) { int i, j; @@ -1009,22 +1009,24 @@ static bool permuteGpuIds(int *g, int n, int last, struct rcclRomeModel* ref, st } if (i < ref->nGpus) return false; // match NBIO - for (i = 0; i < ref->nGpus; i++) { - for (j = 0; j < ref->nGpus; j++) { - if (i == j) continue; - bool nbio_ref = (ref->gpuIds[i]&0xf0000) == (ref->gpuIds[j]&0xf0000); - bool nbio_topo = (topo->gpuIds[g[i]]&0xf0000) == (topo->gpuIds[g[j]]&0xf0000); - if (nbio_ref != nbio_topo) break; - if (nbio_ref && ((ref->gpuIds[i]-ref->gpuIds[j])*(topo->gpuIds[g[i]]-topo->gpuIds[g[j]]) < 0)) break; + if (nbio) { + for (i = 0; i < ref->nGpus; i++) { + for (j = 0; j < ref->nGpus; j++) { + if (i == j) continue; + bool nbio_ref = (ref->gpuIds[i]&0xf0000) == (ref->gpuIds[j]&0xf0000); + bool nbio_topo = (topo->gpuIds[g[i]]&0xf0000) == (topo->gpuIds[g[j]]&0xf0000); + if (nbio_ref != nbio_topo) break; + if (nbio_ref && ((ref->gpuIds[i]-ref->gpuIds[j])*(topo->gpuIds[g[i]]-topo->gpuIds[g[j]]) < 0)) break; + } + if (j < ref->nGpus) break; } - if (j < ref->nGpus) break; + if (i < ref->nGpus) return false; } - if (i < ref->nGpus) return false; return true; } else { for (int i = n; i <= last; i++) { std::swap(g[n], g[i]); - if (permuteGpuIds(g, n+1, last, ref, topo, time)) return true; + if (permuteGpuIds(g, n+1, last, ref, topo, time, nbio)) return true; std::swap(g[n], g[i]); } } @@ -1057,12 +1059,23 @@ static ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopo int time = 0; struct timeval tvs, tve; gettimeofday(&tvs, NULL); + + // check if GPUs are directly connected to CPU + bool match_nbio = true; + for (i = 0; i < romeTopo.nGpus; i++) { + int cpu, gpu; + NCCLCHECK(ncclTopoIdToIndex(system, CPU, romeTopo.gpuNuma[i], &cpu)); + NCCLCHECK(ncclTopoIdToIndex(system, GPU, romeTopo.gpuIds[i], &gpu)); + if (system->nodes[GPU].nodes[gpu].paths[CPU][cpu].count > 2) break; + } + if (i < romeTopo.nGpus) match_nbio = false; + for (i = 0; i < sizeof(romeTopoModels)/sizeof(romeTopoModels[0]); i++) { if (romeTopo.nCpus != romeTopoModels[i].nCpus || romeTopo.nGpus != romeTopoModels[i].nGpus || romeTopo.nNics != romeTopoModels[i].nNics || romeTopo.nLinks != romeTopoModels[i].nLinks) continue; if (strcmp(romeTopoModels[i].pattern, pattern)) continue; for (int j = 0; j < ngpus; j++) g[j] = (j+2)%ngpus; - if (permuteGpuIds(g, 0, ngpus-1, romeTopoModels+i, &romeTopo, &time)) break; + if (permuteGpuIds(g, 0, ngpus-1, romeTopoModels+i, &romeTopo, &time, match_nbio)) break; } gettimeofday(&tve, NULL); float t = (tve.tv_sec - tvs.tv_sec)*1E3 + (tve.tv_usec - tvs.tv_usec)/1E3; @@ -1072,11 +1085,8 @@ static ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopo } char line[1024]; -#ifdef ENABLE_TRACE - sprintf(line, "Found matching Rome model index %d in %.2fms (%d iter) with GPU mapping: ", i, t, time); -#else + //sprintf(line, "Found matching Rome model index %d in %.2fms (%d iter) with GPU mapping: ", i, t, time); sprintf(line, "Found matching Rome model index %d with GPU mapping: ", i); -#endif int offset = strlen(line); for (int k = 0; k < ngpus; k++) { sprintf(line+offset, "%d ", g[k]); From 020ccb6730d0dcf5a1f9c4e59781e3c146b1edb5 Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Tue, 23 Feb 2021 21:41:26 +0000 Subject: [PATCH 2/4] Update tuning parameters for XGMI and NET [ROCm/rccl commit: e820a943e9dd9f4b98778a0e63e1db97a45e680e] --- projects/rccl/src/graph/connect.cc | 2 +- projects/rccl/src/graph/tuning.cc | 7 +++---- 2 files changed, 4 insertions(+), 5 deletions(-) diff --git a/projects/rccl/src/graph/connect.cc b/projects/rccl/src/graph/connect.cc index 737f3218e8..ce60f41121 100644 --- a/projects/rccl/src/graph/connect.cc +++ b/projects/rccl/src/graph/connect.cc @@ -269,7 +269,7 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa if (gcn == 908) nc = std::max(nc, 4); if (comm->topo->nodes[NET].count == 0 && comm->topo->type == RCCL_TOPO_CR8G) nc = nChannels*4; if (!nnets) nnets = comm->topo->nodes[NET].count; - if (nnets && comm->topo->type == RCCL_TOPO_4P2H_ROME) nc = 4*nnets; + if (nnets && comm->topo->type == RCCL_TOPO_4P2H_ROME) nc = (nnets > 3 ? 2 : 4)*nnets; int end = std::min((int)ncclMaxNchannels(), std::max(nc, ncclMinNchannels())); // Duplication should be complete now diff --git a/projects/rccl/src/graph/tuning.cc b/projects/rccl/src/graph/tuning.cc index dd56f8f34c..7982c41ee1 100644 --- a/projects/rccl/src/graph/tuning.cc +++ b/projects/rccl/src/graph/tuning.cc @@ -54,7 +54,7 @@ ncclResult_t parseList(const char* str, const char* elems[], int nelems, int* li // Latencies in us, Bandwidths in GB/s // Tree { LL, LL128, Simple } , Ring { LL, LL128, Simple } -static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 37.9, 37.9, 40.4 }, { 20.5, 20.5, 27.9 }, { 37.9, 37.9, 40.4 } }; +static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 39.0, 39.0, 49.0 }, { 21.0, 21.0, 30.0 }, { 37.9, 37.9, 40.4 } }; // NVLink, PCI, Network #define NCCL_HW_NVLINK 0 @@ -63,11 +63,11 @@ static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 37.9 // Tree/Simple is the latency a 256kB chunk, which is ~ base lat + 256k/12GB/s (+ 256k/12GB/s for the network). static const float hwLat [3][NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { /* NVLINK */ - { /* Tree (LL/LL128/Simple)*/ { 1.2, 1.2, 3.8 }, /* Ring (LL/LL128/Simple)*/ { 2.3, 2.3, 2.7 }, /* CollNet (LL/LL128/Simple)*/ { 1.2, 1.2, 3.8 } }, + { /* Tree (LL/LL128/Simple)*/ { 2.5, 2.5, 5.5 }, /* Ring (LL/LL128/Simple)*/ { 2.5, 2.5, 5 }, /* CollNet (LL/LL128/Simple)*/ { 1.2, 1.2, 3.8 } }, /* PCI */ { /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 1.3, 1.3, 1.9 }, /* CollNet (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 } }, /* NET */ - { /* Tree (LL/LL128/Simple)*/ { 9.8, 9.8, 19.5 }, /* Ring (LL/LL128/Simple)*/ { 2.0, 2.0, 4.5 }, /* CollNet (LL/LL128/Simple)*/ { 9.8, 9.8, 19.5 } } + { /* Tree (LL/LL128/Simple)*/ { 28.0, 28.0, 66.0 }, /* Ring (LL/LL128/Simple)*/ { 8.5, 8.5, 19.0 }, /* CollNet (LL/LL128/Simple)*/ { 9.8, 9.8, 19.5 } } }; // LL128 max BW (per channel) for the different collectives @@ -158,7 +158,6 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom comm->latencies[coll][a][p] = baseLat[a][p]; float intraLat = hwLat[intraHw[a]][a][p]; float interLat = hwLat[NCCL_HW_NET][a][p]; - if (nNodes > 1 && p == NCCL_PROTO_LL) intraLat *= 1.8; if (a == NCCL_ALGO_RING) { float lat = hwLat[hw[a]][a][p]; if ((coll == ncclFuncReduce || coll == ncclFuncBroadcast)) { From 310dc0e738e138ba0532336b9dee8fc27bc3cc3e Mon Sep 17 00:00:00 2001 From: Eiden Yoshida Date: Wed, 24 Feb 2021 19:01:56 -0700 Subject: [PATCH 3/4] Add ubuntu gfx908 target to CI [ROCm/rccl commit: 25a7dad765de512477d007b445c5d4fad4a67f7f] --- projects/rccl/.jenkins/precheckin.groovy | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rccl/.jenkins/precheckin.groovy b/projects/rccl/.jenkins/precheckin.groovy index f02ef151dc..27a723e9b2 100644 --- a/projects/rccl/.jenkins/precheckin.groovy +++ b/projects/rccl/.jenkins/precheckin.groovy @@ -55,7 +55,7 @@ ci: { propertyList = auxiliary.appendPropertyList(propertyList) - def jobNameList = ["compute-rocm-dkms-no-npi-hipclang":([sles15sp1:['4gfx906'],centos8:['4gfx908'],centos7:['4gfx906'],ubuntu18:['4gfx906']])] + def jobNameList = ["compute-rocm-dkms-no-npi-hipclang":([sles15sp1:['4gfx906'],centos8:['4gfx908'],centos7:['4gfx906'],ubuntu18:['4gfx906', '4gfx908']])] jobNameList = auxiliary.appendJobNameList(jobNameList) From 4c42a55c53a57240273b746f146d90aad0b9903f Mon Sep 17 00:00:00 2001 From: Stanley Tsang Date: Thu, 25 Feb 2021 23:59:44 +0000 Subject: [PATCH 4/4] Temporarily disabling multiprocess unit tests [ROCm/rccl commit: 5d3b1fa4d0b9259d1707b0269bbcb66a2d6d8473] --- projects/rccl/.jenkins/common.groovy | 1 - projects/rccl/install.sh | 2 -- 2 files changed, 3 deletions(-) diff --git a/projects/rccl/.jenkins/common.groovy b/projects/rccl/.jenkins/common.groovy index 037c63a0c0..dcc77c3273 100644 --- a/projects/rccl/.jenkins/common.groovy +++ b/projects/rccl/.jenkins/common.groovy @@ -22,7 +22,6 @@ def runTestCommand (platform, project) set -x cd ${project.paths.project_build_prefix}/build/release/test ${sudo} NCCL_DEBUG=INFO HSA_FORCE_FINE_GRAIN_PCIE=1 ./UnitTests --gtest_filter="BroadcastCorrectnessSweep*:*float32*" --gtest_output=xml --gtest_color=yes - ${sudo} NCCL_DEBUG=INFO HSA_FORCE_FINE_GRAIN_PCIE=1 ./UnitTestsMultiProcess --gtest_filter="BroadcastMultiProcessCorrectnessSweep*:*float32*" --gtest_output=xml --gtest_color=yes """ platform.runCommand(this, command) diff --git a/projects/rccl/install.sh b/projects/rccl/install.sh index 40946b4612..62a22f43cf 100755 --- a/projects/rccl/install.sh +++ b/projects/rccl/install.sh @@ -199,10 +199,8 @@ if ($run_tests); then if (test -f "./test/UnitTests"); then if ($run_tests_all); then ./test/UnitTests - ./test/UnitTestsMultiProcess else ./test/UnitTests --gtest_filter="BroadcastCorrectnessSweep*:*float32*" - ./test/UnitTestsMultiProcess --gtest_filter="BroadcastMultiProcessCorrectnessSweep*:*float32*" fi else echo "Unit tests have not been built yet; please re-run script with -t to build unit tests."