From dcf623f2ec54211596b420851713f6b513c193ca Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Wed, 22 Nov 2023 15:07:36 -0800
Subject: [PATCH] Add special handling of gfx940 (#976)

* Add special handling of gfx940

* Update ring base

[ROCm/rccl commit: 50b2dd9fd7fe3981a4303b9a44390b47b3b1d812]
---
 projects/rccl/src/graph/paths.cc            | 52 ++++++++++++++++++++-
 projects/rccl/src/graph/rome_models.cc      | 36 ++------------
 projects/rccl/tools/topo_expl/topo_expl.cpp |  2 +
 3 files changed, 56 insertions(+), 34 deletions(-)

diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc
index 7f373b2f4f..acff580177 100644
--- a/projects/rccl/src/graph/paths.cc
+++ b/projects/rccl/src/graph/paths.cc
@@ -11,6 +11,7 @@
 #include "comm.h"
 #include "net.h"
 #include "channel.h"
+#include "xml.h"
 
 // Pre-compute GPU->NIC, GPU->GPU and NIC->GPU paths
 
@@ -542,6 +543,30 @@ ncclResult_t ncclTopoGetPxnRanks(struct ncclComm* comm, int** intermediateRanks,
   return ncclSuccess;
 }
 
+static bool rcclPathOverride(struct ncclTopoSystem* system, uint64_t distance) {
+  int i, j;
+
+  for (i = 0; i < system->nodes[GPU].count; i++) {
+    for (j = 0; j < system->nodes[NET].count; j++) {
+      if (system->nodes[NET].nodes[j].net.busId - system->nodes[GPU].nodes[i].id == distance)
+        break;
+    }
+    if (j >= system->nodes[NET].count)
+      break;
+  }
+  if (i >= system->nodes[GPU].count) {
+    for (i = 0; i < system->nodes[GPU].count; i++) {
+      for (j = 0; j < system->nodes[NET].count; j++) {
+        if (system->nodes[NET].nodes[j].net.busId - system->nodes[GPU].nodes[i].id == distance)
+          system->nodes[GPU].nodes[i].paths[NET][j].type = PATH_PXB;
+      }
+    }
+    return true;
+  } else {
+    return false;
+  }
+}
+
 ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm* comm) {
   // Precompute paths between GPUs/NICs.
 
@@ -600,6 +625,26 @@ ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm
     }
   }
 
+  // Special handling of gfx94x
+
+#if !defined(TOPO_EXPL)
+  char strValue[1024];
+  NCCLCHECK(ncclTopoGetStrFromSys("/sys/devices/virtual/dmi/id", "bios_version", strValue));
+  if (strncmp("Hyper-V UEFI Release", strValue, 20) == 0) {
+#endif
+    int arch, vendor, model;
+    NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model));
+    if (arch == NCCL_TOPO_CPU_ARCH_X86 && vendor == NCCL_TOPO_CPU_VENDOR_INTEL &&
+      IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx94") &&
+      ((system->nodes[GPU].count == 8 && system->nodes[NET].count == 8 && system->nodes[GPU].count == system->nRanks) ||
+      (system->nodes[GPU].count != system->nRanks))) {
+      if (!rcclPathOverride(system, 0x100000))
+        rcclPathOverride(system, 0x1000);
+    }
+#if !defined(TOPO_EXPL)
+  }
+#endif
+
   // Update paths for NICs (no GPU Direct, PXN, ...)
   for (int n=0; n<system->nodes[NET].count; n++) {
     struct ncclTopoNode* netNode = system->nodes[NET].nodes+n;
@@ -639,7 +684,7 @@ ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm
   return ncclSuccess;
 }
 
-RCCL_PARAM(EnableIntranet, "ENABLE_INTRANET", 0);
+RCCL_PARAM(EnableIntranet, "ENABLE_INTRANET", -2);
 
 ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm* comm) {
   int *domains;
@@ -726,7 +771,10 @@ ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm*
     INFO(NCCL_GRAPH, "GDR is available on all GPUs");
   }
 
-  if (rcclParamEnableIntranet()) {
+  // Special handling of gfx94x
+  if (rcclParamEnableIntranet() == 1 || (rcclParamEnableIntranet() == -2 &&
+    IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx94") &&
+    system->nodes[GPU].count == 8 && system->nodes[NET].count == 8)) {
     remove = 0;
     system->type |= RCCL_TOPO_FORCE_INTRA;
   }
diff --git a/projects/rccl/src/graph/rome_models.cc b/projects/rccl/src/graph/rome_models.cc
index e5015e08a8..6f47e8fac1 100644
--- a/projects/rccl/src/graph/rome_models.cc
+++ b/projects/rccl/src/graph/rome_models.cc
@@ -548,7 +548,7 @@ static struct rcclRomeModel rome_model_79 = {
   .nicIds = { },
   .gpuNuma = { 0, 0, 0, 0, 1, 1, 1, 1, },
   .nicNuma = { },
-  .connMatrix = { 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, },
+  .connMatrix = { 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, },
   .gdrLevel = { },
   .pattern = "4040",
   .ringBase = "0 1 2 3 4 5 6 7|0 1 2 3 4 5 7 6|0 2 4 1 3 6 5 7|0 2 4 6 1 7 3 5|0 3 1 5 2 7 4 6|0 3 5 1 6 2 7 4|0 4 1 7 3 6 2 5|7 6 5 4 3 2 1 0|6 7 5 4 3 2 1 0|7 5 6 3 1 4 2 0|5 3 7 1 6 4 2 0|6 4 7 2 5 1 3 0|4 7 2 6 1 5 3 0|5 2 6 3 7 1 4 0",
@@ -574,37 +574,11 @@ static struct rcclRomeModel rome_model_81 = {
   .nicIds = { 0x7000, 0x1d000, 0x33000, 0x57000, 0x9a000, 0xaa000, 0xba000, 0xda000, },
   .gpuNuma = { 0, 0, 0, 0, 1, 1, 1, 1, },
   .nicNuma = { 0, 0, 0, 0, 1, 1, 1, 1, },
-  .connMatrix = { 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, 2, 2, 2, 2, 2, 2, 2, 2, 0, },
+  .connMatrix = { 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, },
   .gdrLevel = { PATH_PXB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PXB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PXB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PXB, },
   .pattern = "4444",
-  .ringBase = "N0 0 1 2 3 4 5 6 7 N7|N1 1 0 2 4 3 5 7 6 N6|N2 2 5 0 3 6 1 7 4 N4|N3 3 7 0 4 2 1 6 5 N5|N4 4 6 2 7 3 0 5 1 N1|N5 5 4 7 1 3 2 6 0 N0|N6 6 3 1 4 0 7 5 2 N2|N7 7 2 0 6 4 1 5 3 N3",
-  .options = "noCpuCheck=1",
-};
-
-static struct rcclRomeModel rome_model_82 = {
-  .nGpus = 8, .nCpus = 2, .nNics = 0, .nLinks = 7,
-  .gpuIds = { 0xc000, 0x22000, 0x38000, 0x5c000, 0x9f000, 0xaf000, 0xbf000, 0xdf000, },
-  .nicIds = { },
-  .gpuNuma = { 0, 0, 0, 0, 1, 1, 1, 1, },
-  .nicNuma = { },
-  .connMatrix = { 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, },
-  .gdrLevel = { },
-  .pattern = "4040",
-  .ringBase = "0 1 2 3 4 5 6 7|0 1 2 3 4 5 7 6|0 2 4 1 3 6 5 7|0 2 4 6 1 7 3 5|0 3 1 5 2 7 4 6|0 3 5 1 6 2 7 4|0 4 1 7 3 6 2 5|7 6 5 4 3 2 1 0|6 7 5 4 3 2 1 0|7 5 6 3 1 4 2 0|5 3 7 1 6 4 2 0|6 4 7 2 5 1 3 0|4 7 2 6 1 5 3 0|5 2 6 3 7 1 4 0",
-  .options = "noCpuCheck=1,mscclEnabled=1,disableNumaMatching=1",
-};
-
-static struct rcclRomeModel rome_model_83 = {
-  .nGpus = 8, .nCpus = 2, .nNics = 8, .nLinks = 7,
-  .gpuIds = { 0xc000, 0x22000, 0x38000, 0x5c000, 0x9f000, 0xaf000, 0xbf000, 0xdf000, },
-  .nicIds = { 0x7000, 0x1d000, 0x33000, 0x57000, 0x9a000, 0xaa000, 0xba000, 0xda000, },
-  .gpuNuma = { 0, 0, 0, 0, 1, 1, 1, 1, },
-  .nicNuma = { 0, 0, 0, 0, 1, 1, 1, 1, },
-  .connMatrix = { 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, },
-  .gdrLevel = { PATH_PHB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PHB, PATH_PHB, },
-  .pattern = "4444",
-  .ringBase = "N0 0 1 2 3 4 5 6 7 N7|N1 1 0 2 4 3 5 7 6 N6|N2 2 5 0 3 6 1 7 4 N4|N3 3 7 0 4 2 1 6 5 N5|N4 4 6 2 7 3 0 5 1 N1|N5 5 4 7 1 3 2 6 0 N0|N6 6 3 1 4 0 7 5 2 N2|N7 7 2 0 6 4 1 5 3 N3",
-  .options = "noCpuCheck=1,disableNumaMatching=1",
+  .ringBase = "N0 0 1 2 3 4 5 6 7 N7|N1 1 0 2 4 3 5 7 6 N6|N2 2 5 0 3 7 1 6 4 N4|N3 3 6 1 5 2 7 4 0 N0|N4 4 7 0 6 5 1 3 2 N2|N5 5 4 6 3 0 7 2 1 N1|N6 6 2 0 4 1 7 5 3 N3|N7 7 3 1 4 2 6 0 5 N5|N0 0 1 2 3 4 5 6 7 N7|N1 1 0 2 4 3 5 7 6 N6|N2 2 5 0 3 7 1 6 4 N4|N3 3 6 1 5 2 7 4 0 N0|N4 4 7 0 6 5 1 3 2 N2|N5 5 4 6 3 0 7 2 1 N1|N6 6 2 0 4 1 7 5 3 N3|N7 7 3 1 4 2 6 0 5 N5",
+  .options = "noCpuCheck=1,mscclEnabled=1",
 };
 
 static struct rcclRomeModel romeTopoModels[] = {
@@ -649,8 +623,6 @@ static struct rcclRomeModel romeTopoModels[] = {
   rome_model_79,
   rome_model_80,
   rome_model_81,
-  rome_model_82,
-  rome_model_83,
 };
 
 /* Parse user defined rings. Format is like :
diff --git a/projects/rccl/tools/topo_expl/topo_expl.cpp b/projects/rccl/tools/topo_expl/topo_expl.cpp
index b675200f15..cca1f4fdd2 100644
--- a/projects/rccl/tools/topo_expl/topo_expl.cpp
+++ b/projects/rccl/tools/topo_expl/topo_expl.cpp
@@ -154,6 +154,8 @@ NodeModelDesc model_descs[] = {
   {1, "topo_8p_940.xml",        "Single node gfx940 8P"},
   {2, "topo_4p_940.xml",        "2 nodes gfx940 4P"},
   {4, "topo_8p_940.xml",        "4 nodes gfx940 8P"},
+  {1, "topo_8p_940vm.xml",      "single node gfx940 VM"},
+  {2, "topo_8p_940vm.xml",      "2 nodes gfx940 VM"},
 };
 
 NCCL_PARAM(MaxCTAs, "MAX_CTAS", MAXCHANNELS);