Revert "Use nontemporal in slow path and add XGMI sys type (#575)" (#590)

This reverts commit e04bba619a.

[ROCm/rccl commit: 9089e68a99]

projects/rccl/src/graph/search.cc

@@ -1171,6 +1171,7 @@ ncclResult_t ncclTopoGetIntraNetDev(struct ncclTopoSystem* system, int rank, str
 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
   for (int i=0; i<ngpus; i++) {
     if (system->nodes[GPU].nodes[i].gpu.dev == cudaDev1) {
@@ -1187,44 +1188,42 @@ ncclResult_t ncclTopoGetLinkType(struct ncclTopoSystem* system, int cudaDev1, in
       }
     }
   }
-  // try intermediate GPUs
-  if (maxInter) {
-    // check if there are intermediate GPUs that are connected to both
-    bool res1, res2, res3;
-    int j;
-    for (j=0; j<nInter; j++) {
-      ncclTopoGetLinkType(system, inter[j], inter[j+1], &res1, 0);
-      if (!res1) break;
+  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;
     }
-    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++) {
-      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;
-      }
+    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++) {
+    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;
     }
   }
-  *isXGMI = false;
   return ncclSuccess;
 }