Merge remote-tracking branch 'origin/develop' into 2.8.3

[ROCm/rccl commit: ab1e7a0318]

projects/rccl/src/clique/CliqueManager.cc

@@ -46,8 +46,8 @@ int*               CliqueManager::m_staticGpuBarrierMem             = NULL;
 
 // Define some environment variables that affect clique-based kernels
 RCCL_PARAM(EnableClique, "ENABLE_CLIQUE", 0);                                  // Opt-in environment variable for clique-based kernels
-RCCL_PARAM(AllReduceCliqueByteLimit, "CLIQUE_ALLREDUCE_BYTE_LIMIT", 2097152);  // Max number of bytes to use clique-based kernels for all reduce
-RCCL_PARAM(AllReduceNumChannels,     "CLIQUE_ALLREDUCE_NCHANNELS", 4);         // Number of channels to use for all-reduce
+RCCL_PARAM(AllReduceCliqueByteLimit, "CLIQUE_ALLREDUCE_BYTE_LIMIT", 16777216); // Max number of bytes to use clique-based kernels for all reduce
+RCCL_PARAM(AllReduceNumChannels,     "CLIQUE_ALLREDUCE_NCHANNELS", 0);         // Number of channels to use for all-reduce. (0 for auto-select)
 
 CliqueManager::CliqueManager(int          const  rank,
                              int          const  numRanks,
@@ -315,7 +315,22 @@ ncclResult_t CliqueManager::GetNumChannelsToUse(ncclFunc_t const coll,
   *numChannelstoUse = 1;
 
   if (coll == ncclFuncAllReduce) {
-    *numChannelstoUse = std::min((int)rcclParamAllReduceNumChannels(), totalNumChannels);
+    if (rcclParamAllReduceNumChannels() == 0)
+    {
+      // NOTE: These are currently based on collected data and not necessarily ideal for all hardware
+      int numChannels;
+      if (totalBytes <= 65536) numChannels = 1;
+      else if (totalBytes <= 262144) numChannels = 2;
+      else if (totalBytes <= 524288) numChannels = 4;
+      else if (totalBytes <= 2097152) numChannels = 8;
+      else numChannels = 11;
+
+      *numChannelstoUse = std::min(numChannels, totalNumChannels);
+    }
+    else
+    {
+      *numChannelstoUse = std::min((int)rcclParamAllReduceNumChannels(), totalNumChannels);
+    }
   }
 
   return ncclSuccess;
@@ -337,9 +352,6 @@ ncclResult_t CliqueManager::SetCliqueArgs(ncclWorkElem* args)
   int opIndex = args->opCount % NCCL_MAX_OPS;
   args->clique.ptrs = &m_pinnedCliquePtrs[opIndex];
 
-  // Determine number of channels to use for this collective
-  args->clique.nChannels = rcclParamAllReduceNumChannels();
-
   return ncclSuccess;
 }
 

projects/rccl/src/collectives/device/gather.h

@@ -0,0 +1,75 @@
+/*************************************************************************
+ * Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
+ * Modifications Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#include "devcomm.h"
+#include "primitives.h"
+#include "collectives.h"
+
+template<int UNROLL, class FUNC, typename T>
+__attribute__((noinline))
+__device__ void ncclGatherKernel(struct CollectiveArgs* args) {
+  const int tid = threadIdx.x;
+  const int nthreads = args->coll.nThreads;
+  const int nChannels = args->coll.nChannels;
+  struct ncclDevComm* comm = args->comm;
+  struct ncclChannel* channel = comm->channels+blockIdx.x;
+  struct ncclRing* ring = &channel->ring;
+  const ssize_t size = args->coll.count;
+  const int nranks = comm->nRanks;
+  const int bid = args->coll.bid;
+  const int rank = ring->devUserRanks[0];
+  const int stepSize = comm->buffSizes[NCCL_PROTO_SIMPLE] / (sizeof(T)*NCCL_STEPS);
+  const int chunkSize = stepSize * GATHER_CHUNKSTEPS;
+  const int peersPerChan = DIVUP(nranks, nChannels);
+  const ssize_t loopSize = (peersPerChan == 1 ? (nChannels/nranks)*(ssize_t)chunkSize : (ssize_t)chunkSize);
+  const int root = args->coll.root;
+
+  // Compute pointers
+  const T * __restrict__ thisInput = (const T*)args->sendbuff;
+  T * __restrict__ thisOutput = (T*)args->recvbuff;
+
+  for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) {
+    for (int i = 0; i < peersPerChan; i++) {
+      if ((peersPerChan == 1 && blockIdx.x >= (nChannels/nranks)*nranks) ||
+        (peersPerChan > 1 && blockIdx.x*peersPerChan+i >= nranks))
+        continue;
+      int realChunkSize = min(chunkSize, DIVUP(size-gridOffset, (peersPerChan == 1 ? (nChannels/nranks) : 1)));
+      ALIGN_SIZE(realChunkSize, nthreads*sizeof(uint64_t)/sizeof(T));
+      ssize_t chunkOffset = gridOffset + (peersPerChan == 1 ? (bid/nranks)*realChunkSize : 0);
+      int nelem = min(realChunkSize, size-chunkOffset);
+      if ((blockIdx.x*peersPerChan+i)%nranks == 0 && rank == root) {
+        const T* sendbuff = thisInput+chunkOffset;
+        T* recvbuff = thisOutput+chunkOffset+rank*size;
+        if (tid < nthreads && sendbuff != recvbuff) {
+          // local copy
+          ReduceOrCopyMulti<UNROLL, FUNC, T, 1, 1, 1, 1>(tid, nthreads, 1, &sendbuff, 1, &recvbuff, nelem);
+        }
+      }
+      else {
+        int peerSend = (rank+(blockIdx.x*peersPerChan)+i)%nranks;
+        int peerRecv = (2*nranks+rank-((blockIdx.x*peersPerChan)%nranks)-(i%nranks))%nranks;
+        int peerNone = -1;
+        if (rank == root) {
+          ncclPrimitives<UNROLL, GATHER_CHUNKSTEPS/GATHER_SLICESTEPS, GATHER_SLICESTEPS, T, 1, 1, 0, FUNC>
+            prims(tid, nthreads, &peerRecv, &peerNone, NULL, stepSize, channel, comm);
+
+          ssize_t recv_offset = chunkOffset + peerRecv*size;
+          prims.recv(thisOutput+recv_offset, nelem);
+        }
+        else {
+          if (peerSend == root) {
+            ncclPrimitives<UNROLL, GATHER_CHUNKSTEPS/GATHER_SLICESTEPS, GATHER_SLICESTEPS, T, 1, 1, 0, FUNC>
+              prims(tid, nthreads, &peerNone, &peerSend, NULL, stepSize, channel, comm);
+
+            ssize_t send_offset = chunkOffset;
+            prims.send(thisInput+send_offset, nelem);
+          }
+        }
+      }
+    }
+  }
+}

projects/rccl/src/collectives/device/scatter.h

@@ -0,0 +1,75 @@
+/*************************************************************************
+ * Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
+ * Modifications Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#include "devcomm.h"
+#include "primitives.h"
+#include "collectives.h"
+
+template<int UNROLL, class FUNC, typename T>
+__attribute__((noinline))
+__device__ void ncclScatterKernel(struct CollectiveArgs* args) {
+  const int tid = threadIdx.x;
+  const int nthreads = args->coll.nThreads;
+  const int nChannels = args->coll.nChannels;
+  struct ncclDevComm* comm = args->comm;
+  struct ncclChannel* channel = comm->channels+blockIdx.x;
+  struct ncclRing* ring = &channel->ring;
+  const ssize_t size = args->coll.count;
+  const int nranks = comm->nRanks;
+  const int bid = args->coll.bid;
+  const int rank = ring->devUserRanks[0];
+  const int stepSize = comm->buffSizes[NCCL_PROTO_SIMPLE] / (sizeof(T)*NCCL_STEPS);
+  const int chunkSize = stepSize * SCATTER_CHUNKSTEPS;
+  const int peersPerChan = DIVUP(nranks, nChannels);
+  const ssize_t loopSize = (peersPerChan == 1 ? (nChannels/nranks)*(ssize_t)chunkSize : (ssize_t)chunkSize);
+  const int root = args->coll.root;
+
+  // Compute pointers
+  const T * __restrict__ thisInput = (const T*)args->sendbuff;
+  T * __restrict__ thisOutput = (T*)args->recvbuff;
+
+  for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) {
+    for (int i = 0; i < peersPerChan; i++) {
+      if ((peersPerChan == 1 && blockIdx.x >= (nChannels/nranks)*nranks) ||
+        (peersPerChan > 1 && blockIdx.x*peersPerChan+i >= nranks))
+        continue;
+      int realChunkSize = min(chunkSize, DIVUP(size-gridOffset, (peersPerChan == 1 ? (nChannels/nranks) : 1)));
+      ALIGN_SIZE(realChunkSize, nthreads*sizeof(uint64_t)/sizeof(T));
+      ssize_t chunkOffset = gridOffset + (peersPerChan == 1 ? (bid/nranks)*realChunkSize : 0);
+      int nelem = min(realChunkSize, size-chunkOffset);
+      if ((blockIdx.x*peersPerChan+i)%nranks == 0 && rank == root) {
+        const T* sendbuff = thisInput+chunkOffset+rank*size;
+        T* recvbuff = thisOutput+chunkOffset;
+        if (tid < nthreads && sendbuff != recvbuff) {
+          // local copy
+          ReduceOrCopyMulti<UNROLL, FUNC, T, 1, 1, 1, 1>(tid, nthreads, 1, &sendbuff, 1, &recvbuff, nelem);
+        }
+      }
+      else {
+        int peerSend = (rank+(blockIdx.x*peersPerChan)+i)%nranks;
+        int peerRecv = (2*nranks+rank-((blockIdx.x*peersPerChan)%nranks)-(i%nranks))%nranks;
+        int peerNone = -1;
+        if (rank == root) {
+          ncclPrimitives<UNROLL, SCATTER_CHUNKSTEPS/SCATTER_SLICESTEPS, SCATTER_SLICESTEPS, T, 1, 1, 0, FUNC>
+            prims(tid, nthreads, &peerNone, &peerSend, NULL, stepSize, channel, comm);
+
+          ssize_t send_offset = chunkOffset + peerSend*size;
+          prims.send(thisInput+send_offset, nelem);
+        }
+        else {
+          if (peerRecv == root) {
+            ncclPrimitives<UNROLL, SCATTER_CHUNKSTEPS/SCATTER_SLICESTEPS, SCATTER_SLICESTEPS, T, 1, 1, 0, FUNC>
+              prims(tid, nthreads, &peerRecv, &peerNone, NULL, stepSize, channel, comm);
+
+            ssize_t recv_offset = chunkOffset;
+            prims.recv(thisOutput+recv_offset, nelem);
+          }
+        }
+      }
+    }
+  }
+}

projects/rccl/src/graph/rome_models.h

@@ -213,6 +213,28 @@ static struct rcclRomeModel rome_model_44 = {
   .ringBase = "5 4 7 6 2 1 3 0|5 6 7 4 1 0 2 3",
 };
 
+static struct rcclRomeModel rome_model_45 = {
+  .nGpus = 8, .nCpus = 7, .nNics = 0, .nLinks = 3,
+  .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, },
+  .nicIds = { },
+  .gpuNuma = { 1, 2, 2, 3, 5, 5, 6, 7, },
+  .nicNuma = { },
+  .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 = "00102010002010",
+  .ringBase = "0 1 2 3 4 5 6 7|0 2 5 7 4 6 1 3|0 3 1 6 4 7 5 2|0 7 6 5 4 3 2 1",
+};
+
+static struct rcclRomeModel rome_model_46 = {
+  .nGpus = 8, .nCpus = 7, .nNics = 1, .nLinks = 3,
+  .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, },
+  .nicIds = { 0xe1000, },
+  .gpuNuma = { 1, 2, 2, 3, 5, 5, 6, 7, },
+  .nicNuma = { 4, },
+  .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 = "00102010012010",
+  .ringBase = "6 5 7 4 1 2 3 0|7 4 6 5 1 0 3 2",
+};
+
 static struct rcclRomeModel romeTopoModels[] = {
   rome_model_22,
   rome_model_25,
@@ -230,4 +252,6 @@ static struct rcclRomeModel romeTopoModels[] = {
   rome_model_40,
   rome_model_42,
   rome_model_44,
+  rome_model_45,
+  rome_model_46,
 };

projects/rccl/src/graph/search.cc

@@ -462,7 +462,7 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
         for (int i = 0; i<system->nodes[GPU].count; i++)
           if (paths[i].count < paths[f].count) f = i;
         int t = 1 << 10;
-        NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, FORCED_ORDER_PCI, &t, NET, n, f));
+        NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, (f == 0) ? FORCED_ORDER_PCI : 0, &t, NET, n, f));
         if (t == -1) *time = -1;
       }
 
@@ -998,8 +998,10 @@ static bool permuteGpuIds(int *g, int n, int last, struct rcclRomeModel* ref, st
     if (i < ref->nGpus) return false;
     // match XGMI connection
     for (i = 0; i < ref->nGpus; i++) {
-      for (j = 0; j < ref->nGpus; j++)
+      for (j = 0; j < ref->nGpus; j++) {
         if (ref->connMatrix[i*ref->nGpus+j] != topo->connMatrix[g[i]*ref->nGpus+g[j]]) break;
+        if ((ref->gpuIds[i]-ref->gpuIds[j])*(topo->gpuIds[g[i]]-topo->gpuIds[g[j]]) < 0) break;
+      }
       if (j < ref->nGpus) break;
     }
     if (i < ref->nGpus) return false;

projects/rccl/src/graph/xml.cc

@@ -492,7 +492,10 @@ ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml*
     char portSpeedStr[MAX_STR_LEN];
     float portSpeed;
     NCCLCHECK(ncclTopoGetStrFromSys(path, "../max_link_speed", portSpeedStr));
-    sscanf(portSpeedStr, "%f GT/s", &portSpeed);
+    if (portSpeedStr[0])
+      sscanf(portSpeedStr, "%f GT/s", &portSpeed);
+    else
+      portSpeed = deviceSpeed;
     NCCLCHECK(xmlSetAttr(pciNode, "link_speed", portSpeed < deviceSpeed ? portSpeedStr : deviceSpeedStr));
   }
   NCCLCHECK(xmlGetAttrIndex(pciNode, "link_width", &index));
@@ -502,7 +505,11 @@ ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml*
     NCCLCHECK(ncclTopoGetStrFromSys(path, "max_link_width", strValue));
     int deviceWidth = strtol(strValue, NULL, 0);
     NCCLCHECK(ncclTopoGetStrFromSys(path, "../max_link_width", strValue));
-    int portWidth = strtol(strValue, NULL, 0);
+    int portWidth;
+    if (strValue[0])
+      portWidth = strtol(strValue, NULL, 0);
+    else
+      portWidth = deviceWidth;
     NCCLCHECK(xmlSetAttrInt(pciNode, "link_width", std::min(deviceWidth,portWidth)));
   }
   struct ncclXmlNode* parent = pciNode->parent;