Sylvain Jeaugey
920dbe5b35
Optimization for Tree allreduce on A100. Improve aggregation performance. Use shared buffers for inter-node send/recv. Add NVTX profiling hooks. Accelerate alltoall connections by merging communication for all channels. Add support for one hop communication through NVLink, for faster send/recv communication on cubemesh topologies like DGX-1. Improve alltoall scheduling to better balance intra/inter node communication. Increase send/recv parallelism by 8x, each warp sending or receiving to a different peer. Net: move to v4. Net: make flush operation asynchronous to accelerate alltoall. Net: define maximum number of requests. Fix hang when using LL128 protocol after 2^31 steps. Fix #379 : topology injection failing when using less GPUs than described in the XML. Fix #394 : protocol mismatch causing hangs or crashes when using one GPU per node.
58 lignes
1.8 KiB
C++
58 lignes
1.8 KiB
C++
/*************************************************************************
|
|
* Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved.
|
|
*
|
|
* See LICENSE.txt for license information
|
|
************************************************************************/
|
|
|
|
#include "core.h"
|
|
|
|
#define MAXWIDTH 20
|
|
#define PREFIXLEN 15
|
|
#define STRLENGTH (PREFIXLEN+5*MAXWIDTH)
|
|
void dumpLine(int* values, int nranks, const char* prefix) {
|
|
int prefixlen = strlen(prefix);
|
|
char line[STRLENGTH+1];
|
|
line[STRLENGTH] = '\0';
|
|
memset(line, ' ', STRLENGTH);
|
|
strncpy(line, prefix, PREFIXLEN);
|
|
for (int i=0; i<nranks && i<MAXWIDTH; i++) sprintf(line+prefixlen+4*i, " %3d", values[i]);
|
|
INFO(NCCL_INIT,"%s", line);
|
|
}
|
|
|
|
ncclResult_t ncclBuildRings(int nrings, int* rings, int rank, int nranks, int* prev, int* next) {
|
|
for (int r=0; r<nrings; r++) {
|
|
char prefix[40];
|
|
/*sprintf(prefix, "[%d] Channel %d Prev : ", rank, r);
|
|
dumpLine(prev+r*nranks, nranks, prefix);
|
|
sprintf(prefix, "[%d] Channel %d Next : ", rank, r);
|
|
dumpLine(next+r*nranks, nranks, prefix);*/
|
|
|
|
int current = rank;
|
|
for (int i=0; i<nranks; i++) {
|
|
rings[r*nranks+i] = current;
|
|
current = next[r*nranks+current];
|
|
}
|
|
sprintf(prefix, "Channel %02d/%02d : ", r, nrings);
|
|
if (rank == 0) dumpLine(rings+r*nranks, nranks, prefix);
|
|
if (current != rank) {
|
|
WARN("Error : ring %d does not loop back to start (%d != %d)", r, current, rank);
|
|
return ncclInternalError;
|
|
}
|
|
// Check that all ranks are there
|
|
for (int i=0; i<nranks; i++) {
|
|
int found = 0;
|
|
for (int j=0; j<nranks; j++) {
|
|
if (rings[r*nranks+j] == i) {
|
|
found = 1;
|
|
break;
|
|
}
|
|
}
|
|
if (found == 0) {
|
|
WARN("Error : ring %d does not contain rank %d", r, i);
|
|
return ncclInternalError;
|
|
}
|
|
}
|
|
}
|
|
return ncclSuccess;
|
|
}
|