AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

Merge remote-tracking branch 'nccl/master' into HEAD

Explorar el Código
Este commit está contenido en:
Wenkai Du
2020-04-27 17:16:54 +00:00
padre cf5070f6c0 f36540f55a
commit c4edc257b0
Se han modificado 8 ficheros con 50 adiciones y 47 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
src/collectives/device/all_reduce.h
+33 -33
Ver fichero
@@ -44,11 +44,11 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
/////////////// begin AllReduce steps ///////////////
ssize_t offset;
int nelem;
int slice;
int chunk;
// step 0: push data to next GPU
slice = ring->devUserRanks[nranks-1];
offset = chunkOffset + slice * realChunkSize;
chunk = ring->devUserRanks[nranks-1];
offset = chunkOffset + chunk * realChunkSize;
nelem = min(realChunkSize, size-offset);
INIT_COUNTER;
@@ -57,8 +57,8 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
// k-2 steps: reduce and copy to next GPU
for (int j=2; j<nranks; ++j) {
slice = ring->devUserRanks[nranks-j];
offset = chunkOffset + slice * realChunkSize;
chunk = ring->devUserRanks[nranks-j];
offset = chunkOffset + chunk * realChunkSize;
nelem = min(realChunkSize, size-offset);
INIT_COUNTER;
@@ -68,8 +68,8 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
// step k-1: reduce this buffer and data, which will produce the final
// result that we store in this data and push to the next GPU
slice = ring->devUserRanks[0];
offset = chunkOffset + slice * realChunkSize;
chunk = ring->devUserRanks[0];
offset = chunkOffset + chunk * realChunkSize;
nelem = min(realChunkSize, size-offset);
INIT_COUNTER;
@@ -78,8 +78,8 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
// k-2 steps: copy to next GPU
for (int j=1; j<nranks-1; ++j) {
slice = ring->devUserRanks[nranks-j];
offset = chunkOffset + slice * realChunkSize;
chunk = ring->devUserRanks[nranks-j];
offset = chunkOffset + chunk * realChunkSize;
nelem = min(realChunkSize, size-offset);
INIT_COUNTER;
@@ -88,8 +88,8 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
}
// Make final copy from buffer to dest.
slice = ring->devUserRanks[1];
offset = chunkOffset + slice * realChunkSize;
chunk = ring->devUserRanks[1];
offset = chunkOffset + chunk * realChunkSize;
nelem = min(realChunkSize, size-offset);
// Final wait/copy.
@@ -250,19 +250,19 @@ __device__ void ncclAllReduceRingLLKernel(struct CollectiveArgs* args) {
/////////////// begin AllReduce steps ///////////////
ssize_t offset;
int nelem;
int slice;
int chunk;
// step 0: push data to next GPU
slice = ring->devUserRanks[nranks-1];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[nranks-1];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
LLprims.send(thisInput+offset, nelem);
// k-2 steps: reduce and copy to next GPU
for (int j=2; j<nranks; ++j) {
slice = ring->devUserRanks[nranks-j];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[nranks-j];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
LLprims.recvReduceSend(thisInput+offset, nelem);
@@ -270,24 +270,24 @@ __device__ void ncclAllReduceRingLLKernel(struct CollectiveArgs* args) {
// step k-1: reduce this buffer and data, which will produce the final
// result that we store in this data and push to the next GPU
slice = ring->devUserRanks[0];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[0];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
LLprims.recvReduceCopySend(thisInput+offset, thisOutput+offset, nelem);
// k-2 steps: copy to next GPU
for (int j=1; j<nranks-1; ++j) {
slice = ring->devUserRanks[nranks-j];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[nranks-j];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
LLprims.recvCopySend(thisOutput+offset, nelem);
}
// Make final copy from buffer to dest.
slice = ring->devUserRanks[1];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[1];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
// Here we need to copy from buffer to this output.
@@ -443,19 +443,19 @@ __device__ void ncclAllReduceRingLL128Kernel(struct CollectiveArgs* args) {
/////////////// begin AllReduce steps ///////////////
ssize_t offset;
int nelem;
int slice;
int chunk;
// step 0: push data to next GPU
slice = ring->devUserRanks[nranks-1];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[nranks-1];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
LLprims.send(thisInput+offset, nelem);
// k-2 steps: reduce and copy to next GPU
for (int j=2; j<nranks; ++j) {
slice = ring->devUserRanks[nranks-j];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[nranks-j];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
LLprims.recvReduceSend(thisInput+offset, nelem);
@@ -463,24 +463,24 @@ __device__ void ncclAllReduceRingLL128Kernel(struct CollectiveArgs* args) {
// step k-1: reduce this buffer and data, which will produce the final
// result that we store in this data and push to the next GPU
slice = ring->devUserRanks[0];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[0];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
LLprims.recvReduceCopySend(thisInput+offset, thisOutput+offset, nelem);
// k-2 steps: copy to next GPU
for (int j=1; j<nranks-1; ++j) {
slice = ring->devUserRanks[nranks-j];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[nranks-j];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
LLprims.recvCopySend(thisOutput+offset, nelem);
}
// Make final copy from buffer to dest.
slice = ring->devUserRanks[1];
offset = gridOffset + (slice*args->nChannels+bid) * chunkSize;
chunk = ring->devUserRanks[1];
offset = gridOffset + (chunk*args->nChannels+bid) * chunkSize;
nelem = min(chunkSize, size-offset);
// Here we need to copy from buffer to this output.
src/graph/search.cc
+5 -2
Ver fichero
@@ -944,7 +944,10 @@ ncclResult_t ncclTopoDumpGraphs(struct ncclTopoSystem* system, int ngraphs, stru
return ncclSuccess;
}
ncclResult_t ncclTopoGetNetDev(struct ncclTopoGraph* graph, int dir, int channelId, int* dev) {
*dev = graph->inter[(channelId%graph->nChannels)*2+dir];
ncclResult_t ncclTopoGetNetDev(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int rank, int channelId, int* dev) {
int channel = channelId%graph->nChannels;
int ngpus = system->nodes[GPU].count;
int index = graph->intra[channel*ngpus] == rank ? 0 : 1;
*dev = graph->inter[channel*2+index];
return ncclSuccess;
}
src/graph/topo.cc
+3 -4
Ver fichero
@@ -313,8 +313,8 @@ ncclResult_t ncclTopoAddGpu(struct ncclXmlNode* xmlGpu, struct ncclTopoSystem* s
return ncclSuccess;
}
struct kvDict kvDictPciClass[] = { { "0x060400", PCI }, { "0x068000", NVS }, { "0x068001", CPU }, { "0x030200", GPU }, { "0x030000", GPU }, { "0x038000", GPU }, { "0x020700", NIC }, { "0x020000", NIC }, { NULL, 0 } };
struct kvDict kvDictPciGen[] = { { "2.5 GT/s", 15 }, { "5 GT/s", 30 }, { "8 GT/s", 60 }, { "16 GT/s", 120 }, { NULL, 0 } }; // x100 Mbps per lane
struct kvDict kvDictPciClass[] = { { "0x060400", PCI }, { "0x068000", NVS }, { "0x068001", CPU }, { "0x030200", GPU }, { "0x030000", GPU }, { "0x038000", GPU }, { "0x020700", NIC }, { "0x020000", NIC }, { NULL, PCI /* Default fallback value */ } };
struct kvDict kvDictPciGen[] = { { "2.5 GT/s", 15 }, { "5 GT/s", 30 }, { "8 GT/s", 60 }, { "16 GT/s", 120 }, { NULL, 60 /* Default fallback */ } }; // x100 Mbps per lane
ncclResult_t ncclTopoAddPci(struct ncclXmlNode* xmlPci, struct ncclTopoSystem* system, struct ncclTopoNode* parent) {
const char* str;
@@ -366,8 +366,6 @@ ncclResult_t ncclTopoAddPci(struct ncclXmlNode* xmlPci, struct ncclTopoSystem* s
// Manage cases where speed was not indicated in /sys
if (width == 0) width = 16;
if (strlen(str) == 0 || strcasecmp(str, "Unknown speed") == 0) str = "8 GT/s";
NCCLCHECK(kvConvertToInt(str, &speed, kvDictPciGen)); // Values in 100Mbps, per lane (we want GB/s in the end)
NCCLCHECK(ncclTopoConnectNodes(node, parent, LINK_PCI, width*speed/80.0));
@@ -589,6 +587,7 @@ ncclResult_t ncclTopoGetSystem(struct ncclComm* comm, struct ncclTopoSystem** sy
NCCLCHECK(int64ToBusId(comm->peerInfo[r].busId, busId));
struct ncclXmlNode* node;
NCCLCHECK(ncclTopoFillGpu(xml, busId, &node));
if (node == NULL) continue;
NCCLCHECK(xmlSetAttrInt(node, "rank", r));
NCCLCHECK(xmlInitAttrInt(node, "gdr", comm->peerInfo[r].gdrSupport));
}
src/graph/xml.cc
+1 -1
Ver fichero
@@ -578,7 +578,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, nvmlDevice_t nvm
NCCLCHECK(xmlSetAttrInt(gpuNode, "dev", dev));
}
NCCLCHECK(xmlGetAttrInt(gpuNode, "dev", &dev));
if (dev == -1) return ncclSuccess;
if (dev == -1) { *gpuNodeRet = NULL; return ncclSuccess; }
NCCLCHECK(xmlGetAttrIndex(gpuNode, "sm", &index));
if (index == -1) {
src/graph/xml.h
+3 -2
Ver fichero
@@ -218,8 +218,9 @@ static ncclResult_t kvConvertToInt(const char* str, int* value, struct kvDict* d
}
d++;
}
WARN("KV Convert to int : could not find value of '%s' in dictionary", str);
return ncclInternalError;
INFO(NCCL_GRAPH, "KV Convert to int : could not find value of '%s' in dictionary, falling back to %d", str, d->value);
*value = d->value;
return ncclSuccess;
}
static ncclResult_t kvConvertToStr(int value, const char** str, struct kvDict* dict) {
struct kvDict* d = dict;
src/include/graph.h
+1 -1
Ver fichero
@@ -27,7 +27,7 @@ void ncclTopoFree(struct ncclTopoSystem* system);
ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm* comm);
// Query topology
ncclResult_t ncclTopoGetNetDev(struct ncclTopoGraph* graph, int dir, int channelId, int* net);
ncclResult_t ncclTopoGetNetDev(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int rank, int channelId, int* net);
ncclResult_t ncclTopoCheckP2p(struct ncclTopoSystem* system, int64_t id1, int64_t id2, int* p2p);
ncclResult_t ncclTopoCheckGdr(struct ncclTopoSystem* topo, int64_t busId, int netDev, int read, int* useGdr);
src/transport/coll_net.cc
+2 -2
Ver fichero
@@ -85,7 +85,7 @@ ncclResult_t collNetSendSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph*
NCCLCHECK(ncclCalloc(&sendResources, 1));
send->transportResources = sendResources;
NCCLCHECK(ncclTopoGetNetDev(graph, 1, channelId, &sendResources->netDev));
NCCLCHECK(ncclTopoGetNetDev(topo, graph, myInfo->rank, channelId, &sendResources->netDev));
NCCLCHECK(ncclTopoCheckGdr(topo, myInfo->busId, sendResources->netDev, 1, &sendResources->useGdr));
int sendSize = sizeof(struct ncclSendMem);
@@ -111,7 +111,7 @@ ncclResult_t collNetRecvSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph*
NCCLCHECK(ncclCalloc(&recvResources, 1));
recv->transportResources = recvResources;
NCCLCHECK(ncclTopoGetNetDev(graph, 0, channelId, &recvResources->netDev));
NCCLCHECK(ncclTopoGetNetDev(topo, graph, myInfo->rank, channelId, &recvResources->netDev));
NCCLCHECK(ncclTopoCheckGdr(topo, myInfo->busId, recvResources->netDev, 0, &recvResources->useGdr));
int sendSize = sizeof(struct ncclSendMem);
src/transport/net.cc
+2 -2
Ver fichero
@@ -62,7 +62,7 @@ ncclResult_t netSendSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph* gra
NCCLCHECK(ncclCalloc(&resources, 1));
send->transportResources = resources;
NCCLCHECK(ncclTopoGetNetDev(graph, 1, channelId, &resources->netDev));
NCCLCHECK(ncclTopoGetNetDev(topo, graph, myInfo->rank, channelId, &resources->netDev));
NCCLCHECK(ncclTopoCheckGdr(topo, myInfo->busId, resources->netDev, 1, &resources->useGdr));
int sendSize = sizeof(struct ncclSendMem);
@@ -85,7 +85,7 @@ ncclResult_t netRecvSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph* gra
NCCLCHECK(ncclCalloc(&resources, 1));
recv->transportResources = resources;
NCCLCHECK(ncclTopoGetNetDev(graph, 0, channelId, &resources->netDev));
NCCLCHECK(ncclTopoGetNetDev(topo, graph, myInfo->rank, channelId, &resources->netDev));
NCCLCHECK(ncclTopoCheckGdr(topo, myInfo->busId, resources->netDev, 0, &resources->useGdr));
int sendSize = sizeof(struct ncclSendMem);