Trim NICs when all GPUs are connected by XGMI (#430)

* Trim NICs when all GPUs are connected by XGMI

* Only enable clique with maximum of 2 hops

[ROCm/rccl commit: 29c729d8b6]

projects/rccl/src/graph/search.cc

@@ -1046,7 +1046,8 @@ ncclResult_t ncclTopoGetIntraNetDev(struct ncclTopoSystem* system, int rank, str
   return ncclSuccess;
 }
 
-ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, int cudaDev2, bool* isXGMI, bool direct_only) {
+ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, int cudaDev2, bool* isXGMI, int maxInter, int nInter, int *inter) {
+  int interGpus[MAX_XGMI_INTER_GPUS+1];
   int ngpus = system->nodes[GPU].count;
   *isXGMI = false;
   // check for direct XGMI connection
@@ -1065,15 +1066,39 @@ ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, in
       }
     }
   }
-  if (direct_only) return ncclSuccess;
-  // check if there is intermediate GPU that is connected to both
+  if (maxInter == 0) return ncclSuccess;
+  // check if there are intermediate GPUs that are connected to both
+  bool res1, res2, res3;
+  int j;
+  for (j=0; j<nInter; j++) {
+    bool res1;
+    ncclTopoGetLinkType(system, inter[j], inter[j+1], &res1, 0);
+    if (!res1) break;
+  }
+  if (j<nInter) return ncclSuccess;
+  if (nInter > 0 && inter != nullptr) {
+    ncclTopoGetLinkType(system, inter[nInter], cudaDev2, &res2, 0);
+    if (res2) {
+      *isXGMI = true;
+      return ncclSuccess;
+    }
+    memcpy(interGpus+1, inter+1, sizeof(int)*nInter);
+  }
+  interGpus[0] = cudaDev1;
+  // add one more intermediate GPU recursively util reaching max depth
+  nInter++;
+  if (nInter+2 > ngpus || nInter > MAX_XGMI_INTER_GPUS || nInter > maxInter) return ncclSuccess;
   for (int i=0; i<ngpus; i++) {
-    if (system->nodes[GPU].nodes[i].gpu.dev == cudaDev1 || system->nodes[GPU].nodes[i].gpu.dev == cudaDev2)
-      continue;
-    bool res1, res2;
-    ncclTopoGetLinkType(system, system->nodes[GPU].nodes[i].gpu.dev, cudaDev1, &res1, true);
-    ncclTopoGetLinkType(system, system->nodes[GPU].nodes[i].gpu.dev, cudaDev2, &res2, true);
-    if (res1 && res2) {
+    int dev = system->nodes[GPU].nodes[i].gpu.dev;
+    // skip duplicated GPU
+    if (dev == cudaDev2) continue;
+    for (j=0; j<nInter; j++)
+      if (dev == interGpus[j]) break;
+    if (j<nInter) continue;
+    // check connectivity with intermediate GPUs
+    interGpus[nInter] = dev;
+    ncclTopoGetLinkType(system, cudaDev1, cudaDev2, &res3, maxInter, nInter, interGpus);
+    if (res3) {
       *isXGMI = true;
       return ncclSuccess;
     }