From b5b6c6acdd40b816e79fcffb251346ca73dd7bcd Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Thu, 9 Apr 2020 17:14:07 -0700 Subject: [PATCH 1/4] Fix bug #307 : wrong NIC selection on the reduction tree. The reduction tree (tree up) was inverting the NICs to use, causing performance issue in cases where we are using different NICs on a given channel. --- src/graph/search.cc | 7 +++++-- src/include/graph.h | 2 +- src/transport/coll_net.cc | 4 ++-- src/transport/net.cc | 4 ++-- 4 files changed, 10 insertions(+), 7 deletions(-) diff --git a/src/graph/search.cc b/src/graph/search.cc index b4c3e3510d..1bbb7d3371 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -804,7 +804,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; } diff --git a/src/include/graph.h b/src/include/graph.h index b27ea3523c..18144404b9 100644 --- a/src/include/graph.h +++ b/src/include/graph.h @@ -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); diff --git a/src/transport/coll_net.cc b/src/transport/coll_net.cc index 73e9fdd084..435c88dd1d 100644 --- a/src/transport/coll_net.cc +++ b/src/transport/coll_net.cc @@ -84,7 +84,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); @@ -110,7 +110,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); diff --git a/src/transport/net.cc b/src/transport/net.cc index db82a403f4..288ad92d4a 100644 --- a/src/transport/net.cc +++ b/src/transport/net.cc @@ -60,7 +60,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); @@ -83,7 +83,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); From a783484ab5003e96774c12d14f555ef969a410fa Mon Sep 17 00:00:00 2001 From: aokomoriuta Date: Wed, 29 Jan 2020 22:05:12 +0900 Subject: [PATCH 2/4] Fix wrong variable name "slice" to "chunk" https://github.com/NVIDIA/nccl/issues/287 --- src/collectives/device/all_reduce.h | 66 ++++++++++++++--------------- 1 file changed, 33 insertions(+), 33 deletions(-) diff --git a/src/collectives/device/all_reduce.h b/src/collectives/device/all_reduce.h index 173b5fa695..4e04f8829a 100644 --- a/src/collectives/device/all_reduce.h +++ b/src/collectives/device/all_reduce.h @@ -37,19 +37,19 @@ __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); prims.send(thisInput+offset, nelem); // k-2 steps: reduce and copy to next GPU for (int j=2; jdevUserRanks[nranks-j]; - offset = chunkOffset + slice * realChunkSize; + chunk = ring->devUserRanks[nranks-j]; + offset = chunkOffset + chunk * realChunkSize; nelem = min(realChunkSize, size-offset); prims.recvReduceSend(thisInput+offset, nelem); @@ -57,24 +57,24 @@ __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); prims.directRecvReduceCopySend(thisInput+offset, thisOutput+offset, offset, nelem); // k-2 steps: copy to next GPU for (int j=1; jdevUserRanks[nranks-j]; - offset = chunkOffset + slice * realChunkSize; + chunk = ring->devUserRanks[nranks-j]; + offset = chunkOffset + chunk * realChunkSize; nelem = min(realChunkSize, size-offset); prims.directRecvCopySend(thisOutput+offset, offset, nelem); } // 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. @@ -225,19 +225,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; jdevUserRanks[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); @@ -245,24 +245,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; jdevUserRanks[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. @@ -413,19 +413,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; jdevUserRanks[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); @@ -433,24 +433,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; jdevUserRanks[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. From 23a9fbb788ff1150496238026bdf21eef741bbfd Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Thu, 16 Apr 2020 14:27:50 -0700 Subject: [PATCH 3/4] Improve robustness of PCI detection Fallback to default values when class/speed is unknown. --- src/graph/topo.cc | 6 ++---- src/graph/xml.h | 5 +++-- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/src/graph/topo.cc b/src/graph/topo.cc index 5cd8d4e9b6..3767ba92b3 100644 --- a/src/graph/topo.cc +++ b/src/graph/topo.cc @@ -303,8 +303,8 @@ ncclResult_t ncclTopoAddGpu(struct ncclXmlNode* xmlGpu, struct ncclTopoSystem* s return ncclSuccess; } -struct kvDict kvDictPciClass[] = { { "0x060400", PCI }, { "0x068000", NVS }, { "0x068001", CPU }, { "0x03", GPU }, { "0x02", 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 }, { "0x03", GPU }, { "0x02", 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; @@ -356,8 +356,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)); diff --git a/src/graph/xml.h b/src/graph/xml.h index fa04527803..22e016f3bc 100644 --- a/src/graph/xml.h +++ b/src/graph/xml.h @@ -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; From f36540f55a15683a121b6c330657af442b85c796 Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Fri, 17 Apr 2020 10:03:14 -0700 Subject: [PATCH 4/4] Fix crash when only a subset of GPUs are visible within a container. Fixes #326. --- src/graph/topo.cc | 1 + src/graph/xml.cc | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/src/graph/topo.cc b/src/graph/topo.cc index 3767ba92b3..ac6b111581 100644 --- a/src/graph/topo.cc +++ b/src/graph/topo.cc @@ -520,6 +520,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)); } diff --git a/src/graph/xml.cc b/src/graph/xml.cc index 550cfcd0cc..f138d0bbbd 100644 --- a/src/graph/xml.cc +++ b/src/graph/xml.cc @@ -569,7 +569,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) {