Refine and add new Rome models (#548)

src/graph/paths.cc

@@ -699,7 +699,7 @@ static ncclResult_t ncclTopoGetNchannels(struct ncclTopoSystem* system, int g /*
     // Local rank
     path = system->nodes[GPU].nodes[peer].paths[GPU]+g;
     if (path->type == PATH_NVL) {
-      float nvlWidth = ncclTopoNVLinkSpeed(system->nodes[GPU].nodes[g].gpu.cudaCompCap);
+      float nvlWidth = ncclTopoXGMISpeed(system->nodes[GPU].nodes[g].gpu.cudaCompCap);
       *nChannels = 2*std::max(1, (int)(path->width / nvlWidth));
     } else {
       *nChannels = 2;

src/graph/rome_models.cc

@@ -903,7 +903,7 @@ static ncclResult_t parseRomeSystem(struct ncclTopoSystem* system, struct rcclRo
       }
       if (!link->remNode) continue;
       if (link->type != LINK_NVL) continue;
-      romeTopo->connMatrix[i*romeTopo->nGpus+n] = link->width/VEGA_XGMI_WIDTH;
+      romeTopo->connMatrix[i*romeTopo->nGpus+n] = link->width/ncclTopoXGMISpeed(node->gpu.gcn);
       count ++;
     }
     if (romeTopo->nLinks < count) romeTopo->nLinks = count;

src/graph/topo.cc

@@ -543,7 +543,7 @@ ncclResult_t ncclTopoAddXGMI(struct ncclXmlNode* node, struct ncclTopoSystem* sy
       }
     }
     if (remote) {
-      int nvlSpeed = VEGA_XGMI_WIDTH;
+      float nvlSpeed = ncclTopoXGMISpeed(gpu->gpu.gcn);
       NCCLCHECK(ncclTopoConnectNodes(gpu, remote, LINK_NVL, count*nvlSpeed));
       if (remote->type != GPU) {
         NCCLCHECK(ncclTopoConnectNodes(remote, gpu, LINK_NVL, count*nvlSpeed));

src/graph/topo.h

@@ -25,6 +25,7 @@
 #define ARM_WIDTH 6.0
 #define NET_WIDTH 12.0           // 100Gbit
 #define VEGA_XGMI_WIDTH 24.0
+#define MI200_XGMI_WIDTH 36.0
 
 // Intel CPU convert GPU P2P traffic into 64B PCI TLPs, so GPU
 // to GPU traffic consumes more PCI bandwidth.
@@ -183,13 +184,8 @@ static ncclResult_t ncclTopoRankToIndex(struct ncclTopoSystem* system, int rank,
   return ncclInternalError;
 }
 
-// Returns NVLink speed in GB/s
-static float ncclTopoNVLinkSpeed(int cudaCompCap) {
-  return
-    cudaCompCap == 86 ? SM86_NVLINK_WIDTH :
-    cudaCompCap >= 80 ? SM80_NVLINK_WIDTH :
-    cudaCompCap >= 70 ? SM70_NVLINK_WIDTH :
-    cudaCompCap >= 60 ? SM60_NVLINK_WIDTH :
-    SM80_NVLINK_WIDTH;
+// Returns XGMI speed in GB/s
+static float ncclTopoXGMISpeed(int gcn) {
+  return gcn == 910 ? MI200_XGMI_WIDTH : VEGA_XGMI_WIDTH;
 }
 #endif