From 135d47d125cf16064fc63e8ee946174e7ec8afdd Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Tue, 27 Jul 2021 08:30:08 -0700
Subject: [PATCH] topo_expl: fix build after switching to rocm-smi-lib (#405)

* topo_expl: fix build after switching to rocm-smi-lib

* Use minimal of 4 channels for gfx908
---
 src/init.cc                   |  2 ++
 tools/topo_expl/Makefile      |  2 +-
 tools/topo_expl/topo_expl.cpp |  4 ++--
 tools/topo_expl/utils.cpp     | 15 +++++++++++++++
 4 files changed, 20 insertions(+), 3 deletions(-)

diff --git a/src/init.cc b/src/init.cc
index 6da023d14d..5579cfa361 100644
--- a/src/init.cc
+++ b/src/init.cc
@@ -865,6 +865,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
   int idx;
   NCCLCHECK(ncclTopoIdToIndex(comm->topo, GPU, myInfo->busId, &idx));
   allGather3Data[rank].nc = 2;
+  if (comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 908)
+    allGather3Data[rank].nc = std::max(4/ringGraph.nChannels, 2);
   if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G))
     allGather3Data[rank].nc = 4;
   if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910)
diff --git a/tools/topo_expl/Makefile b/tools/topo_expl/Makefile
index 7480cae94b..1edb7f0edf 100644
--- a/tools/topo_expl/Makefile
+++ b/tools/topo_expl/Makefile
@@ -6,7 +6,7 @@ endif
 HIPCC = $(HIP_PATH)/bin/hipcc
 
 EXE = topo_expl
-CXXFLAGS = -g -O3 -Iinclude -I../../src -I../../src/include -I../../src/graph/ -DTOPO_EXPL -DENABLE_TRACE -lnuma
+CXXFLAGS = -g -O3 -Iinclude -I../../src -I../../src/include -I../../src/graph/ -I/opt/rocm/rocm_smi/include/ -DTOPO_EXPL -DENABLE_TRACE -lnuma
 
 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 ../../src/graph/rome_models.cc
diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp
index 52796ee21f..622949b2df 100644
--- a/tools/topo_expl/topo_expl.cpp
+++ b/tools/topo_expl/topo_expl.cpp
@@ -184,8 +184,8 @@ int main(int argc,char* argv[])
   for (int i = 0; i < nranks; i++) {
     comm[i].rank = i;
     comm[i].nRanks = nranks;
-    NCCLCHECK(ncclCalloc(&comm[i].connectSend, comm->nRanks));
-    NCCLCHECK(ncclCalloc(&comm[i].connectRecv, comm->nRanks));
+    NCCLCHECK(ncclCalloc(&comm[i].connectSend, NCCL_MAX_CONNS*comm->nRanks));
+    NCCLCHECK(ncclCalloc(&comm[i].connectRecv, NCCL_MAX_CONNS*comm->nRanks));
     comm[i].p2pSendCount = comm[i].p2pRecvCount = 0;
     NCCLCHECK(ncclCalloc(&comm[i].p2pSends, comm->nRanks));
     NCCLCHECK(ncclCalloc(&comm[i].p2pRecvs, comm->nRanks));
diff --git a/tools/topo_expl/utils.cpp b/tools/topo_expl/utils.cpp
index 269d642092..8ecebbe690 100644
--- a/tools/topo_expl/utils.cpp
+++ b/tools/topo_expl/utils.cpp
@@ -29,6 +29,7 @@
 #include "coll_net.h"
 #include "model.h"
 #include "utils.h"
+#include "rocm_smi/rocm_smi.h"
 
 const char* ncclFuncStr[NCCL_NUM_FUNCTIONS+1] = { "Broadcast", "Reduce", "AllGather", "ReduceScatter", "AllReduce", "SendRecv" };
 const char* ncclAlgoStr[NCCL_NUM_ALGORITHMS] = { "Tree", "Ring", "CollNet" };
@@ -436,6 +437,8 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t
   int idx;
   NCCLCHECK(ncclTopoIdToIndex(comm->topo, GPU, myInfo->busId, &idx));
   allGather3Data[rank].nc = 2;
+  if (comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 908)
+    allGather3Data[rank].nc = std::max(4/ringGraph.nChannels, 2);
   if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G))
     allGather3Data[rank].nc = 4;
   if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910)
@@ -843,3 +846,15 @@ affinity_restore:
   TRACE(NCCL_INIT, "rank %d nranks %d - DONE", rank, nranks);
   return ncclSuccess;
 }
+
+ncclResult_t rocm_smi_init() {
+  return ncclSuccess;
+}
+
+ncclResult_t rocm_smi_getDeviceIndexByPciBusId(const char* pciBusId, uint32_t* deviceIndex) {
+  return ncclSuccess;
+}
+
+ncclResult_t rocm_smi_getLinkInfo(int srcDev, int dstDev, RSMI_IO_LINK_TYPE* rsmi_type, int *hops, int *bw) {
+  return ncclSuccess;
+}