diff --git a/src/include/enqueue.h b/src/include/enqueue.h index a715a44c43..03d44c51ff 100644 --- a/src/include/enqueue.h +++ b/src/include/enqueue.h @@ -33,7 +33,7 @@ static inline size_t ncclFuncSendCount(ncclFunc_t func, int nRanks, size_t count static inline size_t ncclFuncRecvCount(ncclFunc_t func, int nRanks, size_t count) { return func == ncclFuncAllGather ? nRanks*count : count; } -rccl_static_inline size_t ncclFuncMaxSendRecvCount(ncclFunc_t func, int nRanks, size_t count) { +rccl_static inline size_t ncclFuncMaxSendRecvCount(ncclFunc_t func, int nRanks, size_t count) { return func == ncclFuncAllGather || func == ncclFuncReduceScatter ? nRanks*count : count; } diff --git a/src/include/rccl_common.h b/src/include/rccl_common.h index 549df17ef3..9b245685fb 100644 --- a/src/include/rccl_common.h +++ b/src/include/rccl_common.h @@ -78,5 +78,4 @@ ncclResult_t rcclGetAlgoInfo(struct ncclComm* comm, ncclFunc_t coll, uint64_t co int* algo, int* protocol, int* maxChannels); ncclResult_t rcclFuncMaxSendRecvCount(ncclFunc_t func, int nRanks, size_t count, size_t& maxCount); - #endif \ No newline at end of file diff --git a/src/rccl_wrap.cc b/src/rccl_wrap.cc index 33f4f2afb1..7ed13a6aeb 100644 --- a/src/rccl_wrap.cc +++ b/src/rccl_wrap.cc @@ -23,7 +23,7 @@ THE SOFTWARE. #include "rccl_common.h" #include "comm.h" #include "graph/topo.h" - +#include "enqueue.h" void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, struct ncclTaskColl* info) { // Honor user input for protocol choice static int userProtocolInput = -2; @@ -72,7 +72,6 @@ void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, s } } -extern size_t ncclFuncMaxSendRecvCount(ncclFunc_t func, int nRanks, size_t count); extern ncclResult_t getAlgoInfo( struct ncclComm* comm, struct ncclTaskColl* task, int collNetSupport, int nvlsSupport, int numPipeOps, ncclSimInfo_t* simInfo = NULL diff --git a/tools/topo_expl/Makefile b/tools/topo_expl/Makefile index d4527e2a9d..dfab36f19c 100644 --- a/tools/topo_expl/Makefile +++ b/tools/topo_expl/Makefile @@ -10,7 +10,7 @@ CXXFLAGS = -g -ffunction-sections -fdata-sections -Wl,--gc-sections -fgpu-rdc -I files = $(EXE).cpp model.cpp utils.cpp hipify_rccl/graph/topo.cc hipify_rccl/graph/rings.cc hipify_rccl/graph/paths.cc hipify_rccl/graph/trees.cc ../../src/misc/param.cc \ hipify_rccl/graph/search.cc hipify_rccl/graph/connect.cc hipify_rccl/graph/tuning.cc hipify_rccl/graph/xml.cc ../../src/misc/nvmlwrap_stub.cc hipify_rccl/graph/rome_models.cc hipify_rccl/graph/archinfo.cc \ - hipify_rccl/collectives.cc hipify_rccl/enqueue.cc ../../src/rccl_wrap.cc + hipify_rccl/collectives.cc hipify_rccl/register.cc hipify_rccl/enqueue.cc ../../src/rccl_wrap.cc all: hipify $(EXE) @@ -25,6 +25,7 @@ hipify: cp -a ../../src/device/*.h hipify_rccl/device/include cp -a ../../src/device/network/unpack/*.h hipify_rccl/include/network/unpack cp -a ../../src/enqueue.cc hipify_rccl/ + cp -a ../../src/register/register.cc hipify_rccl/ cp -a ../../src/collectives.cc hipify_rccl/ cp -a ../../src/misc/archinfo.cc hipify_rccl/graph/ hipify-perl -inplace -quiet-warnings hipify_rccl/include/*.h diff --git a/tools/topo_expl/model.cpp b/tools/topo_expl/model.cpp index 9c4449f280..12f22be219 100644 --- a/tools/topo_expl/model.cpp +++ b/tools/topo_expl/model.cpp @@ -92,7 +92,7 @@ ncclResult_t p2pCanConnect(int* ret, struct ncclComm* comm, struct ncclTopoGrap // Check topology / p2p level. int intermediateRank; - NCCLCHECK(ncclTopoCheckP2p(comm->topo, info1->rank, info2->rank, ret, NULL, &intermediateRank)); + NCCLCHECK(ncclTopoCheckP2p(comm, comm->topo, info1->rank, info2->rank, ret, NULL, &intermediateRank)); if (*ret == 0) return ncclSuccess; if (intermediateRank != -1) { if (useMemcpy) *ret = 0; @@ -214,7 +214,7 @@ ncclResult_t netSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st int64_t netId; if (connIndex == NCCL_CONN_IDX_P2P_NET) NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, myInfo->rank, graph, channelId, 1, &netId, &req.netDev)); if (req.netDev < 0) NCCLCHECK(ncclTopoGetNetDev(comm, myInfo->rank, graph, channelId, peerInfo->rank, &netId, &req.netDev, &proxyRank)); - NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, netId, 1, &req.useGdr)); + NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->rank, netId, 1, &req.useGdr)); if (proxyRank == myInfo->rank) { INFO(NCCL_INIT|NCCL_NET,"Channel %02d/%d : %d[%lx] -> %d[%lx] [send] via NET/%s/%d%s%s", channelId, connIndex, myInfo->rank, myInfo->busId, peerInfo->rank, peerInfo->busId, comm->ncclNet->name, req.netDev, @@ -242,7 +242,7 @@ ncclResult_t netRecvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st int64_t netId; if (connIndex == NCCL_CONN_IDX_P2P_NET) NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, myInfo->rank, graph, channelId, 0, &netId, &req.netDev)); if (req.netDev < 0) NCCLCHECK(ncclTopoGetNetDev(comm, myInfo->rank, graph, channelId, myInfo->rank, &netId, &req.netDev, &proxyRank)); - NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, netId, 0, &req.useGdr)); + NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->rank, netId, 0, &req.useGdr)); INFO(NCCL_INIT|NCCL_NET,"Channel %02d/%d : %d[%lx] -> %d[%lx] [receive] via NET/%s/%d%s%s", channelId, connIndex, peerInfo->rank, peerInfo->busId, myInfo->rank, myInfo->busId, comm->ncclNet->name, req.netDev, req.useGdr ? "/GDRDMA" : "", req.shared ? "/Shared" : ""); @@ -266,7 +266,7 @@ ncclResult_t collNetSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph int netDev, useGdr = 0, proxy; NCCLCHECK(ncclTopoGetNetDev(comm, myInfo->rank, graph, channelId, peerInfo->rank, NULL, &netDev, &proxy)); - NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, netDev, 1, &useGdr)); + NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->rank, netDev, 1, &useGdr)); INFO(NCCL_INIT|NCCL_NET,"Coll %02d : %d [send] via COLLNET/%s/%d%s", channelId, myInfo->rank, "SHARP", netDev, useGdr ? "/GDRDMA" : ""); return ncclSuccess; @@ -276,7 +276,7 @@ ncclResult_t collNetRecvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph int netDev, useGdr = 0, proxy; NCCLCHECK(ncclTopoGetNetDev(comm, myInfo->rank, graph, channelId, peerInfo->rank, NULL, &netDev, &proxy)); - NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, netDev, 0, &useGdr)); + NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->rank, netDev, 0, &useGdr)); INFO(NCCL_INIT|NCCL_NET,"Coll %02d : %d [receive] via COLLNET/%s/%d%s", channelId, myInfo->rank, "SHARP", netDev, useGdr ? "/GDRDMA" : ""); return ncclSuccess; diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp index 88cb3a38de..56455c833b 100644 --- a/tools/topo_expl/topo_expl.cpp +++ b/tools/topo_expl/topo_expl.cpp @@ -268,8 +268,7 @@ int main(int argc,char* argv[]) for (int a=0; a= 0 && time < minTime) { algorithm = a; protocol = p;