AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Merge pull request #320 from wenkaidu/nbio

Sfoglia il codice sorgente 
Match NBIO only when GPUs and NICs are directly connected to CPU

[ROCm/rccl commit: 7191d0959b]
This commit is contained in:
Wenkai Du
2021-02-25 17:49:09 -08:00
committato da GitHub
parent fb7478e3a2 4c42a55c53
commit 897cb8d618
8 ha cambiato i file con 45 aggiunte e 28 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
projects/rccl/.jenkins/common.groovy
-1
Vedi File
@@ -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)
projects/rccl/.jenkins/precheckin.groovy
+1 -1
Vedi File
@@ -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)
projects/rccl/install.sh
-2
Vedi File
@@ -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."
projects/rccl/src/graph/connect.cc
+1 -1
Vedi File
@@ -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
projects/rccl/src/graph/paths.cc
+13 -2
Vedi File
@@ -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 <system->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 <system->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;
projects/rccl/src/graph/rome_models.h
+1 -1
Vedi File
@@ -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[] = {
projects/rccl/src/graph/search.cc
+26 -16
Vedi File
@@ -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]);
projects/rccl/src/graph/tuning.cc
+3 -4
Vedi File
@@ -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)) {