AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

topo_expl: fix build after switching to rocm-smi-lib (#405)

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

* Use minimal of 4 channels for gfx908
This commit is contained in:
Wenkai Du
2021-07-27 08:30:08 -07:00
committed by GitHub
parent 56801656f3
commit 135d47d125
4 changed files with 20 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/init.cc
+2
View File
@@ -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)
tools/topo_expl/Makefile
+1 -1
View File
@@ -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
tools/topo_expl/topo_expl.cpp
+2 -2
View File
@@ -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));
tools/topo_expl/utils.cpp
+15
View File
@@ -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;
}