Improve GPU direct RDMA handling on Rome (#294)

src/graph/paths.cc

@@ -414,14 +414,15 @@ ncclResult_t ncclTopoCheckGdr(struct ncclTopoSystem* system, int64_t busId, int
 
   // Check if we are close enough that it makes sense to enable GDR
   int netGdrLevel = PATH_PXB;
-#ifdef TOPO_EXPL
-  int arch, vendor, model;
-  NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model));
-  if (arch == NCCL_TOPO_CPU_ARCH_X86 && vendor == NCCL_TOPO_CPU_VENDOR_AMD && model == NCCL_TOPO_CPU_TYPE_ROME)
-    netGdrLevel = PATH_PHB;
-#endif
   NCCLCHECK(ncclGetLevel(&ncclTopoUserGdrLevel, NULL, "NCCL_NET_GDR_LEVEL"));
   if (ncclTopoUserGdrLevel != -2) netGdrLevel = ncclTopoUserGdrLevel;
+  else {
+    int arch, vendor, model;
+    NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model));
+    if((system->nodes[GPU].nodes[g].id & 0xf0000) == (system->nodes[NET].nodes[n].net.busId & 0xf0000))
+      netGdrLevel = PATH_PHB;
+  }
+
   int distance = gpu->paths[NET][n].type;
   if (distance > netGdrLevel) {
     INFO(NCCL_NET,"GPU Direct RDMA Disabled for GPU %lx / HCA %d (distance %d > %d)", busId, netDev, distance, netGdrLevel);

src/graph/rome_models.h

@@ -29,6 +29,7 @@ struct rcclRomeModel {
   int nNics;
   int nLinks;
   int64_t gpuIds[MAX_ROME_GPUS];
+  int64_t nicIds[MAX_ROME_NICS];
   int64_t gpuNuma[MAX_ROME_GPUS];
   int64_t nicNuma[MAX_ROME_NICS];
   int connMatrix[MAX_ROME_GPUS*MAX_ROME_GPUS];
@@ -39,6 +40,7 @@ struct rcclRomeModel {
 static struct rcclRomeModel rome_model_22 = {
   .nGpus = 8, .nCpus = 4, .nNics = 1, .nLinks = 2,
   .gpuIds = { 0x3000, 0x43000, 0x26000, 0xc3000, 0x83000, 0x23000, 0xc6000, 0xa3000, },
+  .nicIds = { 0xe1000, },
   .gpuNuma = { 1, 0, 1, 2, 3, 1, 2, 3, },
   .nicNuma = { 2, },
   .connMatrix = { 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 1, 1, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, },
@@ -49,6 +51,7 @@ static struct rcclRomeModel rome_model_22 = {
 static struct rcclRomeModel rome_model_25 = {
   .nGpus = 8, .nCpus = 4, .nNics = 2, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, },
+  .nicIds = { 0x61000, 0xa1000, },
   .gpuNuma = { 0, 1, 1, 1, 2, 2, 2, 3, },
   .nicNuma = { 0, 3, },
   .connMatrix = { 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, },
@@ -59,6 +62,7 @@ static struct rcclRomeModel rome_model_25 = {
 static struct rcclRomeModel rome_model_27 = {
   .nGpus = 8, .nCpus = 4, .nNics = 2, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, },
+  .nicIds = { 0x61000, 0xa1000, },
   .gpuNuma = { 0, 1, 1, 1, 2, 2, 2, 3, },
   .nicNuma = { 0, 3, },
   .connMatrix = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, },
@@ -69,6 +73,7 @@ static struct rcclRomeModel rome_model_27 = {
 static struct rcclRomeModel rome_model_29 = {
   .nGpus = 8, .nCpus = 4, .nNics = 1, .nLinks = 3,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, },
+  .nicIds = { 0xe1000, },
   .gpuNuma = { 0, 1, 1, 1, 2, 2, 3, 3, },
   .nicNuma = { 2, },
   .connMatrix = { 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, },
@@ -79,6 +84,7 @@ static struct rcclRomeModel rome_model_29 = {
 static struct rcclRomeModel rome_model_31 = {
   .nGpus = 8, .nCpus = 8, .nNics = 2, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, },
+  .nicIds = { 0x61000, 0xa1000, },
   .gpuNuma = { 1, 2, 2, 3, 4, 5, 5, 7, },
   .nicNuma = { 0, 6, },
   .connMatrix = { 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, },
@@ -89,6 +95,7 @@ static struct rcclRomeModel rome_model_31 = {
 static struct rcclRomeModel rome_model_33 = {
   .nGpus = 8, .nCpus = 8, .nNics = 2, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, },
+  .nicIds = { 0x61000, 0xa1000, },
   .gpuNuma = { 1, 2, 2, 3, 4, 5, 5, 7, },
   .nicNuma = { 0, 6, },
   .connMatrix = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, },
@@ -99,6 +106,7 @@ static struct rcclRomeModel rome_model_33 = {
 static struct rcclRomeModel rome_model_30 = {
   .nGpus = 8, .nCpus = 8, .nNics = 0, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, },
+  .nicIds = { },
   .gpuNuma = { 1, 2, 2, 3, 4, 5, 5, 7, },
   .nicNuma = { },
   .connMatrix = { 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, },
@@ -109,6 +117,7 @@ static struct rcclRomeModel rome_model_30 = {
 static struct rcclRomeModel rome_model_32 = {
   .nGpus = 8, .nCpus = 8, .nNics = 0, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, },
+  .nicIds = { },
   .gpuNuma = { 1, 2, 2, 3, 4, 5, 5, 7, },
   .nicNuma = { },
   .connMatrix = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, },
@@ -119,6 +128,7 @@ static struct rcclRomeModel rome_model_32 = {
 static struct rcclRomeModel rome_model_24 = {
   .nGpus = 8, .nCpus = 4, .nNics = 0, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, },
+  .nicIds = { },
   .gpuNuma = { 0, 1, 1, 1, 2, 2, 2, 3, },
   .nicNuma = { },
   .connMatrix = { 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 1, 0, },
@@ -129,6 +139,7 @@ static struct rcclRomeModel rome_model_24 = {
 static struct rcclRomeModel rome_model_26 = {
   .nGpus = 8, .nCpus = 4, .nNics = 0, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xe3000, 0xc3000, 0xc6000, 0x83000, },
+  .nicIds = { },
   .gpuNuma = { 0, 1, 1, 1, 2, 2, 2, 3, },
   .nicNuma = { },
   .connMatrix = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, },
@@ -139,6 +150,7 @@ static struct rcclRomeModel rome_model_26 = {
 static struct rcclRomeModel rome_model_23 = {
   .nGpus = 8, .nCpus = 4, .nNics = 0, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, },
+  .nicIds = { },
   .gpuNuma = { 0, 1, 1, 1, 2, 2, 3, 3, },
   .nicNuma = { },
   .connMatrix = { 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, },
@@ -149,6 +161,7 @@ static struct rcclRomeModel rome_model_23 = {
 static struct rcclRomeModel rome_model_38 = {
   .nGpus = 8, .nCpus = 7, .nNics = 0, .nLinks = 2,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, },
+  .nicIds = { },
   .gpuNuma = { 1, 2, 2, 3, 5, 5, 6, 7, },
   .nicNuma = { },
   .connMatrix = { 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, 0, 0, 1, 0, },
@@ -159,6 +172,7 @@ static struct rcclRomeModel rome_model_38 = {
 static struct rcclRomeModel rome_model_28 = {
   .nGpus = 8, .nCpus = 4, .nNics = 0, .nLinks = 3,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, },
+  .nicIds = { },
   .gpuNuma = { 0, 1, 1, 1, 2, 2, 3, 3, },
   .nicNuma = { },
   .connMatrix = { 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, },
@@ -169,6 +183,7 @@ static struct rcclRomeModel rome_model_28 = {
 static struct rcclRomeModel rome_model_40 = {
   .nGpus = 8, .nCpus = 4, .nNics = 1, .nLinks = 3,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, },
+  .nicIds = { 0xe1000, },
   .gpuNuma = { 0, 1, 1, 1, 2, 2, 3, 3, },
   .nicNuma = { 2, },
   .connMatrix = { 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 1, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, },
@@ -179,6 +194,7 @@ static struct rcclRomeModel rome_model_40 = {
 static struct rcclRomeModel rome_model_42 = {
   .nGpus = 8, .nCpus = 7, .nNics = 1, .nLinks = 3,
   .gpuIds = { 0x43000, 0x23000, 0x26000, 0x3000, 0xc3000, 0xc6000, 0xa3000, 0x83000, },
+  .nicIds = { 0xe1000, },
   .gpuNuma = { 1, 2, 2, 3, 5, 5, 6, 7, },
   .nicNuma = { 4, },
   .connMatrix = { 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 1, 1, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 1, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, },
@@ -189,11 +205,12 @@ static struct rcclRomeModel rome_model_42 = {
 static struct rcclRomeModel rome_model_44 = {
   .nGpus = 8, .nCpus = 4, .nNics = 1, .nLinks = 3,
   .gpuIds = { 0x63000, 0x43000, 0x27000, 0x3000, 0xe3000, 0xc3000, 0xa3000, 0x83000, },
+  .nicIds = { 0xc4000, },
   .gpuNuma = { 0, 0, 1, 1, 2, 2, 3, 3, },
   .nicNuma = { 2, },
   .connMatrix = { 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0, 0, 0, 0, 1, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 1, 0, 0, 0, 0, 1, 1, 1, 0, },
   .pattern = "20202120",
-  .ringBase = "4 5 7 6 2 1 3 0|7 4 5 6 2 0 1 3",
+  .ringBase = "5 4 7 6 2 1 3 0|5 6 7 4 1 0 2 3",
 };
 
 static struct rcclRomeModel romeTopoModels[] = {

src/graph/search.cc

@@ -905,6 +905,7 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo
     }
     if (j >= romeTopo->nNics) {
       net_map[j] = i;
+      romeTopo->nicIds[romeTopo->nNics] = system->nodes[NET].nodes[i].net.busId;
       (romeTopo->nNics)++;
       if (romeTopo->nNics >= MAX_ROME_NICS) break;
     }
@@ -941,6 +942,9 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo
     fprintf(file, "  .gpuIds = { ");
     for (int i = 0; i < romeTopo->nGpus; i ++) fprintf(file, "0x%lx, ", romeTopo->gpuIds[i]);
     fprintf(file, "},\n");
+    fprintf(file, "  .nicIds = { ");
+    for (int i = 0; i < romeTopo->nNics; i ++) fprintf(file, "0x%lx, ", romeTopo->nicIds[i]);
+    fprintf(file, "},\n");
     fprintf(file, "  .gpuNuma = { ");
     for (int i = 0; i < romeTopo->nGpus; i ++) fprintf(file, "%ld, ", romeTopo->gpuNuma[i]);
     fprintf(file, "},\n");
@@ -1038,13 +1042,18 @@ static ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopo
   }
 
   char line[1024];
+#ifdef ENABLE_TRACE
   sprintf(line, "Found matching Rome model index %d in %.2fms (%d iter) with GPU mapping: ", i, t, time);
+#else
+  sprintf(line, "Found matching Rome model index %d with GPU mapping: ", i);
+#endif
   int offset = strlen(line);
   for (int k = 0; k < ngpus; k++) {
     sprintf(line+offset, "%d ", g[k]);
     offset = strlen(line);
   }
   INFO(NCCL_GRAPH, "%s", line);
+
   // create 4P2H based on reference and remapped ids
   NCCLCHECK(parseGraph(romeTopoModels[i].ringBase, system, graph, g, romeTopo.nNics, net_map));
   return ncclSuccess;

src/graph/topo.cc

@@ -266,7 +266,7 @@ ncclResult_t ncclTopoSortSystem(struct ncclTopoSystem* system) {
   return ncclSuccess;
 }
 
-ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* system, struct ncclTopoNode* nic) {
+ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* system, struct ncclTopoNode* nic, int64_t busId) {
   int dev;
   NCCLCHECK(xmlGetAttrInt(xmlNet, "dev", &dev));
 
@@ -286,6 +286,7 @@ ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* s
   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 (ncclCollNet && xmlGetAttrInt(xmlNet, "coll", &net->net.collSupport) != ncclSuccess) net->net.collSupport = 0;
+  net->net.busId = busId;
   ncclDebugNoWarn = 0;
 
   NCCLCHECK(ncclTopoConnectNodes(nic, net, LINK_NET, net->net.width));
@@ -293,14 +294,14 @@ ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* s
   return ncclSuccess;
 }
 
-ncclResult_t ncclTopoAddNic(struct ncclXmlNode* xmlNic, struct ncclTopoSystem* system, struct ncclTopoNode* nic) {
+ncclResult_t ncclTopoAddNic(struct ncclXmlNode* xmlNic, struct ncclTopoSystem* system, struct ncclTopoNode* nic, int64_t busId) {
   for (int s=0; s<xmlNic->nSubs; s++) {
     struct ncclXmlNode* xmlNet = xmlNic->subs[s];
     if (strcmp(xmlNet->name, "net") != 0) continue;
     int index;
     NCCLCHECK(xmlGetAttrIndex(xmlNet, "dev", &index));
     if (index == -1) continue;
-    NCCLCHECK(ncclTopoAddNet(xmlNet, system, nic));
+    NCCLCHECK(ncclTopoAddNet(xmlNet, system, nic, busId));
   }
   return ncclSuccess;
 }
@@ -354,7 +355,7 @@ ncclResult_t ncclTopoAddPci(struct ncclXmlNode* xmlPci, struct ncclTopoSystem* s
       NCCLCHECK(ncclTopoCreateNode(system, &nicNode, type, busId));
       node = nicNode; // Connect it to parent later on
     }
-    NCCLCHECK(ncclTopoAddNic(xmlNic, system, nicNode));
+    NCCLCHECK(ncclTopoAddNic(xmlNic, system, nicNode, busId));
   } else if (type == PCI) {
     NCCLCHECK(ncclTopoCreateNode(system, &node, type, busId));
     for (int s=0; s<xmlPci->nSubs; s++) {
@@ -421,7 +422,7 @@ ncclResult_t ncclTopoAddCpu(struct ncclXmlNode* xmlCpu, struct ncclTopoSystem* s
         NCCLCHECK(ncclTopoConnectNodes(cpu, nic, LINK_PCI, LOC_WIDTH));
         NCCLCHECK(ncclTopoConnectNodes(nic, cpu, LINK_PCI, LOC_WIDTH));
       }
-      NCCLCHECK(ncclTopoAddNic(node, system, nic));
+      NCCLCHECK(ncclTopoAddNic(node, system, nic, 0));
     }
   }
   return ncclSuccess;

src/graph/topo.h

@@ -100,6 +100,7 @@ struct ncclTopoNode {
       int gdrSupport;
       int collSupport;
       int maxChannels;
+      int64_t busId;
     }net;
     struct {
       int arch;