From 8852e5418137aca9c227cfce12557c45c8a87c90 Mon Sep 17 00:00:00 2001
From: Wenkai Du <wenkai.du@amd.com>
Date: Tue, 3 Mar 2020 11:42:40 -0800
Subject: [PATCH] topo_expl: update to 2.6

[ROCm/rccl commit: 6f54b23503b203b5d01e3be0ae41b8d967e539b8]
---
 projects/rccl/src/graph/topo.cc               |   2 +-
 projects/rccl/tools/topo_expl/Makefile        |   3 +-
 projects/rccl/tools/topo_expl/include/model.h | 425 +++---------------
 projects/rccl/tools/topo_expl/include/nccl.h  |   6 +-
 projects/rccl/tools/topo_expl/include/utils.h |  36 +-
 projects/rccl/tools/topo_expl/model.cpp       |  44 +-
 .../rccl/tools/topo_expl/models/topo_4p1h.xml |  43 ++
 .../tools/topo_expl/models/topo_4p1h_1.xml    |  46 ++
 .../rccl/tools/topo_expl/models/topo_4p2h.xml |  77 ++++
 .../tools/topo_expl/models/topo_4p2h_1.xml    |  77 ++++
 .../tools/topo_expl/models/topo_4p2h_2nic.xml |  82 ++++
 .../rccl/tools/topo_expl/models/topo_4p3l.xml |  52 +++
 .../rccl/tools/topo_expl/models/topo_8p6l.xml | 107 +++++
 .../tools/topo_expl/models/topo_8p_pcie.xml   |  53 +++
 .../tools/topo_expl/models/topo_8p_pcie_1.xml |  53 +++
 .../topo_expl/models/topo_8p_pcie_2nic.xml    |  58 +++
 projects/rccl/tools/topo_expl/topo_expl.cpp   | 126 +++---
 projects/rccl/tools/topo_expl/utils.cpp       | 194 +++++---
 18 files changed, 940 insertions(+), 544 deletions(-)
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_4p1h.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_4p1h_1.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_4p2h.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_4p2h_1.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_4p2h_2nic.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_4p3l.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_8p6l.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_8p_pcie.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_8p_pcie_1.xml
 create mode 100644 projects/rccl/tools/topo_expl/models/topo_8p_pcie_2nic.xml

diff --git a/projects/rccl/src/graph/topo.cc b/projects/rccl/src/graph/topo.cc
index b125527230..d94129cbdf 100644
--- a/projects/rccl/src/graph/topo.cc
+++ b/projects/rccl/src/graph/topo.cc
@@ -284,7 +284,7 @@ ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* s
   if (xmlGetAttrInt(xmlNet, "port", &net->net.port) != ncclSuccess) net->net.port = 0;
   if (xmlGetAttrInt(xmlNet, "gdr", &net->net.gdrSupport) != ncclSuccess) net->net.gdrSupport = 0;
   if (xmlGetAttrInt(xmlNet, "maxconn", &net->net.maxChannels) != ncclSuccess) net->net.maxChannels = MAXCHANNELS;
-  if (xmlGetAttrInt(xmlNet, "coll", &net->net.collSupport) != ncclSuccess) net->net.collSupport = 0;
+  if (ncclCollNet && xmlGetAttrInt(xmlNet, "coll", &net->net.collSupport) != ncclSuccess) net->net.collSupport = 0;
   ncclDebugNoWarn = 0;
 
   NCCLCHECK(ncclTopoConnectNodes(nic, net, LINK_NET, net->net.width));
diff --git a/projects/rccl/tools/topo_expl/Makefile b/projects/rccl/tools/topo_expl/Makefile
index e19ded6d27..a2dc9704d6 100644
--- a/projects/rccl/tools/topo_expl/Makefile
+++ b/projects/rccl/tools/topo_expl/Makefile
@@ -8,7 +8,8 @@ HIPCC = $(HIP_PATH)/bin/hipcc
 EXE = topo_expl
 CXXFLAGS = -g -O3 -Iinclude -I../../src/include -I../../src/graph/ -DTOPO_EXPL -DENABLE_TRACE
 
-files = $(EXE).cpp model.cpp utils.cpp ../../src/graph/topo.cc ../../src/graph/rings.cc ../../src/graph/paths.cc ../../src/graph/trees.cc ../../src/graph/search.cc ../../src/graph/connect.cc
+files = $(EXE).cpp model.cpp utils.cpp ../../src/graph/topo.cc ../../src/graph/rings.cc ../../src/graph/paths.cc ../../src/graph/trees.cc \
+	../../src/graph/search.cc ../../src/graph/connect.cc ../../src/graph/tuning.cc ../../src/graph/xml.cc ../../src/misc/nvmlwrap_stub.cc
 
 all: $(EXE)
 
diff --git a/projects/rccl/tools/topo_expl/include/model.h b/projects/rccl/tools/topo_expl/include/model.h
index 116febad1d..3b02c55ea2 100644
--- a/projects/rccl/tools/topo_expl/include/model.h
+++ b/projects/rccl/tools/topo_expl/include/model.h
@@ -23,397 +23,114 @@ THE SOFTWARE.
 #ifndef MODEL_H_
 #define MODEL_H_
 
-class CpuDevices {
-private:
-  char *cpuName;
-  int interCpuWidth;
-  int cpuPciWidth;
-  int p2pPciWidth;
-
-public:
-  CpuDevices(const char *cpuname, const int intercpuwidth, const int cpupciwidth, const int p2ppciwidth) :
-    cpuName((char *)cpuname), interCpuWidth(intercpuwidth), cpuPciWidth(cpupciwidth), p2pPciWidth(p2ppciwidth) {}
-
-  CpuDevices() : cpuName(0), interCpuWidth(0), cpuPciWidth(0), p2pPciWidth(0) {}
-
-  ncclResult_t getCpuWidths(char* name, int* interCpu, int* cpuPci, int* p2pPci) {
-    strcpy(name, cpuName);
-    *interCpu = interCpuWidth;
-    *cpuPci = cpuPciWidth;
-    *p2pPci = p2pPciWidth;
-    return ncclSuccess;
-  }
-};
-
-class GpuDevices {
-private:
-  int nGpus;
-  uint64_t *busIds;
-  char **gpuPciPaths;
-  int *gpuNumaIds;
-  int *connMatrix;
-
-public:
-  GpuDevices(const int ngpus, const uint64_t *busids, const char **gpupcipaths, const int *gpunumaids, const int *connmatrix) :
-    nGpus(ngpus), busIds((uint64_t *)busids), gpuPciPaths((char **)gpupcipaths), gpuNumaIds((int *)gpunumaids), connMatrix((int *)connmatrix) {}
-
-  GpuDevices () : nGpus(0), busIds(0), gpuPciPaths(0), gpuNumaIds(0), connMatrix(0) {}
-
-  int getnDevs() { return nGpus; }
-
-  uint64_t getBusId(int dev) { return busIds[dev]; }
-
-  ncclResult_t getPciPath(char* busId, char** path) {
-    char tempBusId[] = "0000:00:00.0";
-    *path = (char *)malloc(PATH_MAX);
-    int i;
-    for (i = 0; i < nGpus; i++) {
-      NCCLCHECK(int64ToBusId(busIds[i], tempBusId));
-      if (strcmp(busId, tempBusId) == 0)
-        break;
-    }
-    if (i < nGpus)
-      strcpy(*path, gpuPciPaths[i]);
-    else {
-      WARN("Could not find real path of %s", busId);
-      return ncclSystemError;
-    }
-    return ncclSuccess;
-  }
-
-  int p2pCanConnect(int device1, int device2) {
-    // connection matrix are 8 GPUs
-    int dist = connMatrix[device1*8+device2];
-    if (dist == 255)
-      return 0;
-    //if (dist%15 == 0 && dist/15 != 1) {
-    //  return 0;
-    //}
-    return 1;
-  };
-
-  hipError_t getLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) {
-    // connection matrix are 8 GPUs
-    int dist = connMatrix[device1*8+device2];
-
-    if (dist%15 == 0) {
-      *linktype = 4;
-      *hopcount = dist/15;
-    }
-    else if (dist%20 == 0) {
-      *linktype = 2;
-      *hopcount = dist/20;
-    }
-    else if (dist%36 == 0) {
-      *linktype = 1;
-      *hopcount = dist/36;
-    }
-    return hipSuccess;
-  }
-
-  virtual int getNumaId(char *path) {
-    int n;
-    // search for all GPUs
-    for (n = 0; n < nGpus; n++)
-      if (strcmp(path, gpuPciPaths[n]) == 0)
-        break;
-    if (n < nGpus)
-      return gpuNumaIds[n];
-    return -1;
-  }
-};
-
-class NetDevices {
-private:
-  int nNetDevs;
-  char **netPciPaths;
-  uint64_t *netGuids;    // IB ports on same card share the same GUID
-  int *netNumaIds;
-
-public:
-  NetDevices(const int nnetdevs, const char **netpcipaths, const uint64_t *netguids, const int *netnumaids) :
-    nNetDevs(nnetdevs), netPciPaths((char **)netpcipaths), netGuids((uint64_t *)netguids), netNumaIds((int *)netnumaids) {}
-
-  NetDevices() : nNetDevs(0), netPciPaths(0), netGuids(0), netNumaIds(0) {}
-
-  int getnDevs() { return nNetDevs; }
-
-  ncclResult_t getPciPath(int dev, char** path) {
-    *path = (char *)malloc(PATH_MAX);
-    if (dev < nNetDevs)
-      strcpy(*path, netPciPaths[dev]);
-    else {
-      WARN("Could not find real path of %d", dev);
-      return ncclSystemError;
-    }
-    return ncclSuccess;
-  }
-
-  virtual int getNumaId(char *path) {
-    int n;
-    // search for all NICs
-    for (n = 0; n < nNetDevs; n++)
-      if (strcmp(path, netPciPaths[n]) == 0)
-        break;
-    if (n < nNetDevs)
-      return netNumaIds[n];
-    return -1;
-  }
-
-  uint64_t getIbGuid(char* path) {
-    int n;
-    for (n = 0; n < nNetDevs; n++)
-      if (strcmp(path, netPciPaths[n]) == 0)
-        break;
-    if (n < nNetDevs)
-      return netGuids[n];
-    WARN("Invalid IB path %s", path);
-    return 0;
-  }
-};
+#include <vector>
+#include "topo.h"
+#include "xml.h"
+#include "utils.h"
 
 class NodeModel {
 private:
-  CpuDevices cpus;
-  GpuDevices gpus;
-  NetDevices netdevs;
 
 public:
-  int nodeId;
-  int currRank;
-  int firstRank;
+  std::vector<struct ncclTopoSystem*> systems;
   uint64_t hostHash;  // auto-generated
   uint64_t pidHash;   // auto-generated
-  char description[256];
+  int nodeId;
+  int firstRank;
+  int currRank;
 
-  int rankToCudaDev(int rank) { return rank - firstRank; }
-
-  int getnGpus() { return gpus.getnDevs(); }
-
-  int getnNetDevs() { return netdevs.getnDevs(); }
-
-  ncclResult_t getGpuPciPath(char* busId, char** path) {
-    return gpus.getPciPath(busId, path);
+  NodeModel(const char *xml_file) {
+    char filename[PATH_MAX];
+    ssize_t count = readlink("/proc/self/exe", filename, PATH_MAX);
+    while (--count > 0) {
+      if (filename[count] == '/') {
+        filename[count+1] = 0;
+        break;
+      }
+    };
+    strcat(filename, "models/");
+    strcat(filename, xml_file);
+    struct ncclTopoSystem* system;
+    ncclTopoGetSystem(filename, &system);
+    systems.push_back(system);
+    for (int i=0; i<getNumGpus()-1; i++) {
+      ncclTopoGetSystem(filename, &system);
+      systems.push_back(system);
+    }
+    hostHash = ((uint64_t)rand() << 32) | rand();
+    pidHash = ((uint64_t)rand() << 32) | rand();
   }
 
-  ncclResult_t getNetPciPath(int dev, char** path) {
+  struct ncclTopoSystem* getSystem(int rank) { return systems[rank-firstRank]; }
 
-    return netdevs.getPciPath(dev, path);
+  int getNumGpus() {
+    return systems[0]->nodes[GPU].count;
   }
 
-  uint64_t getGpuBusId(int dev) {
-    return gpus.getBusId(dev);
+  int rankToCudaDev(int rank) {
+    for (int i=0; i<getNumGpus(); i++) {
+      if (rank == systems[0]->nodes[GPU].nodes[i].gpu.rank)
+        return systems[0]->nodes[GPU].nodes[i].gpu.dev;
+    }
+    return -1;
   }
 
-  int p2pCanConnect(int device1, int device2) { return gpus.p2pCanConnect(device1, device2); }
-
-  hipError_t getLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) {
-    return gpus.getLinkTypeAndHopCount(device1, device2, linktype, hopcount);
+  int64_t getGpuBusId(int rank) {
+    for (int i=0; i<getNumGpus(); i++) {
+      if (rank == systems[0]->nodes[GPU].nodes[i].gpu.rank)
+        return systems[0]->nodes[GPU].nodes[i].id;
+    }
+    return -1;
   }
 
-  uint64_t getIbGuid(char* path) {
-    return netdevs.getIbGuid(path);
+  int busIdToCudaDev(int64_t busId) {
+    for (int i=0; i<getNumGpus(); i++)
+      if (systems[0]->nodes[GPU].nodes[i].id == busId)
+        return systems[0]->nodes[GPU].nodes[i].gpu.dev;
+    return -1;
   }
 
+  void setRanks() {
+    for (int r=0; r<getNumGpus(); r++)
+      for (int i=0; i<getNumGpus(); i++)
+        systems[r]->nodes[GPU].nodes[i].gpu.rank += firstRank;
+  }
+
+  int p2pCanConnect(int device1, int device2) { return 1; }
   int shmCanConnect(int device1, int device2) { return 1; }
   int netCanConnect(int device1, int device2) { return 1; }
 
-  virtual int getNumaId(char *path) {
-    int numa = gpus.getNumaId(path);
-    if (numa != -1) return numa;
-    numa = netdevs.getNumaId(path);
-    if (numa != -1) return numa;
-    WARN("Invalid path %s for getNumaId", path);
-    return 0;
-  }
-
-  virtual ncclResult_t getCpuWidths(char* name, int* interCpu, int* cpuPci, int* p2pPci) {
-    return cpus.getCpuWidths(name, interCpu, cpuPci, p2pPci);
-  }
-
-  NodeModel(CpuDevices cpu, GpuDevices gpu, NetDevices net, const char *desc) :
-    cpus(cpu), gpus(gpu), netdevs(net) {
-      strncpy(description, desc, 256);
-  }
-
-  NodeModel() {}
-
   ~NodeModel() {}
 };
 
 class NetworkModel {
 private:
-  int nNodes;
   int nRanks;
-  NodeModel nodes[NCCL_TOPO_MAX_NODES];
+  std::vector<NodeModel*> nodes;
 
 public:
-  void AddNode(NodeModel node) {
-    nodes[nNodes] = node;
-    nodes[nNodes].nodeId = nNodes;
-    nodes[nNodes].firstRank = nRanks;
-    nodes[nNodes].hostHash = ((uint64_t)rand() << 32) | rand();
-    nodes[nNodes].pidHash = ((uint64_t)rand() << 32) | rand();
-    nNodes++;
-    nRanks += node.getnGpus();
+  void AddNode(NodeModel* node) {
+    node->nodeId = nodes.size();
+    node->firstRank = nRanks;
+    node->setRanks();
+    nRanks += node->getNumGpus();
+    nodes.push_back(node);
   }
 
-  int GetNNodes() { return nNodes; }
-
-  int GetNRanks() { return nRanks; }
-
   NodeModel* GetNode(int rank) {
-    int node_id;
-
-    if(rank < 0 || rank >= nRanks)
-      return 0;
-
-    for(node_id = nNodes-1; node_id >= 0; node_id--)
-      if(rank >= nodes[node_id].firstRank) break;
-
-    if (node_id >= 0) {
-      nodes[node_id].currRank = rank;
-      return nodes+node_id;
+    for (auto & node : nodes) {
+      if (rank >= node->firstRank && rank < node->firstRank+node->getNumGpus()) {
+        node->currRank = rank;
+        return node;
+      }
     }
-    else
-      return 0;
+    return NULL;
   }
 
-  NetworkModel() : nNodes(0), nRanks(0) {}
-};
+  int GetNNodes() { return nodes.size(); }
+  int GetNRanks() { return nRanks; }
 
-
-const static uint64_t busIds_8[] = { 0x1d000, 0x20000, 0x23000, 0x26000, 0x3f000, 0x43000, 0x46000, 0x49000 };
-
-const static char* gpuPciPaths_8[] = {
-  "/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:08.0/0000:1b:00.0/0000:1c:00.0/0000:1d:00.0",
-  "/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:0c.0/0000:1e:00.0/0000:1f:00.0/0000:20:00.0",
-  "/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:10.0/0000:21:00.0/0000:22:00.0/0000:23:00.0",
-  "/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:14.0/0000:24:00.0/0000:25:00.0/0000:26:00.0",
-  "/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:04.0/0000:3d:00.0/0000:3e:00.0/0000:3f:00.0",
-  "/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:0c.0/0000:41:00.0/0000:42:00.0/0000:43:00.0",
-  "/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:10.0/0000:44:00.0/0000:45:00.0/0000:46:00.0",
-  "/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:14.0/0000:47:00.0/0000:48:00.0/0000:49:00.0",
-};
-
-const static int gpuPciNumaIds_8[] = { 0, 0, 0, 0, 0, 0, 0, 0 };
-
-const static char* netPciPaths_1[] = {
-  "/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:04.0/0000:1a:00.0",
-};
-
-const static char* netPciPaths_1_1[] = {
-  "/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:08.0/0000:4c:00.0",
-};
-
-const static uint64_t netGuids_1[] = {
-  0xb8599f030007053aL,
-};
-
-const static int netPciNumaIds_1[] = { 0 };
-
-const static char* netPciPaths_2[] = {
-  "/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:04.0/0000:1a:00.0",
-  "/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:08.0/0000:4c:00.0",
-};
-
-const static uint64_t netGuids_2[] = {
-  0xb8599f030007053aL,
-  0x506b4b030027bbf2L,
-};
-
-const static int netPciNumaIds_2[] = { 0, 0 };
-
-const static uint64_t rome_busIds_8[] = { 0x63000, 0x23000, 0x26000, 0x03000, 0xe3000, 0xc3000, 0xc6000, 0xa3000 };
-
-const static char* rome_gpuPciPaths_8[] = {
-  "/sys/devices/pci0000:60/0000:60:03.1/0000:61:00.0/0000:62:00.0/0000:63:00.0",
-  "/sys/devices/pci0000:20/0000:20:01.1/0000:21:00.0/0000:22:00.0/0000:23:00.0",
-  "/sys/devices/pci0000:20/0000:20:03.1/0000:24:00.0/0000:25:00.0/0000:26:00.0",
-  "/sys/devices/pci0000:00/0000:00:01.1/0000:01:00.0/0000:02:00.0/0000:03:00.0",
-  "/sys/devices/pci0000:e0/0000:e0:03.1/0000:e1:00.0/0000:e2:00.0/0000:e3:00.0",
-  "/sys/devices/pci0000:c0/0000:c0:01.1/0000:c1:00.0/0000:c2:00.0/0000:c3:00.0",
-  "/sys/devices/pci0000:c0/0000:c0:03.1/0000:c4:00.0/0000:c5:00.0/0000:c6:00.0",
-  "/sys/devices/pci0000:a0/0000:a0:03.1/0000:a1:00.0/0000:a2:00.0/0000:a3:00.0",
-};
-
-const static int rome_gpuPciNumaIds_8[] = { 0, 0, 0, 0, 4, 4, 4, 4 };
-
-const static char* rome_netPciPaths_1[] = {
-  "/sys/devices/pci0000:40/0000:40:01.1/0000:41:00.0",
-};
-
-const static uint64_t rome_netGuids_1[] = {
-  0xb8599f030007053aL,
-};
-
-const static int rom_netPciNumaIds_1[] = { 0 };
-
-const static char* rome_netPciPaths_2[] = {
-  "/sys/devices/pci0000:40/0000:40:01.1/0000:41:00.0",
-  "/sys/devices/pci0000:80/0000:80:01.1/0000:81:00.0",
-};
-
-const static uint64_t rome_netGuids_2[] = {
-  0xb8599f030007053aL,
-  0x506b4b030027bbf2L,
-};
-
-const static int rom_netPciNumaIds_2[] = { 0, 4 };
-
-const int conn_mat_pcie[64] = {
-  0 , 40, 40, 40, 40, 40, 40, 40,
-  40, 0 , 40, 40, 40, 40, 40, 40,
-  40, 40, 0 , 40, 40, 40, 40, 40,
-  40, 40, 40, 0 , 40, 40, 40, 40,
-  40, 40, 40, 40, 0 , 40, 40, 40,
-  40, 40, 40, 40, 40, 0 , 40, 40,
-  40, 40, 40, 40, 40, 40, 0 , 40,
-  40, 40, 40, 40, 40, 40, 40, 0 ,
-};
-
-const int conn_mat_4p2h[64] = {
-  0 , 15, 15, 30, 40, 40, 40, 40,
-  15, 0 , 30, 15, 40, 40, 40, 40,
-  15, 30, 0 , 15, 40, 40, 40, 40,
-  30, 15, 15, 0 , 40, 40, 40, 40,
-  40, 40, 40, 40, 0 , 15, 15, 30,
-  40, 40, 40, 40, 15, 0 , 30, 15,
-  40, 40, 40, 40, 15, 30, 0 , 15,
-  40, 40, 40, 40, 30, 15, 15, 0 ,
-};
-
-const int conn_mat_8p6l[64] = {
-  0 , 15, 15, 15, 15, 30, 15, 15,
-  15, 0 , 15, 15, 30, 15, 15, 15,
-  15, 15, 0 , 15, 15, 15, 15, 30,
-  15, 15, 15, 0 , 15, 15, 30, 15,
-  15, 30, 15, 15, 0 , 15, 15, 15,
-  30, 15, 15, 15, 15, 0 , 15, 15,
-  15, 15, 15, 30, 15, 15, 0 , 15,
-  15, 15, 30, 15, 15, 15, 15, 0 ,
-};
-
-const int conn_mat_8p6l_1[64] = {
- 0 , 15, 15, 30, 15, 15, 15, 15,
- 15, 0 , 30, 15, 15, 15, 15, 15,
- 15, 30, 0 , 15, 15, 15, 15, 15,
- 30, 15, 15, 0 , 15, 15, 15, 15,
- 15, 15, 15, 15, 0 , 15, 15, 30,
- 15, 15, 15, 15, 15, 0 , 30, 15,
- 15, 15, 15, 15, 15, 30, 0 , 15,
- 15, 15, 15, 15, 30, 15, 15, 0 ,
-};
-
-const int conn_mat_rome[64] = {
-  0 , 40, 40, 40, 72, 72, 72, 72,
-  40, 0 , 40, 40, 72, 72, 72, 72,
-  40, 40, 0 , 40, 72, 72, 72, 72,
-  40, 40, 40, 0 , 72, 72, 72, 72,
-  72, 72, 72, 72, 0 , 40, 40, 40,
-  72, 72, 72, 72, 40, 0 , 40, 40,
-  72, 72, 72, 72, 40, 40, 0 , 40,
-  72, 72, 72, 72, 40, 40, 40, 0 ,
+  NetworkModel() : nRanks(0) {}
 };
 
 #endif
\ No newline at end of file
diff --git a/projects/rccl/tools/topo_expl/include/nccl.h b/projects/rccl/tools/topo_expl/include/nccl.h
index 23ffec5283..225b150127 100644
--- a/projects/rccl/tools/topo_expl/include/nccl.h
+++ b/projects/rccl/tools/topo_expl/include/nccl.h
@@ -12,11 +12,11 @@
 #include <hip/hip_fp16.h>
 
 #define NCCL_MAJOR 2
-#define NCCL_MINOR 5
-#define NCCL_PATCH 6
+#define NCCL_MINOR 6
+#define NCCL_PATCH 2
 #define NCCL_SUFFIX ""
 
-#define NCCL_VERSION_CODE 2506
+#define NCCL_VERSION_CODE 2602
 #define NCCL_VERSION(X,Y,Z) ((X) * 1000 + (Y) * 100 + (Z))
 
 #define RCCL_BFLOAT16 1
diff --git a/projects/rccl/tools/topo_expl/include/utils.h b/projects/rccl/tools/topo_expl/include/utils.h
index 26e85c26d6..9337e18aee 100644
--- a/projects/rccl/tools/topo_expl/include/utils.h
+++ b/projects/rccl/tools/topo_expl/include/utils.h
@@ -13,32 +13,34 @@ struct allGather1Data_t {
   struct ncclComm* comm;
 };
 
-struct allGather3Data_t {
+// AllGather3 - begin
+struct ncclGraphInfo {
+  int sameChannels;
+  float speedIntra;
+  float speedInter;
+  int typeIntra;
+};
+
+struct allGather3Data_t{
   int cudaCompCap;
   int fullCudaCompCap;
-  int nvlink;
   int nChannels;
-  struct {
-    int sameChannels;
-    int speedIntra;
-    int speedInter;
-    int nvlink;
-  } tree;
-  struct {
-    int sameChannels;
-    int speedIntra;
-    int speedInter;
-    int nvlink;
-  } ring;
+  struct ncclGraphInfo tree;
+  struct ncclGraphInfo ring;
+  struct ncclGraphInfo collNet;
   struct ncclTopoRanks topoRanks;
 };
 
+ncclResult_t ncclTopoGetSystem(const char* xmlTopoFile, struct ncclTopoSystem** system);
+
+ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem** topoSystem);
+
 ncclResult_t bootstrapAllGather(struct ncclComm* comm, struct allGather1Data_t * allGather1Data);
 
-ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data,
-  struct allGather3Data_t *allGather3Data, struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph);
+ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data, struct allGather3Data_t *allGather3Data,
+  struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph);
 
 ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t *allGather3Data,
-  struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph);
+  struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph);
 
 #endif
\ No newline at end of file
diff --git a/projects/rccl/tools/topo_expl/model.cpp b/projects/rccl/tools/topo_expl/model.cpp
index 9492526dc9..db5eb881bf 100644
--- a/projects/rccl/tools/topo_expl/model.cpp
+++ b/projects/rccl/tools/topo_expl/model.cpp
@@ -41,24 +41,14 @@ THE SOFTWARE.
 #include <sys/stat.h>
 #include <unistd.h>
 #include "model.h"
+#include "topo.h"
 
 extern NodeModel *node_model;
 
-static ncclResult_t dummyNetDevices(int* ndev) {
-  *ndev = node_model->getnNetDevs();
-  return ncclSuccess;
-}
-
-static ncclResult_t dummyNetPciPath(int dev, char** path) {
-  node_model->getNetPciPath(dev, path);
-  return ncclSuccess;
-}
-
 ncclNet_t ncclNetDummy = {
   "IB",
   0,
-  dummyNetDevices,
-  dummyNetPciPath,
+  0,
   0,
   0,
   0,
@@ -76,24 +66,9 @@ ncclNet_t ncclNetDummy = {
 
 ncclNet_t* ncclNet = &ncclNetDummy;
 
-ncclResult_t wrapNvmlDeviceGetHandleByPciBusId(const char* pciBusId, nvmlDevice_t* device) {
-  return ncclSuccess;
-}
-
 /* Convert a PCI busId string into a local cudaDev device index (cf. CUDA_VISIBLE_DEVICES) */
 int busIdToCudaDev(int64_t busId) {
-  int cudaDev;
-
-  for (cudaDev = 0; cudaDev < node_model->getnGpus(); cudaDev++) {
-    if (node_model->getGpuBusId(cudaDev) == busId)
-      break;
-  }
-
-  if (cudaDev < node_model->getnGpus())
-    return cudaDev;
-  else
-    WARN("Invalid busId %lx", busId);
-  return 0;
+  return node_model->busIdToCudaDev(busId);
 }
 
 /* Determine if two peers can communicate with P2P */
@@ -177,6 +152,8 @@ ncclResult_t netSendSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph* gra
   int netDev, useGdr = 0;
 
   NCCLCHECK(ncclTopoGetNetDev(graph, 1, channelId, &netDev));
+  NCCLCHECK(ncclTopoCheckGdr(topo, myInfo->busId, netDev, 1, &useGdr));
+
   INFO(NCCL_INIT|NCCL_NET,"Ring %02d : %d[%lx] -> %d[%lx] [send] via NET/%s/%d%s", channelId, myInfo->rank, myInfo->busId, peerInfo->rank, peerInfo->busId, ncclNetName(), netDev,
       useGdr ? "/GDRDMA" : "");
   return ncclSuccess;
@@ -188,15 +165,8 @@ ncclResult_t netRecvSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph* gra
   int netDev, useGdr = 0;
 
   NCCLCHECK(ncclTopoGetNetDev(graph, 0, channelId, &netDev));
-  // Check if we are close enough that it makes sense to enable GDR
-  int netGdrLevel = ncclParamNetGdrLevel();
-  int distance;
-  NCCLCHECK(ncclTopoNetDistance(topo, myInfo->busId, netDev, &distance));
-  if (distance >= netGdrLevel) {
-    INFO(NCCL_NET,"NET/%s : GPU Direct RDMA Disabled for GPU %lx / HCA %d (distance %d >= %d)", ncclNetName(), myInfo->busId, netDev, distance, netGdrLevel);
-  }
-  else
-    useGdr = 1;
+  NCCLCHECK(ncclTopoCheckGdr(topo, myInfo->busId, netDev, 0, &useGdr));
+
   INFO(NCCL_INIT|NCCL_NET,"Ring %02d : %d[%lx] -> %d[%lx] [receive] via NET/%s/%d%s", channelId, peerInfo->rank, peerInfo->busId, myInfo->rank, myInfo->busId, ncclNetName(), netDev,
       useGdr ? "/GDRDMA" : "");
   return ncclSuccess;
diff --git a/projects/rccl/tools/topo_expl/models/topo_4p1h.xml b/projects/rccl/tools/topo_expl/models/topo_4p1h.xml
new file mode 100644
index 0000000000..6878edf115
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_4p1h.xml
@@ -0,0 +1,43 @@
+<system version="1">
+  <cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="0" sm="30" rank="0" gdr="1">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="1" sm="30" rank="1" gdr="1">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="2" sm="30" rank="2" gdr="1">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="3" sm="30" rank="3" gdr="1">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:1a:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_4p1h_1.xml b/projects/rccl/tools/topo_expl/models/topo_4p1h_1.xml
new file mode 100644
index 0000000000..bc7dcaa606
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_4p1h_1.xml
@@ -0,0 +1,46 @@
+<system version="1">
+  <cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1a:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="0" sm="30" rank="0" gdr="0">
+          <xgmi target="0000:3d:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:b1:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:3d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="1" sm="30" rank="1" gdr="0">
+          <xgmi target="0000:1a:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:88:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:01:00.0" class="0x020000" link_speed="8 GT/s" link_width="4">
+      <nic>
+        <net name="eno1" dev="0" speed="10000" port="0" guid="0x0" maxconn="65536" gdr="0"/>
+      </nic>
+    </pci>
+    <nic>
+      <net name="virbr0" dev="1" speed="10000" port="0" guid="0x1" maxconn="65536" gdr="0"/>
+    </nic>
+  </cpu>
+  <cpu numaid="1" affinity="fffc00,0fffc000" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:86:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:88:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="2" sm="30" rank="2" gdr="0">
+          <xgmi target="0000:3d:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:b1:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:af:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:b1:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="3" sm="30" rank="3" gdr="0">
+          <xgmi target="0000:1a:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:88:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_4p2h.xml b/projects/rccl/tools/topo_expl/models/topo_4p2h.xml
new file mode 100644
index 0000000000..03ce1dff5c
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_4p2h.xml
@@ -0,0 +1,77 @@
+<system version="1">
+  <cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="0" sm="30" rank="0" gdr="1">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="1" sm="30" rank="1" gdr="1">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="2" sm="30" rank="2" gdr="1">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="3" sm="30" rank="3" gdr="1">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:1a:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+    <pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="4" sm="30" rank="4" gdr="1">
+            <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="5" sm="30" rank="5" gdr="1">
+            <xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="6" sm="30" rank="6" gdr="1">
+            <xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="7" sm="30" rank="7" gdr="1">
+            <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_4p2h_1.xml b/projects/rccl/tools/topo_expl/models/topo_4p2h_1.xml
new file mode 100644
index 0000000000..fb2dea7d9b
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_4p2h_1.xml
@@ -0,0 +1,77 @@
+<system version="1">
+  <cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="0" sm="30" rank="0" gdr="1">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="1" sm="30" rank="1" gdr="1">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="2" sm="30" rank="2" gdr="1">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="3" sm="30" rank="3" gdr="1">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+    </pci>
+    <pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="4" sm="30" rank="4" gdr="1">
+            <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="5" sm="30" rank="5" gdr="1">
+            <xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="6" sm="30" rank="6" gdr="1">
+            <xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="7" sm="30" rank="7" gdr="1">
+            <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:4c:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_4p2h_2nic.xml b/projects/rccl/tools/topo_expl/models/topo_4p2h_2nic.xml
new file mode 100644
index 0000000000..9ec580f726
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_4p2h_2nic.xml
@@ -0,0 +1,82 @@
+<system version="1">
+  <cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="0" sm="30" rank="0" gdr="1">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="1" sm="30" rank="1" gdr="1">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="2" sm="30" rank="2" gdr="1">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="3" sm="30" rank="3" gdr="1">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:1a:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+    <pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="4" sm="30" rank="4" gdr="1">
+            <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="5" sm="30" rank="5" gdr="1">
+            <xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="6" sm="30" rank="6" gdr="1">
+            <xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="7" sm="30" rank="7" gdr="1">
+            <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:4c:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_1" dev="1" speed="100000" port="1" guid="0xb8599f030007053a" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_4p3l.xml b/projects/rccl/tools/topo_expl/models/topo_4p3l.xml
new file mode 100644
index 0000000000..0549d11388
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_4p3l.xml
@@ -0,0 +1,52 @@
+<system version="1">
+  <cpu numaid="0" affinity="0000ffff,0000ffff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="0" sm="30" rank="0" gdr="0">
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="1" sm="30" rank="1" gdr="0">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="2" sm="30" rank="2" gdr="0">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+      <pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="3" sm="30" rank="3" gdr="0">
+            <xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
+            <xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
+          </gpu>
+        </pci>
+      </pci>
+    </pci>
+    <pci busid="0000:5e:00.0" class="0x060400" link_speed="8 GT/s" link_width="8">
+      <pci busid="0000:60:00.0" class="0x020000" link_speed="2.5 GT/s" link_width="1">
+        <nic>
+          <net name="enp96s0f0" dev="0" speed="10000" port="0" guid="0x0" maxconn="65536" gdr="0"/>
+        </nic>
+      </pci>
+    </pci>
+    <nic>
+      <net name="veth608058c" dev="1" speed="10000" port="0" guid="0x1" maxconn="65536" gdr="0"/>
+    </nic>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_8p6l.xml b/projects/rccl/tools/topo_expl/models/topo_8p6l.xml
new file mode 100644
index 0000000000..e54ac50a1b
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_8p6l.xml
@@ -0,0 +1,107 @@
+<system version="1">
+  <cpu numaid="0" affinity="00000000,00000000,00000000,00000000,00000000,000000ff,ffff0000,00ffffff" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
+    <pci busid="0000:61:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
+      <pci busid="0000:63:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="0" sm="30" rank="0" gdr="0">
+          <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:41:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
+      <pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="1" sm="30" rank="1" gdr="0">
+          <xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:22:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
+      <pci busid="0000:24:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="2" sm="30" rank="2" gdr="0">
+          <xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:01:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
+      <pci busid="0000:03:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="3" sm="30" rank="3" gdr="0">
+          <xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+  </cpu>
+  <cpu numaid="1" affinity="00000000,00000000,00000000,00000000,00000000,ffffff00,0000ffff,ff000000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
+    <pci busid="0000:e1:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
+      <pci busid="0000:e3:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="4" sm="30" rank="4" gdr="0">
+          <xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:c1:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
+      <pci busid="0000:c3:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="5" sm="30" rank="5" gdr="0">
+          <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:a1:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
+      <pci busid="0000:a3:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="6" sm="30" rank="6" gdr="0">
+          <xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:81:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
+      <pci busid="0000:83:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+        <gpu dev="7" sm="30" rank="7" gdr="0">
+          <xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
+          <xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
+        </gpu>
+      </pci>
+    </pci>
+    <pci busid="0000:c4:00.0" class="0x020000" link_speed="2.5 GT/s" link_width="1">
+      <nic>
+        <net name="enp196s0" dev="0" speed="10000" port="0" guid="0x0" maxconn="65536" gdr="0"/>
+      </nic>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_8p_pcie.xml b/projects/rccl/tools/topo_expl/models/topo_8p_pcie.xml
new file mode 100644
index 0000000000..521d1ab695
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_8p_pcie.xml
@@ -0,0 +1,53 @@
+<system version="1">
+  <cpu numaid="0" affinity="00ff00ff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="0" sm="30" rank="0" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="1" sm="30" rank="1" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="2" sm="30" rank="2" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="3" sm="30" rank="3" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:1a:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_0" dev="0" speed="100000" port="1" guid="0x3a050700039f59b8" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+    <pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="4" sm="30" rank="4" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="5" sm="30" rank="5" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="6" sm="30" rank="6" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="7" sm="30" rank="7" gdr="1"/>
+        </pci>
+      </pci>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_8p_pcie_1.xml b/projects/rccl/tools/topo_expl/models/topo_8p_pcie_1.xml
new file mode 100644
index 0000000000..350f11f816
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_8p_pcie_1.xml
@@ -0,0 +1,53 @@
+<system version="1">
+  <cpu numaid="0" affinity="00ff00ff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="0" sm="30" rank="0" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="1" sm="30" rank="1" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="2" sm="30" rank="2" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="3" sm="30" rank="3" gdr="1"/>
+        </pci>
+      </pci>
+    </pci>
+    <pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="4" sm="30" rank="4" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="5" sm="30" rank="5" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="6" sm="30" rank="6" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="7" sm="30" rank="7" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:4c:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xb8599f030007053a" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/models/topo_8p_pcie_2nic.xml b/projects/rccl/tools/topo_expl/models/topo_8p_pcie_2nic.xml
new file mode 100644
index 0000000000..dae7b1c0f1
--- /dev/null
+++ b/projects/rccl/tools/topo_expl/models/topo_8p_pcie_2nic.xml
@@ -0,0 +1,58 @@
+<system version="1">
+  <cpu numaid="0" affinity="00ff00ff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
+    <pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="0" sm="30" rank="0" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="1" sm="30" rank="1" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="2" sm="30" rank="2" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="3" sm="30" rank="3" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:1a:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_0" dev="0" speed="100000" port="1" guid="0x3a050700039f59b8" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+    <pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+      <pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="4" sm="30" rank="4" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="5" sm="30" rank="5" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="6" sm="30" rank="6" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
+        <pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
+          <gpu dev="7" sm="30" rank="7" gdr="1"/>
+        </pci>
+      </pci>
+      <pci busid="0000:4c:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
+        <nic>
+          <net name="mlx5_1" dev="1" speed="100000" port="1" guid="0xb8599f030007053a" maxconn="262144" gdr="1"/>
+        </nic>
+      </pci>
+    </pci>
+  </cpu>
+</system>
diff --git a/projects/rccl/tools/topo_expl/topo_expl.cpp b/projects/rccl/tools/topo_expl/topo_expl.cpp
index 98f175cccc..4f2df2ef4e 100644
--- a/projects/rccl/tools/topo_expl/topo_expl.cpp
+++ b/projects/rccl/tools/topo_expl/topo_expl.cpp
@@ -29,7 +29,6 @@ THE SOFTWARE.
 #include "net.h"
 #include "graph.h"
 #include "argcheck.h"
-#include "cpuset.h"
 #include <sched.h>
 #include <fcntl.h>
 #include <unistd.h>
@@ -65,18 +64,18 @@ bool cmdOptionExists(char** begin, char** end, const std::string& option) {
 
 const char *model_descriptions[] = {
   "single node VEGA20 4P1H",
+  "single node VEGA20 4P1H Alt. Model",
   "single node VEGA20 4P2H",
   "single node gfx908 4P3L",
   "single node gfx908 8P6L",
-  "single node gfx908 8P6L Alt. Connection",
-  "single node 8 VEGA20 PCIe on Rome",
-  "single node gfx908 8P6L on Rome",
+  "single node 8 VEGA20 PCIe",
   "4 nodes with 8 GPUs PCIe 1 NIC",
+  "4 nodes with 8 GPUs PCIe 1 NIC 2nd PLX Bridge",
   "4 nodes with 8 GPUs PCIe 2 NIC",
   "2 nodes VEGA20 4P1H",
   "4 nodes with 8 VEGA20 GPUs XGMI 4P2H 1 NIC",
-  "4 nodes 8 GPUs PCIe 2 NICs on Rome",
-  "3 nodes 8 GPUs PCIe + 1 Rome 8 GPUs PCIe + 2 nodes gfx908 4P3L",
+  "4 nodes with 8 VEGA20 GPUs XGMI 4P2H 1 NIC 2nd Hive",
+  "4 nodes with 8 VEGA20 GPUs XGMI 4P2H 2 NIC",
   NULL,
 };
 
@@ -97,97 +96,75 @@ int main(int argc,char* argv[])
   if (mi)
     model_id = atol(mi);
 
-  // CPU, GPU and NIC devices on Skylake
-  CpuDevices skylake("Skylake", SKL_QPI_WIDTH, SKL_CPUPCI_WIDTH, SKL_PCI_WIDTH);
-  GpuDevices vg20_pcie(8, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_pcie);
-  GpuDevices vg20_4p1h(4, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_4p2h);
-  GpuDevices vg20_4p2h(8, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_4p2h);
-  GpuDevices gfx908_4p3l(4, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_8p6l);
-  GpuDevices gfx908_8p6l(8, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_8p6l);
-  GpuDevices gfx908_8p6l_1(8, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_8p6l_1);
-  NetDevices nic_1(1, netPciPaths_1, netGuids_1, netPciNumaIds_1);
-  NetDevices nic_1_1(1, netPciPaths_1_1, netGuids_1, netPciNumaIds_1);
-  NetDevices nic_2(2, netPciPaths_2, netGuids_2, netPciNumaIds_2);
-
-  // CPU, GPU and NIC devices on Rome
-  CpuDevices rome("Rome", ROME_QPI_WIDTH, ROME_CPUPCI_WIDTH, ROME_PCI_WIDTH);
-  GpuDevices vg20_pcie_rome(8, rome_busIds_8, rome_gpuPciPaths_8, rome_gpuPciNumaIds_8, conn_mat_rome);
-  NetDevices nic_1_rome(1, rome_netPciPaths_1, rome_netGuids_1, rom_netPciNumaIds_1);
-  NetDevices nic_2_rome(2, rome_netPciPaths_2, rome_netGuids_2, rom_netPciNumaIds_2);
-
-  // 8 GPUs PCIe 1 NIC
-  NodeModel model_8pcie_1nic(skylake, vg20_pcie, nic_1, "Skylake 8 GPUs PCIe");
-
-  // 8 GPUs PCIe 2 NIC
-  NodeModel model_8pcie_2nic(skylake, vg20_pcie, nic_2, "Skylake 8 GPUs PCIe 2 NIC");
-
-  // VEGA20 4P1H, use VEGA20 4P2H model
-  NodeModel model_vg20_4p1h_1nic(skylake, vg20_4p1h, nic_1, "Skylake VEGA20 4P1H");
-
-  // VEGA20 GPUs XGMI 4P2H 1 NIC
-  NodeModel model_vg20_4p2h_1nic(skylake, vg20_4p2h, nic_1_1, "Skylake VEGA20 4P2H");
-
-  // gfx908 4P3L
-  NodeModel model_gfx908_4p_1nic(skylake, gfx908_4p3l, nic_1, "Skylake gfx908 4P3L");
-
-  // gfx908 8P6L
-  NodeModel model_gfx908_8p_1nic(skylake, gfx908_8p6l, nic_1, "Skylake gfx908 8P6L");
-
-  // gfx908 8P6L alternative connection
-  NodeModel model_gfx908_8p_1nic_1(skylake, gfx908_8p6l_1, nic_1, "Skylake gfx908 8P6L Alt. Connection");
-
-  // 8 GPUs PCIe on Rome
-  NodeModel model_8pcie_1nic_rome(rome, vg20_pcie_rome, nic_1_rome, "Rome 8 GPUs PCIe");
-
-  // 8 GPUs PCIe 2 NICs on Rome
-  NodeModel model_8pcie_2nic_rome(rome, vg20_pcie_rome, nic_2_rome, "Rome 8 GPUs PCIe 2 NICs");
-
-  // gfx908 8P6L on Rome
-  NodeModel model_gfx908_8p_1nic_rome(rome, gfx908_8p6l, nic_1, "Rome gfx908 8P6L");
-
   NetworkModel network;
+  NodeModel* node;
 
   switch(model_id) {
     case 0:
-      network.AddNode(model_vg20_4p1h_1nic);
+      node = new NodeModel("topo_4p1h.xml");
+      network.AddNode(node);
       break;
     case 1:
-      network.AddNode(model_vg20_4p2h_1nic);
+      node = new NodeModel("topo_4p1h_1.xml");
+      network.AddNode(node);
       break;
     case 2:
-      network.AddNode(model_gfx908_4p_1nic);
+      node = new NodeModel("topo_4p2h.xml");
+      network.AddNode(node);
       break;
     case 3:
-      network.AddNode(model_gfx908_8p_1nic);
+      node = new NodeModel("topo_4p3l.xml");
+      network.AddNode(node);
       break;
     case 4:
-      network.AddNode(model_gfx908_8p_1nic_1);
+      node = new NodeModel("topo_8p6l.xml");
+      network.AddNode(node);
       break;
     case 5:
-      network.AddNode(model_8pcie_1nic_rome);
+      node = new NodeModel("topo_8p_pcie.xml");
+      network.AddNode(node);
       break;
     case 6:
-      network.AddNode(model_gfx908_8p_1nic_rome);
+      for (int i=0; i<4; i++) {
+        node = new NodeModel("topo_8p_pcie.xml");
+        network.AddNode(node);
+      }
       break;
     case 7:
-      for (int i = 0; i < 4; i ++) network.AddNode(model_8pcie_1nic);
+      for (int i=0; i<4; i++) {
+        node = new NodeModel("topo_8p_pcie_1.xml");
+        network.AddNode(node);
+      }
       break;
     case 8:
-      for (int i = 0; i < 4; i ++) network.AddNode(model_8pcie_2nic);
+      for (int i=0; i<4; i++) {
+        node = new NodeModel("topo_8p_pcie_2nic.xml");
+        network.AddNode(node);
+      }
       break;
     case 9:
-      for (int i = 0; i < 2; i ++) network.AddNode(model_vg20_4p1h_1nic);
+      for (int i=0; i<2; i++) {
+        node = new NodeModel("topo_4p1h.xml");
+        network.AddNode(node);
+      }
       break;
     case 10:
-      for (int i = 0; i < 4; i ++) network.AddNode(model_vg20_4p2h_1nic);
+      for (int i=0; i<4; i++) {
+        node = new NodeModel("topo_4p2h.xml");
+        network.AddNode(node);
+      }
       break;
     case 11:
-      for (int i = 0; i < 4; i ++) network.AddNode(model_8pcie_2nic_rome);
+      for (int i=0; i<4; i++) {
+        node = new NodeModel("topo_4p2h_1.xml");
+        network.AddNode(node);
+      }
       break;
     case 12:
-      for (int i = 0; i < 3; i ++) network.AddNode(model_8pcie_1nic);
-      network.AddNode(model_8pcie_1nic_rome);
-      for (int i = 0; i < 2; i ++) network.AddNode(model_gfx908_4p_1nic);
+      for (int i=0; i<4; i++) {
+        node = new NodeModel("topo_4p2h_2nic.xml");
+        network.AddNode(node);
+      }
       break;
     default:
       printf("Invalid model_id %d\n", model_id);
@@ -203,8 +180,8 @@ int main(int argc,char* argv[])
   for (int i = 0; i < nranks; i++) {
     node_model = network.GetNode(i);
     assert(node_model!=0);
-    printf("Rank %d: node %d (%s) GPU busId %lx\n", i, node_model->nodeId,
-      node_model->description, node_model->getGpuBusId(node_model->rankToCudaDev(i)));
+    printf("Rank %d: node %d cudaDev %d GPU busId %lx\n", i, node_model->nodeId,
+      node_model->rankToCudaDev(i), node_model->getGpuBusId(i));
   }
 
   NCCLCHECK(ncclCalloc(&comm, nranks));
@@ -220,21 +197,22 @@ int main(int argc,char* argv[])
     comm[i].nRanks = nranks;
     node_model = network.GetNode(i);
     assert(node_model!=0);
+    comm[i].topo = node_model->getSystem(i);
     bootstrapAllGather(&comm[i], allGather1Data);
   }
 
-  struct ncclTopoGraph treeGraph, ringGraph;
+  struct ncclTopoGraph treeGraph, ringGraph, collNetGraph;
 
   for (int i = 0; i < nranks; i++) {
     node_model = network.GetNode(i);
     assert(node_model!=0);
-    initTransportsRank_1(&comm[i], allGather1Data, allGather3Data, treeGraph, ringGraph);
+    initTransportsRank_1(&comm[i], allGather1Data, allGather3Data, treeGraph, ringGraph, collNetGraph);
   }
 
   for (int i = 0; i < nranks; i++) {
     node_model = network.GetNode(i);
     assert(node_model!=0);
-    initTransportsRank_3(&comm[i], allGather3Data, treeGraph, ringGraph);
+    initTransportsRank_3(&comm[i], allGather3Data, treeGraph, ringGraph, collNetGraph);
   }
 
   free(allGather3Data);
diff --git a/projects/rccl/tools/topo_expl/utils.cpp b/projects/rccl/tools/topo_expl/utils.cpp
index 18d61d0b99..942b4a8a66 100644
--- a/projects/rccl/tools/topo_expl/utils.cpp
+++ b/projects/rccl/tools/topo_expl/utils.cpp
@@ -25,12 +25,19 @@
 #include <sys/types.h>
 #include <sys/stat.h>
 #include <unistd.h>
+#include "xml.h"
+#include "coll_net.h"
 #include "model.h"
 #include "utils.h"
 
 extern NodeModel *node_model;
 
 NCCL_PARAM(CrossNic, "CROSS_NIC", 2);
+NCCL_PARAM(CollNetEnable, "COLLNET_ENABLE", 0);
+NCCL_PARAM(GraphDumpFileRank, "GRAPH_DUMP_FILE_RANK", 0);
+
+thread_local int ncclDebugNoWarn = 0;
+ncclCollNet_t* ncclCollNet = NULL;
 
 // Get current Compute Capability
 int ncclCudaCompCap() {
@@ -43,7 +50,7 @@ ncclResult_t int64ToBusId(int64_t id, char* busId) {
   return ncclSuccess;
 }
 
-ncclResult_t busIdToInt64(char* busId, int64_t* id) {
+ncclResult_t busIdToInt64(const char* busId, int64_t* id) {
   const int size = strlen(busId);
   char* hexStr;
   NCCLCHECK(ncclCalloc(&hexStr, size));
@@ -87,9 +94,9 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long flags, const char *file
   if (ncclDebugLevel == -1) ncclDebugInit();
   if (level == NCCL_LOG_TRACE && ncclDebugLevel != NCCL_LOG_TRACE) return;
   char buffer[1024];
-  size_t len;
-  len = snprintf(buffer, sizeof(buffer),
-                   "[%d:%d] ", node_model->nodeId, node_model->currRank);
+  size_t len = 0;
+  if (node_model) len = snprintf(buffer, sizeof(buffer),
+    "[%d:%d] ", node_model->nodeId, node_model->currRank);
   va_list args;
   va_start(args, fmt);
   vsprintf(buffer+len, fmt, args);
@@ -102,6 +109,16 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long flags, const char *file
   }
 }
 
+ncclResult_t ncclTopoGetSystem(const char* xmlTopoFile, struct ncclTopoSystem** system) {
+  struct ncclXml* xml;
+  NCCLCHECK(ncclCalloc(&xml, 1));
+  NCCLCHECK(ncclTopoGetXmlFromFile(xmlTopoFile, xml));
+  NCCLCHECK(ncclTopoGetSystemFromXml(xml, system));
+  free(xml);
+  return ncclSuccess;
+}
+
+
 ncclResult_t bootstrapAllGather(struct ncclComm* comm, struct allGather1Data_t * allGather1Data) {
   // AllGather1 - begin
   allGather1Data[comm->rank].peerInfo.rank = comm->rank;
@@ -110,12 +127,12 @@ ncclResult_t bootstrapAllGather(struct ncclComm* comm, struct allGather1Data_t *
   allGather1Data[comm->rank].peerInfo.hostHash = node_model->hostHash;
   allGather1Data[comm->rank].peerInfo.pidHash = node_model->pidHash;
   allGather1Data[comm->rank].peerInfo.shmDev = 0x19;
-  allGather1Data[comm->rank].peerInfo.busId = node_model->getGpuBusId(node_model->rankToCudaDev(comm->rank));
+  allGather1Data[comm->rank].peerInfo.busId = node_model->getGpuBusId(comm->rank);
   return ncclSuccess;
 }
 
-ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data,
-  struct allGather3Data_t *allGather3Data, struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph) {
+ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data, struct allGather3Data_t *allGather3Data,
+  struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph) {
   // We use 3 AllGathers
   // 1. { peerInfo, comm }
   // 2. ConnectTransport[nranks], ConnectValue[nranks]
@@ -147,45 +164,70 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t
   // AllGather1 - end
 
   // Topo detection / System graph creation
-  NCCLCHECK(ncclTopoGetSystem(comm, &comm->topo));
+  //NCCLCHECK(ncclTopoGetSystem(comm, &comm->topo));
   // Compute paths between GPUs and NICs
   NCCLCHECK(ncclTopoComputePaths(comm->topo, comm->peerInfo));
   // Remove inaccessible GPUs and unused NICs
   NCCLCHECK(ncclTopoTrimSystem(comm->topo, comm));
   // Recompute paths after trimming
   NCCLCHECK(ncclTopoComputePaths(comm->topo, comm->peerInfo));
-  // Compute max speed to accelerate search
-  NCCLCHECK(ncclTopoGetMaxSpeed(comm->topo));
+  // Init search
+  NCCLCHECK(ncclTopoSearchInit(comm->topo));
   // Print final topology
   NCCLCHECK(ncclTopoPrint(comm->topo));
 
   // Get rings and trees
-  //struct ncclTopoGraph treeGraph;
-  treeGraph.pattern = NCCL_TOPO_PATTERN_SPLIT_TREE;
-  treeGraph.crossNic = ncclParamCrossNic();
-  // We communicate only half the data between node with trees on 2 nodes.
-  NCCLCHECK(ncclTopoCompute(comm->topo, &treeGraph));
-  NCCLCHECK(ncclTopoPrintGraph(comm->topo, &treeGraph));
   //struct ncclTopoGraph ringGraph;
+  ringGraph.id = 0;
   ringGraph.pattern = NCCL_TOPO_PATTERN_RING;
   ringGraph.crossNic = ncclParamCrossNic();
+  ringGraph.collNet = 0;
+  ringGraph.minChannels = 1;
+  ringGraph.maxChannels = MAXCHANNELS/2;
   NCCLCHECK(ncclTopoCompute(comm->topo, &ringGraph));
   NCCLCHECK(ncclTopoPrintGraph(comm->topo, &ringGraph));
 
+  //struct ncclTopoGraph treeGraph;
+  treeGraph.id = 1;
+  treeGraph.pattern = NCCL_TOPO_PATTERN_SPLIT_TREE;
+  treeGraph.crossNic = ncclParamCrossNic();
+  treeGraph.collNet = 0;
+  treeGraph.minChannels = comm->topo->nodes[NET].count != 0 ? 1 : ringGraph.nChannels;
+  treeGraph.maxChannels = ringGraph.nChannels;
+  NCCLCHECK(ncclTopoCompute(comm->topo, &treeGraph));
+  NCCLCHECK(ncclTopoPrintGraph(comm->topo, &treeGraph));
+
+  //struct ncclTopoGraph collNetGraph;
+  collNetGraph.id = 2;
+  collNetGraph.pattern = NCCL_TOPO_PATTERN_TREE;
+  collNetGraph.collNet = 1;
+  collNetGraph.crossNic = ncclParamCrossNic();
+  collNetGraph.minChannels = collNetGraph.maxChannels = ringGraph.nChannels;
+  NCCLCHECK(ncclTopoCompute(comm->topo, &collNetGraph));
+  NCCLCHECK(ncclTopoPrintGraph(comm->topo, &collNetGraph));
+
+  if (comm->rank == ncclParamGraphDumpFileRank()) {
+    struct ncclTopoGraph* graphs[3] = { &ringGraph, &treeGraph, &collNetGraph };
+    NCCLCHECK(ncclTopoDumpGraphs(comm->topo, 3, graphs));
+  }
+
   // AllGather3 - begin
   allGather3Data[rank].cudaCompCap = ncclCudaCompCap();
-  allGather3Data[rank].nvlink = treeGraph.nvlink;
   allGather3Data[rank].nChannels = comm->nChannels = std::min(treeGraph.nChannels, ringGraph.nChannels);
   allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels;
   allGather3Data[rank].tree.speedIntra = treeGraph.speedIntra;
   allGather3Data[rank].tree.speedInter = treeGraph.speedInter;
-  allGather3Data[rank].tree.nvlink = treeGraph.nvlink;
+  allGather3Data[rank].tree.typeIntra = treeGraph.typeIntra;
   allGather3Data[rank].ring.sameChannels = ringGraph.sameChannels;
   allGather3Data[rank].ring.speedIntra = ringGraph.speedIntra;
   allGather3Data[rank].ring.speedInter = ringGraph.speedInter;
-  allGather3Data[rank].ring.nvlink = ringGraph.nvlink;
+  allGather3Data[rank].ring.typeIntra = ringGraph.typeIntra;
+  allGather3Data[rank].collNet.sameChannels = collNetGraph.sameChannels;
+  allGather3Data[rank].collNet.speedIntra = collNetGraph.speedIntra;
+  allGather3Data[rank].collNet.speedInter = collNetGraph.speedInter;
+  allGather3Data[rank].collNet.typeIntra = collNetGraph.typeIntra;
 
-  NCCLCHECK(ncclTopoPreset(comm, &treeGraph, &ringGraph, &allGather3Data[rank].topoRanks));
+  NCCLCHECK(ncclTopoPreset(comm, &treeGraph, &ringGraph, &collNetGraph, &allGather3Data[rank].topoRanks));
   //INFO(NCCL_GRAPH, "%d: nvlink %d nChannels %d tree.sameChannels %d tree.speedIntra %d tree.speedInter %d tree.nvlink %d ring.sameChannels %d ring.speedIntra %d ring.speedInter %d ring.nvlink %d",
   //  rank, allGather3Data[rank].nvlink, allGather3Data[rank].nChannels, allGather3Data[rank].tree.sameChannels, allGather3Data[rank].tree.speedIntra, allGather3Data[rank].tree.speedInter, allGather3Data[rank].tree.nvlink,
   //  allGather3Data[rank].ring.sameChannels, allGather3Data[rank].ring.speedIntra, allGather3Data[rank].ring.speedInter, allGather3Data[rank].ring.nvlink);
@@ -203,14 +245,8 @@ static ncclResult_t selectTransport(struct ncclTopoSystem* topo, struct ncclTopo
     int ret = 0;
     NCCLCHECK(transport->canConnect(&ret, topo, graph, myInfo, peerInfo));
     if (ret) {
-      //cpu_set_t affinitySave;
-      //sched_getaffinity(0, sizeof(cpu_set_t), &affinitySave);
-      //int cudaDev;
-      //CUDACHECK(hipGetDevice(&cudaDev));
-      //setCpuAffinity(cudaDev);
       connector->transportComm = transportComm;
       NCCLCHECK(transportComm->setup(topo, graph, myInfo, peerInfo, connect, connector, buffSize, channelId));
-      //sched_setaffinity(0, sizeof(cpu_set_t), &affinitySave);
       return ncclSuccess;
     }
   }
@@ -265,21 +301,26 @@ static ncclResult_t p2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph,
   return ncclSuccess;
 }
 
+
 ncclResult_t initChannel(struct ncclComm* comm, int channelid) {
   struct ncclChannel* channel = comm->channels+channelid;
   channel->id = channelid;
 
   // Setup intermediate buffering
-  //channel->buffSize = ncclParamBuffsize();
+  //int buffSize = ncclParamBuffsize();
+  int cpuArch, cpuVendor, cpuModel;
+  NCCLCHECK(ncclTopoCpuType(comm->topo, &cpuArch, &cpuVendor, &cpuModel));
+  //channel->buffSize = buffSize != -2 ? buffSize :
+  //  cpuArch == NCCL_TOPO_CPU_ARCH_ARM ? DEFAULT_BUFFER_SIZE_BYTES_ARM : DEFAULT_BUFFER_SIZE_BYTES;
 
   // Ring index to user rank table.
   //NCCLCHECK(ncclCudaCalloc(&channel->ring.devUserRanks, comm->nRanks));
   NCCLCHECK(ncclCalloc(&channel->ring.userRanks, comm->nRanks));
 
   // Communication structures with peers.
-  //NCCLCHECK(ncclCudaCalloc(&channel->devPeers, comm->nRanks));
-  NCCLCHECK(ncclCalloc(&channel->peers, comm->nRanks));
-  for (size_t i=0; i<comm->nRanks; ++i) {
+  //NCCLCHECK(ncclCudaCalloc(&channel->devPeers, comm->nRanks+1)); // The extra one rank is for collnet root (i.e. network)
+  NCCLCHECK(ncclCalloc(&channel->peers, comm->nRanks+1));
+  for (size_t i=0; i<comm->nRanks+1; ++i) {
     channel->peers[i].send.comm = comm;
     channel->peers[i].recv.comm = comm;
   }
@@ -307,8 +348,8 @@ static ncclResult_t setupChannel(struct ncclComm* comm, int channelId, int rank,
   return ncclSuccess;
 }
 
-ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t *allGather3Data, struct ncclTopoGraph& treeGraph,
-  struct ncclTopoGraph& ringGraph) {
+ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t *allGather3Data,
+  struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph) {
   int rank = comm->rank;
   int nranks = comm->nRanks;
   //NCCLCHECK(bootstrapAllGather(comm->bootstrap, allGather3Data, sizeof(*allGather3Data)));
@@ -329,6 +370,15 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
     if (i == comm->rank) comm->node = node;
   }
 
+  char line[1024];
+  sprintf(line, "nodesFirstRank: ");
+  int offset = strlen(line);
+  for (int i=0; i<comm->nNodes; i++) {
+    sprintf(line+offset, "%d ", nodesFirstRank[i]);
+    offset = strlen(line);
+  }
+  INFO(NCCL_INIT, "%s", line);
+
   // Determine the minimum CUDA Compute capability of all GPUs
   int myCompCap = allGather3Data[rank].cudaCompCap;
   int minCompCap = myCompCap, maxCompCap = myCompCap;
@@ -337,9 +387,6 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
     maxCompCap = std::max(allGather3Data[i].cudaCompCap, maxCompCap);
   }
 
-  comm->nvlink = 1;
-  for (int i = 0; i < nranks; i++) comm->nvlink &= allGather3Data[i].nvlink;
-
   int nChannelsOrig = comm->nChannels;
   struct ncclTopoRanks** allTopoRanks;
   NCCLCHECK(ncclCalloc(&allTopoRanks, comm->nRanks));
@@ -350,11 +397,15 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
     treeGraph.sameChannels = std::min(allGather3Data[i].tree.sameChannels, treeGraph.sameChannels);
     treeGraph.speedIntra = std::min(allGather3Data[i].tree.speedIntra, treeGraph.speedIntra);
     treeGraph.speedInter = std::min(allGather3Data[i].tree.speedInter, treeGraph.speedInter);
-    treeGraph.nvlink = std::min(allGather3Data[i].tree.nvlink, treeGraph.nvlink);
+    treeGraph.typeIntra = std::min(allGather3Data[i].tree.typeIntra, treeGraph.typeIntra);
     ringGraph.sameChannels = std::min(allGather3Data[i].ring.sameChannels, ringGraph.sameChannels);
     ringGraph.speedIntra = std::min(allGather3Data[i].ring.speedIntra, ringGraph.speedIntra);
     ringGraph.speedInter = std::min(allGather3Data[i].ring.speedInter, ringGraph.speedInter);
-    ringGraph.nvlink = std::min(allGather3Data[i].ring.nvlink, ringGraph.nvlink);
+    ringGraph.typeIntra = std::min(allGather3Data[i].ring.typeIntra, ringGraph.typeIntra);
+    collNetGraph.sameChannels = std::min(allGather3Data[i].collNet.sameChannels, collNetGraph.sameChannels);
+    collNetGraph.speedIntra = std::min(allGather3Data[i].collNet.speedIntra, collNetGraph.speedIntra);
+    collNetGraph.speedInter = std::min(allGather3Data[i].collNet.speedInter, collNetGraph.speedInter);
+    collNetGraph.typeIntra = std::min(allGather3Data[i].collNet.typeIntra, collNetGraph.typeIntra);
   }
 
   if (comm->nChannels < nChannelsOrig) {
@@ -366,24 +417,23 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
   int *rings;
   NCCLCHECK(ncclCalloc(&rings, nranks*MAXCHANNELS));
 
-  char line[1024];
-  sprintf(line, "nodesFirstRank: ");
-  int offset = strlen(line);
-  for (int i=0; i<comm->nNodes; i++) {
-    sprintf(line+offset, "%d ", nodesFirstRank[i]);
-    offset = strlen(line);
-  }
-  INFO(NCCL_INIT, "%s", line);
-
   NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, allTopoRanks, rings));
+  if (comm->nNodes > 1 &&
+      ncclParamCollNetEnable() == 1 &&
+      collNetSupport()) {
+    NCCLCHECK(ncclTopoConnectCollNet(comm, &collNetGraph, rank));
+  }
 
   free(allTopoRanks);
   free(nodesFirstRank);
+  //free(allGather3Data);
 
   // AllGather3 - end
 
   TRACE(NCCL_INIT, "rank %d nranks %d - BUILT %d TREES/RINGS", rank, nranks, comm->nChannels);
 
+  NCCLCHECK(ncclTopoSetThresholds(comm, minCompCap, maxCompCap, &treeGraph, &ringGraph, &collNetGraph));
+
   line[0]='\0';
   for (int c=0; c<comm->nChannels; c++) {
     struct ncclTree* treeUp = &comm->channels[c].treeUp;
@@ -395,26 +445,56 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
   line[1023] = '\0';
   INFO(NCCL_INIT, "Trees%s", line);
 
-  free(rings);
-
-  // Done with AllGather1 data
-  //free(allGather1Data);
-
-  TRACE(NCCL_INIT, "rank %d nranks %d - DONE", rank, nranks);
+  // Set Affinity to a CPU local the our GPU, so that all memory we allocate
+  // on the host is local.
+  cpu_set_t affinitySave;
+  sched_getaffinity(0, sizeof(cpu_set_t), &affinitySave);
+  NCCLCHECK(ncclTopoSetAffinity(comm->topo, comm->rank));
+  ncclResult_t ret;
 
   // Connect with prev/next for each ring
   struct ncclConnect *connect;
-  NCCLCHECK(ncclCalloc(&connect, 2));
+  NCCLCHECKGOTO(ncclCalloc(&connect, 2), ret, affinity_restore);
   for (int c=0; c<comm->nChannels; c++) {
     struct ncclChannel* channel = comm->channels+c;
-    NCCLCHECK(setupChannel(comm, c, rank, nranks, rings+c*nranks));
+    NCCLCHECKGOTO(setupChannel(comm, c, rank, nranks, rings+c*nranks), ret, affinity_restore);
     if (comm->nRanks == 1) continue;
-    NCCLCHECK(p2pSetup(comm, &ringGraph, channel, 1, &channel->ring.prev, 1, &channel->ring.next));
-    NCCLCHECK(p2pSetup(comm, &treeGraph, channel, NCCL_MAX_TREE_ARITY, channel->treeUp.down, 1, &channel->treeUp.up));
-    NCCLCHECK(p2pSetup(comm, &treeGraph, channel, 1, &channel->treeDn.up, NCCL_MAX_TREE_ARITY, channel->treeDn.down));
+    NCCLCHECKGOTO(p2pSetup(comm, &ringGraph, channel, 1, &channel->ring.prev, 1, &channel->ring.next), ret, affinity_restore);
+    NCCLCHECKGOTO(p2pSetup(comm, &treeGraph, channel, NCCL_MAX_TREE_ARITY, channel->treeUp.down, 1, &channel->treeUp.up), ret, affinity_restore);
+    NCCLCHECKGOTO(p2pSetup(comm, &treeGraph, channel, 1, &channel->treeDn.up, NCCL_MAX_TREE_ARITY, channel->treeDn.down), ret, affinity_restore);
   }
+
+  // Check if we can setup CollNet
+#if 0
+  if (comm->nNodes > 1 &&
+      ncclParamCollNetEnable() == 1 &&
+      collNetSupport()) {
+    int logicChannels = comm->nChannels/2;
+    int collNetSetupFail = 0;
+    const int recvIndex = 0;  // recv GPU index is always 0
+    const int sendIndex = collNetGraph.pattern == NCCL_TOPO_PATTERN_TREE ? 0 : 1;  // send GPU index depends on topo pattern
+    for (int c=0; c<logicChannels; c++) {
+      struct ncclChannel* channelRecv = comm->channels+logicChannels+c;
+      struct ncclChannel* channelSend = comm->channels+c;
+      NCCLCHECK(p2pSetup(comm, &collNetGraph, channelRecv, 1, &channelRecv->collTreeDn.up, 1, channelRecv->collTreeDn.down));
+      NCCLCHECK(p2pSetup(comm, &collNetGraph, channelSend, 1, channelSend->collTreeUp.down, 1, &channelSend->collTreeUp.up));
+      const int recvMaster = collNetGraph.intra[c*comm->localRanks+recvIndex];
+      const int sendMaster = collNetGraph.intra[c*comm->localRanks+sendIndex];
+      if (collNetSetup(comm, &collNetGraph, channelRecv, logicChannels, rank, nranks, recvMaster, sendMaster, comm->nNodes, 1) != 1)
+        collNetSetupFail = 1;
+      if (collNetSetup(comm, &collNetGraph, channelSend, logicChannels, rank, nranks, sendMaster, recvMaster, comm->nNodes, 0) != 1)
+        collNetSetupFail = 1;
+    }
+    // Verify CollNet setup across ranks
+    NCCLCHECK(checkCollNetSetup(comm, rank, collNetSetupFail));
+  }
+#endif
   TRACE(NCCL_INIT, "rank %d nranks %d - CONNECTED %d RINGS AND TREES", rank, nranks, comm->nChannels);
   free(connect);
+  free(rings);
+
+affinity_restore:
+  if (ret != ncclSuccess) return ret;
 
   return ncclSuccess;
 }