collnet: support multiple NICs (#335)

src/enqueue.cc

@@ -283,10 +283,11 @@ static ncclResult_t getAlgoInfo(struct ncclInfo* info) {
   info->algorithm = -1;
   info->protocol = -1;
   int nAlgos = NCCL_NUM_ALGORITHMS;
+  #define SHARP_COLL_SAT_THRESHOLD 16384
 
   // Check collNet support
   int collNetTypeSupport = 0;
-  if (info->comm->collNetSupport)
+  if (info->comm->collNetSupport && info->nBytes < SHARP_COLL_SAT_THRESHOLD*comm->collNetnChannels/2)
     NCCLCHECK(collNetReduceSupport(info->datatype, info->op, &collNetTypeSupport));
   if (collNetTypeSupport != 1) nAlgos--;
   for (int a=0; a<nAlgos; a++) {

src/graph/connect.cc

@@ -20,19 +20,15 @@ ncclResult_t ncclTopoPreset(struct ncclComm* comm,
     struct ncclTopoRanks* topoRanks) {
   int rank = comm->rank;
   int localRanks = comm->localRanks;
-  int nChannels = comm->nChannels;
 
-  for (int c=0; c<nChannels; c++) {
+  for (int c=0; c<comm->nChannels; c++) {
     struct ncclChannel* channel = comm->channels+c;
     channel->ring.prev = channel->ring.next = -1;
     channel->tree.up = -1;
     for (int i=0; i<NCCL_MAX_TREE_ARITY; i++) channel->tree.down[i] = -1;
-    channel->collTree.up = -1;
-    for (int i=0; i<NCCL_MAX_TREE_ARITY; i++) channel->collTree.down[i] = -1;
 
     int* ringIntra = ringGraph->intra+c*localRanks;
     int* treeIntra = treeGraph->intra+c*localRanks;
-    int* collNetIntra = collNetGraph->intra+c*localRanks;
 
     for (int i=0; i<localRanks; i++) {
       if (ringIntra[i] == rank) {
@@ -52,6 +48,23 @@ ncclResult_t ncclTopoPreset(struct ncclComm* comm,
         channel->tree.up         = i == 0 ? -1 : treeIntra[i-1];
         channel->tree.down[0]    = i == localRanks-1 ? -1 : treeIntra[i+1];
       }
+    }
+    topoRanks->ringPrev[c] = channel->ring.prev;
+    topoRanks->ringNext[c] = channel->ring.next;
+  }
+  // Duplicate channels rings/trees
+  struct ncclChannel* channel0 = comm->channels;
+  struct ncclChannel* channel1 = channel0+comm->nChannels;
+  memcpy(channel1, channel0, comm->nChannels*sizeof(struct ncclChannel));
+  // Setup collnet tree
+  for (int c=0; c<comm->collNetnChannels; c++) {
+    struct ncclChannel* channel = comm->channels+c;
+    channel->collTree.up = -1;
+    for (int i=0; i<NCCL_MAX_TREE_ARITY; i++) channel->collTree.down[i] = -1;
+
+    int* collNetIntra = collNetGraph->intra+c*localRanks;
+
+    for (int i=0; i<localRanks; i++) {
       if (collNetIntra[i] == rank) {
         int prev = (i-1+localRanks)%localRanks, next = (i+1)%localRanks;
 
@@ -59,13 +72,7 @@ ncclResult_t ncclTopoPreset(struct ncclComm* comm,
         channel->collTree.down[0] = collNetIntra[next];
       }
     }
-    topoRanks->ringPrev[c] = channel->ring.prev;
-    topoRanks->ringNext[c] = channel->ring.next;
   }
-  // Duplicate channels rings/trees
-  struct ncclChannel* channel0 = comm->channels;
-  struct ncclChannel* channel1 = channel0+nChannels;
-  memcpy(channel1, channel0, nChannels*sizeof(struct ncclChannel));
   return ncclSuccess;
 }
 

src/graph/search.cc

@@ -1153,9 +1153,32 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph
     NCCLCHECK(parseRome4P2H(system, graph));
   }
   if (graph->collNet && graph->nChannels) {
-    graph->nChannels = 1;
-    memcpy(graph->intra+graph->nChannels*ngpus, graph->intra, ngpus*sizeof(int));
-    memcpy(graph->inter+graph->nChannels*2, graph->inter, 2*sizeof(int));
+    struct ncclTopoGraph tmpGraph;
+    memcpy(&tmpGraph, graph, sizeof(struct ncclTopoGraph));
+    int nets[MAXCHANNELS], n = 0;
+    for (int i = 0; i < tmpGraph.nChannels; i++) {
+      int j;
+      for (j = 0; j < n; j++) {
+        if (nets[j] == tmpGraph.inter[i*2])
+          break;
+      }
+      if (j >= n)
+        nets[n++] = tmpGraph.inter[i*2];
+    }
+    for (int i = 0; i < n; i++) {
+      int j;
+      for (j = 0; j < tmpGraph.nChannels; j++) {
+        if (nets[i] == tmpGraph.inter[j*2])
+          break;
+      }
+      if (j < tmpGraph.nChannels) {
+        memcpy(graph->intra+i*ngpus, &tmpGraph.intra[j*ngpus], ngpus*sizeof(int));
+        memcpy(graph->inter+i*2, &tmpGraph.inter[j*2], 2*sizeof(int));
+      }
+    }
+    memcpy(graph->intra+n*ngpus, graph->intra, ngpus*sizeof(int)*n);
+    memcpy(graph->inter+n*2, graph->inter, 2*sizeof(int)*n);
+    graph->nChannels = n;
   }
   if (graph->nChannels) return ncclSuccess;
 
@@ -1286,7 +1309,12 @@ done:
     graph->nChannels = 1;
   }
 
-  if (graph->speedIntra >= 25.0) {
+  if (graph->nChannels && graph->collNet) {
+    // duplicate collnet channels
+    memcpy(graph->intra+graph->nChannels*ngpus, graph->intra, ngpus*sizeof(int)*graph->nChannels);
+    memcpy(graph->inter+graph->nChannels*2, graph->inter, 2*sizeof(int)*graph->nChannels);
+  }
+  else if (graph->speedIntra >= 25.0) {
     int dupChannels = std::min(graph->nChannels*2, graph->maxChannels);
     memcpy(graph->intra+graph->nChannels*ngpus, graph->intra, (dupChannels-graph->nChannels)*ngpus*sizeof(int));
     memcpy(graph->inter+graph->nChannels*2,graph->inter, (dupChannels-graph->nChannels)*2*sizeof(int));

src/graph/tuning.cc

@@ -63,11 +63,11 @@ static const float baseLat  [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 39.0
 // Tree/Simple is the latency a 256kB chunk, which is ~ base lat + 256k/12GB/s (+ 256k/12GB/s for the network).
 static const float hwLat [3][NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] =
 { /* NVLINK */
-  { /* Tree (LL/LL128/Simple)*/ { 2.5, 2.5, 5.5 }, /* Ring (LL/LL128/Simple)*/ { 2.5, 2.5, 5 }, /* CollNet (LL/LL128/Simple)*/ { 1.2, 1.2, 3.8 } },
+  { /* Tree (LL/LL128/Simple)*/ { 2.5, 2.5, 5.5 }, /* Ring (LL/LL128/Simple)*/ { 2.5, 2.5, 5 }, /* CollNet (LL/LL128/Simple)*/ { 1.1, 1.1, 3.3 } },
   /* PCI */
-  { /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 1.3, 1.3, 1.9 }, /* CollNet (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 } },
+  { /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 1.3, 1.3, 1.9 }, /* CollNet (LL/LL128/Simple)*/ { 1.1, 1.1, 1.7 } },
   /* NET */
-  { /* Tree (LL/LL128/Simple)*/ { 28.0, 28.0, 66.0 }, /* Ring (LL/LL128/Simple)*/ { 8.5, 8.5, 19.0 }, /* CollNet (LL/LL128/Simple)*/ { 9.8, 9.8, 19.5 } }
+  { /* Tree (LL/LL128/Simple)*/ { 28.0, 28.0, 66.0 }, /* Ring (LL/LL128/Simple)*/ { 8.5, 8.5, 19.0 }, /* CollNet (LL/LL128/Simple)*/ { 6.5, 6.5, 14.5 } }
 };
 
 // LL128 max BW (per channel) for the different collectives
@@ -148,7 +148,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
         if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_LL128) busBw = std::min(busBw * (nNodes == 1 ? 7.0/9.0 : 0.915 /*120.0/128.0*/), ll128MaxBwPerCh[coll]*graphs[a]->nChannels);
 #endif
         if (a == NCCL_ALGO_COLLNET) busBw *= .9;
-        if (a == NCCL_ALGO_COLLNET && p == NCCL_PROTO_LL) busBw *= 1.0/6.0; // Take into account that GDR read is disabled on both sides
+        if (a == NCCL_ALGO_COLLNET && p == NCCL_PROTO_LL) busBw *= 1.0/2.0; // Take into account that GDR read is disabled on both sides
         if (a == NCCL_ALGO_COLLNET && p == NCCL_PROTO_LL128) busBw = 0;  // CollNet does not support LL128
 
         // Convert bus BW to algorithm BW

src/init.cc

@@ -958,6 +958,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
   allGather3Data[rank].collNet.typeIntra = collNetGraph.typeIntra;
   allGather3Data[rank].collNet.typeInter = collNetGraph.typeInter;
 
+  // CollNet channels are already duplicated
+  comm->collNetnChannels = 2*collNetGraph.nChannels;
   NCCLCHECK(ncclTopoPreset(comm, &treeGraph, &ringGraph, &collNetGraph, &allGather3Data[rank].topoRanks));
 
   NCCLCHECK(bootstrapAllGather(comm->bootstrap, allGather3Data, sizeof(*allGather3Data)));
@@ -1035,9 +1037,9 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
   if (comm->nNodes > 1 &&
       ncclParamCollNetEnable() == 1 &&
       collNetSupport() && collNetGraph.nChannels) {
-    // Force 2 channels for CollNet
-    comm->collNetnChannels = collNetGraph.nChannels = 2;
     NCCLCHECK(ncclTopoConnectCollNet(comm, &collNetGraph, rank));
+  } else {
+    comm->collNetnChannels = 0;
   }
 
   free(allTopoRanks);
@@ -1094,6 +1096,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
   if (comm->nNodes > 1 &&
       ncclParamCollNetEnable() == 1 &&
       collNetSupport() && collNetGraph.nChannels) {
+    for (int c=comm->nChannels; c<comm->collNetnChannels; c++)
+      NCCLCHECK(initChannel(comm, c));;
     int logicChannels = comm->collNetnChannels/2;
     int collNetSetupFail = 0;
     const int recvIndex = 0;  // recv GPU index is always 0