diff --git a/src/collectives/device/all_reduce.h b/src/collectives/device/all_reduce.h index c2abcfff37..095367fd42 100644 --- a/src/collectives/device/all_reduce.h +++ b/src/collectives/device/all_reduce.h @@ -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; jdevUserRanks[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; jdevUserRanks[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; 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); @@ -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; 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. @@ -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; 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); @@ -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; 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. diff --git a/src/graph/search.cc b/src/graph/search.cc index ba4ac472ab..fd30ee88f1 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -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; } diff --git a/src/graph/topo.cc b/src/graph/topo.cc index d94129cbdf..2ee21d7215 100644 --- a/src/graph/topo.cc +++ b/src/graph/topo.cc @@ -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)); } diff --git a/src/graph/xml.cc b/src/graph/xml.cc index 900a1ddfc1..bb9d2d7f15 100644 --- a/src/graph/xml.cc +++ b/src/graph/xml.cc @@ -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) { 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; 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 dc49c6c849..918d46d573 100644 --- a/src/transport/coll_net.cc +++ b/src/transport/coll_net.cc @@ -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); diff --git a/src/transport/net.cc b/src/transport/net.cc index f9c586b894..bc2b0655c4 100644 --- a/src/transport/net.cc +++ b/src/transport/net.cc @@ -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);