diff --git a/src/collectives/device/all_reduce.h b/src/collectives/device/all_reduce.h index bb573428ac..21cb89b39a 100644 --- a/src/collectives/device/all_reduce.h +++ b/src/collectives/device/all_reduce.h @@ -258,7 +258,7 @@ class ncclFunction= nChannels) { // second half of the channels do broadcast + if (blockIdx.x >= nChannels && blockIdx.x < 2*nChannels) { // second half of the channels do broadcast ncclPrimitives prims(tid, nthreads, &tree->up, tree->down, NULL, stepSize, channel, comm, ncclShmem->ptrs, 0); for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) { @@ -455,7 +455,7 @@ class ncclFunction= nChannels) { // second half of the channels do broadcast + if (blockIdx.x >= nChannels && blockIdx.x < 2*nChannels) { // second half of the channels do broadcast ncclLLPrimitives LLprims(tid, nthreads, &tree->up, tree->down, stepLines, channel, comm); for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) { // Down diff --git a/src/enqueue.cc b/src/enqueue.cc index 002fb20db8..92b99b94d0 100644 --- a/src/enqueue.cc +++ b/src/enqueue.cc @@ -308,7 +308,7 @@ static ncclResult_t getAlgoInfo(struct ncclInfo* info) { TRACE(NCCL_COLL, "%ld Bytes -> Algo %d proto %d time %f", info->nBytes, info->algorithm, info->protocol, minTime); int nc = (info->nChannels > 0) ? info->nChannels : - (info->algorithm == NCCL_ALGO_COLLNET) ? comm->nChannels/2 : comm->nChannels; // CollNet uses one channel for up and one channel for down + (info->algorithm == NCCL_ALGO_COLLNET) ? comm->collNetnChannels/2 : comm->nChannels; // CollNet uses one channel for up and one channel for down int nt = comm->maxThreads[info->algorithm][info->protocol]; int threadThreshold = comm->threadThresholds[info->algorithm][info->protocol]; while (info->nBytes < nc*nt*threadThreshold) { @@ -511,7 +511,7 @@ ncclResult_t ncclSaveKernel(struct ncclInfo* info) { proxyArgs.channel = channel; // Adjust pattern for CollNet based on channel index if (nSubChannels == 2) { - info->pattern = (channelId < info->comm->nChannels/nSubChannels) ? ncclPatternCollTreeUp : ncclPatternCollTreeDown; + info->pattern = (channelId < info->comm->collNetnChannels/nSubChannels) ? ncclPatternCollTreeUp : ncclPatternCollTreeDown; } if (proxyArgs.nsteps) NCCLCHECK(ncclProxySaveColl(&proxyArgs, info->pattern, info->root, info->comm->nRanks)); diff --git a/src/graph/connect.cc b/src/graph/connect.cc index 653e560ab6..620c8ce415 100644 --- a/src/graph/connect.cc +++ b/src/graph/connect.cc @@ -174,7 +174,7 @@ ncclResult_t ncclTopoConnectCollNet(struct ncclComm* comm, struct ncclTopoGraph* int depth = nranks/comm->nNodes; int sendIndex = collNetGraph->pattern == NCCL_TOPO_PATTERN_TREE ? 0 : 1; // send GPU index depends on topo pattern int sendEndIndex = (sendIndex+comm->localRanks-1)%comm->localRanks; - for (int c=0; cnChannels/2; c++) { + for (int c=0; ccollNetnChannels/2; c++) { struct ncclChannel* channel = comm->channels+c; // Set root of collTree to id nranks if (rank == collNetGraph->intra[sendIndex+c*comm->localRanks]) { // is master @@ -188,8 +188,8 @@ ncclResult_t ncclTopoConnectCollNet(struct ncclComm* comm, struct ncclTopoGraph* } int recvIndex = 0; // recv GPU index is always 0 int recvEndIndex = (recvIndex+comm->localRanks-1)%comm->localRanks; - for (int c=0; cnChannels/2; c++) { - struct ncclChannel* channel = comm->channels+comm->nChannels/2+c; + for (int c=0; ccollNetnChannels/2; c++) { + struct ncclChannel* channel = comm->channels+comm->collNetnChannels/2+c; // Set root of collTree to id nranks if (rank == collNetGraph->intra[recvIndex+c*comm->localRanks]) { // is master channel->collTree.up = nranks; @@ -198,7 +198,7 @@ ncclResult_t ncclTopoConnectCollNet(struct ncclComm* comm, struct ncclTopoGraph* channel->collTree.down[0] = -1; } channel->collTree.depth = depth; - INFO(NCCL_GRAPH, "CollNet Channel %d rank %d up %d down %d", comm->nChannels/2+c, rank, channel->collTree.up, channel->collTree.down[0]); + INFO(NCCL_GRAPH, "CollNet Channel %d rank %d up %d down %d", comm->collNetnChannels/2+c, rank, channel->collTree.up, channel->collTree.down[0]); } return ncclSuccess; } diff --git a/src/graph/paths.cc b/src/graph/paths.cc index 92ed788b67..92758c154a 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -572,6 +572,6 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) { for (int b=1, mb=(comm->p2pnChannels>>1); bp2pnChannels; b<<=1, mb>>=1) if (c & b) mirror |= mb; comm->p2pChannels[c] = mirror; } - INFO(NCCL_INIT, "%d coll channels, %d p2p channels, %d p2p channels per peer", comm->nChannels, comm->p2pnChannels, comm->p2pnChannelsPerPeer); + INFO(NCCL_INIT, "%d coll channels, %d collnet channels, %d p2p channels, %d p2p channels per peer", comm->nChannels, comm->collNetnChannels, comm->p2pnChannels, comm->p2pnChannelsPerPeer); return ncclSuccess; } diff --git a/src/graph/search.cc b/src/graph/search.cc index 2c00068eb9..d6eb13fc2e 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -1112,6 +1112,8 @@ float speedArray[] = { 42.0, 30.0, 24.0, 21.0, 18.0, 15.0, 12.0, 10.0, 9.0, 7.0, #endif #define NSPEEDS (sizeof(speedArray)/sizeof(float)) +RCCL_PARAM(ModelMatchingDisable, "MODEL_MATCHING_DISABLE", 0); + ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph) { int ngpus = system->nodes[GPU].count; int nnets = system->nodes[NET].count; @@ -1143,13 +1145,18 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph if (graph->nChannels) { system->type |= RCCL_TOPO_4P2H_ROME; } - } else { + } else if (!rcclParamModelMatchingDisable()) { // try to match 8P6L NCCLCHECK(parseChordalRing(system, graph)); if (graph->nChannels) return ncclSuccess; // try to match Rome 4P2H NCCLCHECK(parseRome4P2H(system, graph)); } + if (graph->collNet && graph->nChannels) { + graph->nChannels = 1; + memcpy(graph->intra+graph->nChannels*ngpus, graph->intra, ngpus*sizeof(int)); + memcpy(graph->inter+graph->nChannels*2, graph->inter, 2*sizeof(int)); + } if (graph->nChannels) return ncclSuccess; if (ngpus == 1) if (graph->pattern != NCCL_TOPO_PATTERN_RING) graph->pattern = NCCL_TOPO_PATTERN_TREE; diff --git a/src/include/comm.h b/src/include/comm.h index cbdb66a9ea..665e259c68 100644 --- a/src/include/comm.h +++ b/src/include/comm.h @@ -101,6 +101,8 @@ struct ncclComm { int p2pnChannels; int p2pnChannelsPerPeer; int p2pChannels[MAXCHANNELS]; + //Channels for collnet + int collNetnChannels; // Buffer sizes int buffSizes[NCCL_NUM_PROTOCOLS]; diff --git a/src/init.cc b/src/init.cc index 1ccbd83c8e..deacdba573 100644 --- a/src/init.cc +++ b/src/init.cc @@ -717,7 +717,7 @@ static ncclResult_t checkCollNetSetup(struct ncclComm* comm, int rank, int collN if (collNetSetupFail) { if (rank == 0) WARN("Cannot initialize CollNet, using %s instead", ncclNetName()); // Free collNet resources - for (int r=0; rnChannels; r++) { + for (int r=0; rcollNetnChannels; r++) { struct ncclChannel* channel = comm->channels+r; struct ncclPeer* peer = channel->peers+nranks; if (peer->send.transportResources && peer->send.transportComm) NCCLCHECK(peer->send.transportComm->free(peer->send.transportResources)); @@ -1035,6 +1035,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm if (comm->nNodes > 1 && ncclParamCollNetEnable() == 1 && collNetSupport() && collNetGraph.nChannels) { + // Force 2 channels for CollNet + comm->collNetnChannels = collNetGraph.nChannels = 2; NCCLCHECK(ncclTopoConnectCollNet(comm, &collNetGraph, rank)); } @@ -1092,7 +1094,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm if (comm->nNodes > 1 && ncclParamCollNetEnable() == 1 && collNetSupport() && collNetGraph.nChannels) { - int logicChannels = comm->nChannels/2; + int logicChannels = comm->collNetnChannels/2; int collNetSetupFail = 0; const int recvIndex = 0; // recv GPU index is always 0 const int sendIndex = collNetGraph.pattern == NCCL_TOPO_PATTERN_TREE ? 0 : 1; // send GPU index depends on topo pattern diff --git a/src/transport/coll_net.cc b/src/transport/coll_net.cc index 5544be8889..e23873b65c 100644 --- a/src/transport/coll_net.cc +++ b/src/transport/coll_net.cc @@ -9,12 +9,14 @@ #include "coll_net.h" #include "graph.h" #include +#include struct collNetRecvConnectInfo { collNetHandle_t collNetHandle; }; struct collNetSendConnectInfo { + collNetHandle_t collNetHandle; void* collNetComm; void* mhandles[NCCL_NUM_PROTOCOLS]; struct reqSlot* reqFifo; @@ -39,6 +41,7 @@ struct collNetSendResources { uint64_t llLastCleaning; struct reqSlot* reqFifo; int collNetRank; + uint32_t* curr_hdp_reg; // Curr GPU in ring (for rdma transport use only) }; struct collNetRecvResources { @@ -55,6 +58,7 @@ struct collNetRecvResources { uint64_t llLastCleaning; struct reqSlot* reqFifo; int collNetRank; + uint32_t* curr_hdp_reg; // Curr GPU in ring (for rdma transport use only) }; /* Determine if we can communicate with the peer */ @@ -78,7 +82,28 @@ ncclResult_t collNetSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph for (int p=0; pcomm->buffSizes[p]; if (resources->useGdr) { - NCCLCHECK(ncclCudaCalloc((char**)(&resources->devRecvMem), recvSize)); + NCCLCHECK(ncclCudaCalloc((char**)(&resources->devRecvMem), recvSize, resources->useGdr)); + //CUDACHECK(hipDeviceGetAttribute((int*)&resources->curr_hdp_reg, hipDeviceAttributeHdpMemFlushCntl, myInfo->cudaDev)); + struct data_struct {hsa_agent_t agent; int counter;} out; + out.counter = 0; + out.agent.handle = myInfo->cudaDev; + hsa_iterate_agents([](hsa_agent_t agent, void* data) { + int devId = ((struct data_struct *)data)->agent.handle; + hsa_device_type_t type; + hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); + if(type != HSA_DEVICE_TYPE_GPU) + return HSA_STATUS_SUCCESS; + if(((struct data_struct *)data)->counter!=devId) { + ((struct data_struct *)data)->counter++; + return HSA_STATUS_SUCCESS; + } + ((struct data_struct *)data)->agent = agent; + return HSA_STATUS_SUCCESS; + }, (void*)&out); + hsa_amd_hdp_flush_t hdpinfo; + hsa_status_t err = hsa_agent_get_info(out.agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_HDP_FLUSH, &hdpinfo); + resources->curr_hdp_reg = hdpinfo.HDP_MEM_FLUSH_CNTL; + send->conn.curr_hdp_reg = resources->curr_hdp_reg; } NCCLCHECK(ncclCudaHostCalloc((char**)&resources->recvMem, recvSize)); NCCLCHECK(ncclIbMalloc((void**)&(resources->llData), send->comm->buffSizes[NCCL_PROTO_LL]/2)); @@ -103,7 +128,7 @@ ncclResult_t collNetRecvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph for (int p=0; pcomm->buffSizes[p]; if (resources->useGdr) { - NCCLCHECK(ncclCudaCalloc((char**)(&resources->devRecvMem), recvSize)); + NCCLCHECK(ncclCudaCalloc((char**)(&resources->devRecvMem), recvSize, resources->useGdr)); } NCCLCHECK(ncclCudaHostCalloc((char**)&resources->recvMem, recvSize)); @@ -291,6 +316,11 @@ ncclResult_t collNetSendProxy(struct ncclProxyArgs* args) { size = nFifoLines*2*sizeof(uint32_t); } if (ready) { + // flush HDP if not done + if (resources->curr_hdp_reg && args->hdp_flushed < LOAD(recvTail)) { + args->hdp_flushed = LOAD(recvTail); + STORE(resources->curr_hdp_reg, 1); + } // Data is ready, try to send. int count = size/ncclTypeSize(args->dtype); NCCLCHECK(collNetIallreduce(resources->collNetSendComm, (void*) buff, (void*)(reqFifo[buffSlot].recvBuff), count, args->dtype, args->redOp, sendMhandle, recvMhandle, args->requests+buffSlot)); diff --git a/tools/TopoVisual/extract_topo.awk b/tools/TopoVisual/extract_topo.awk index c5bd729a6e..b9da532341 100755 --- a/tools/TopoVisual/extract_topo.awk +++ b/tools/TopoVisual/extract_topo.awk @@ -1,5 +1,5 @@ #!/usr/bin/gawk -f -# Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2019-2021 Advanced Micro Devices, Inc. All rights reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal @@ -26,8 +26,15 @@ BEGIN { treedns[""]=0 max_treedn=0 conn[""]=0 + has_collnet=0 + max_collnet=0 + max_collnet_rank=0 + collnet[""]=0 + collnet_conn[""]=0 + collnet_conn_type[""]=0 col_start=2 col_p1=col_start+1 + col_p2=col_start+2 col_p3=col_start+3 col_p4=col_start+4 col_p5=col_start+5 @@ -40,6 +47,7 @@ BEGIN { if($3=="NCCL" && $4=="INFO" && col_start==2) { col_start=5 col_p1=col_start+1 + col_p2=col_start+2 col_p3=col_start+3 col_p4=col_start+4 col_p5=col_start+5 @@ -48,7 +56,19 @@ BEGIN { col_p8=col_start+8 } - if($col_start=="Ring" && $col_p4="->" && $col_p6=="->") { + if($5=="NCCL" && $6=="INFO" && col_start==2) { + col_start=7 + col_p1=col_start+1 + col_p2=col_start+2 + col_p3=col_start+3 + col_p4=col_start+4 + col_p5=col_start+5 + col_p6=col_start+6 + col_p7=col_start+7 + col_p8=col_start+8 + } + + if($col_start=="Ring" && $col_p4=="->" && $col_p6=="->") { chan=strtonum($col_p1) rank=strtonum($col_p5) next_rank=strtonum($col_p7) @@ -91,6 +111,31 @@ BEGIN { } while ($col_1!="") } + if($col_start=="CollNet" && $col_p1=="Channel") { + chan=strtonum($col_p2) + rank=strtonum($col_p4) + up_rank=strtonum($col_p6) + collnet[up_rank "," rank "," chan]="1" + if(has_collnet==0) + has_collnet=1 + if(chan>max_collnet) + max_collnet=chan + if(up_rank>max_collnet_rank) + max_collnet_rank=up_rank + } + + if($col_start=="Coll" && $col_p2==":") { + chan=strtonum($col_p1) + rank=strtonum($col_p3) + if($col_p4=="[receive]") + collnet_conn[rank "," chan]=0 + else if($col_p4=="[send]") + collnet_conn[rank "," chan]=1 + else + printf "Error!\n" + collnet_conn_type[rank "," chan]=$col_p6 + } + if($col_p6=="via") { match($col_p1, /([0-9]+)/, ary) chan=strtonum(ary[1]) @@ -180,5 +225,67 @@ END { } printf " }\n\n" } + + for(r=0; has_collnet && r<=max_collnet; r++) { + printf " subgraph collnet_%d {\n", r + num_top_ranks=0 + for(s=0;s c%d_%d [dir=back label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, max_collnet_rank, r, rank, val, color, style + else + printf " c%d_%d -> c%d_%d [label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, max_collnet_rank, r, rank, val, color, style + while(1) { + for(s=0;s c%d_%d [dir=back label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, rank, r, s, val, color, style + else + printf " c%d_%d -> c%d_%d [label=\"%s\",color=\"%s\",style=\"%s\",fontname=\"Helvetica\"];\n", r, rank, r, s, val, color, style + rank=s + break; + } + } + if(s>=max_collnet_rank) { + break; + } + } + } + printf "\n" + for(s=0;s<=max_collnet_rank;s++) { + printf " c%d_%d [label=\"%d\",fontsize=\"28\"];\n", r, s, s + } + printf " }\n\n" + } printf "}\n" } diff --git a/tools/scripts/topo_val.sh b/tools/scripts/topo_val.sh index 07d8d050d5..4d27a33db6 100755 --- a/tools/scripts/topo_val.sh +++ b/tools/scripts/topo_val.sh @@ -21,8 +21,13 @@ DIR="$(cd -P "$(dirname "${BASH_SOURCE[0]}")" && pwd)" -for i in {0..49} +for i in {0..50} do - $DIR/../topo_expl/topo_expl -m $i > "topo_m$i.log" + if [[ $i -eq 50 ]] + then + NCCL_COLLNET_ENABLE=1 $DIR/../topo_expl/topo_expl -m $i > "topo_m$i.log" + else + $DIR/../topo_expl/topo_expl -m $i > "topo_m$i.log" + fi $DIR/../TopoVisual/topo_visual.sh -i "topo_m$i.log" done diff --git a/tools/topo_expl/model.cpp b/tools/topo_expl/model.cpp index 59193349df..4585f3048d 100644 --- a/tools/topo_expl/model.cpp +++ b/tools/topo_expl/model.cpp @@ -159,8 +159,6 @@ ncclResult_t netSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st return ncclSuccess; } -NCCL_PARAM(NetGdrLevel, "NET_GDR_LEVEL", PATH_PHB); - ncclResult_t netRecvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* recv, int channelId) { int netDev, useGdr = 0; @@ -179,6 +177,39 @@ struct ncclTransport netTransport = { { netRecvSetup, NULL, NULL, NULL } }; +/* Determine if two peers can communicate with NET */ +ncclResult_t collNetCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) { + *ret = 1; + return ncclSuccess; +} + +ncclResult_t collNetSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* send, int channelId) { + int netDev, useGdr = 0; + + NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, &netDev)); + NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, 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; +} + +ncclResult_t collNetRecvSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclPeerInfo* myInfo, struct ncclPeerInfo* peerInfo, struct ncclConnect* connectInfo, struct ncclConnector* recv, int channelId) { + int netDev, useGdr = 0; + + NCCLCHECK(ncclTopoGetNetDev(comm->topo, myInfo->rank, graph, channelId, &netDev)); + NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, 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; +} + +struct ncclTransport collNetTransport = { + "COL", + collNetCanConnect, + { collNetSendSetup, NULL, NULL, NULL }, + { collNetRecvSetup, NULL, NULL, NULL } +}; + struct ncclTransport ncclTransports[NTRANSPORTS] = { p2pTransport, shmTransport, diff --git a/tools/topo_expl/models/topo_collnet_n1.xml b/tools/topo_expl/models/topo_collnet_n1.xml new file mode 100644 index 0000000000..6e63daf649 --- /dev/null +++ b/tools/topo_expl/models/topo_collnet_n1.xml @@ -0,0 +1,87 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp index 77b0597e51..d2029f8140 100644 --- a/tools/topo_expl/topo_expl.cpp +++ b/tools/topo_expl/topo_expl.cpp @@ -119,6 +119,7 @@ NodeModelDesc model_descs[] = { {2, "topo_8p_rome_pcie.xml", "2 nodes node 8 VEGA20 PCIe"}, {1, "topo_8p_rome_4nics.xml", "single node 8 gfx908 Rome 4 NICs"}, {4, "topo_8p_rome_4nics.xml", "4 nodes node 8 gfx908 Rome 4 NICs"}, + {4, "topo_collnet_n1.xml", "4 nodes collnet 1 NICs"}, }; int main(int argc,char* argv[]) diff --git a/tools/topo_expl/utils.cpp b/tools/topo_expl/utils.cpp index d8a438bea6..a69cd6c4ca 100644 --- a/tools/topo_expl/utils.cpp +++ b/tools/topo_expl/utils.cpp @@ -41,7 +41,7 @@ NCCL_PARAM(CollNetEnable, "COLLNET_ENABLE", 0); NCCL_PARAM(GraphDumpFileRank, "GRAPH_DUMP_FILE_RANK", 0); thread_local int ncclDebugNoWarn = 0; -ncclCollNet_t* ncclCollNet = NULL; +ncclCollNet_t* ncclCollNet = 0; // Get current Compute Capability int ncclCudaCompCap() { @@ -137,10 +137,107 @@ ncclResult_t bootstrapAllGather(struct ncclComm* comm, struct allGather1Data_t * return ncclSuccess; } +extern struct ncclTransport collNetTransport; + +// All ranks must participate in collNetSetup call +// type: 0 for send, 1 for recv +// return: 0 - unsupported, 1 - supported +// We do not NCCLCHECK this call because we would fall back to P2P network in case CollNet setup fails +static int collNetSetup(struct ncclComm* comm, struct ncclTopoGraph* collNetGraph, struct ncclChannel* channel, int rank, int nranks, int masterRank, int masterPeer, int nMasters, int type) { + int rankInCollNet = -1; + int supported = 0; + int isMaster = (rank == masterRank) ? 1 : 0; + struct { + int collNetRank; + ncclConnect connect; + } sendrecvExchange; + + // check if we can connect to collnet, whose root is the nranks-th rank + struct ncclPeerInfo *myInfo = comm->peerInfo+rank, *peerInfo = comm->peerInfo+nranks; + peerInfo->rank = nranks; + int ret = 1; + if (isMaster) { + NCCLCHECK(collNetTransport.canConnect(&ret, comm->topo, collNetGraph, myInfo, peerInfo)); + } + + // send master receives connect info from peer recv master + if (isMaster && type == 0) { + //NCCLCHECK(bootstrapRecv(comm->bootstrap, masterPeer, &sendrecvExchange, sizeof(sendrecvExchange))); + rankInCollNet = sendrecvExchange.collNetRank; + INFO(NCCL_INIT, "CollNet [send] : rank %d collNetRank %d collNetNranks %d received connect from rank %d", rank, rankInCollNet, nMasters, masterPeer); + } + + // select + struct ncclPeer* root = channel->peers+nranks; + struct ncclConnector* conn = (type == 1) ? &root->recv : &root->send; + struct ncclTransportComm* transportComm = (type == 1) ? &(collNetTransport.recv) : &(collNetTransport.send); + conn->transportComm = transportComm; + // setup + struct ncclConnect myConnect; + if (isMaster && ret > 0) { + NCCLCHECK(transportComm->setup(comm, collNetGraph, myInfo, peerInfo, &myConnect, conn, channel->id)); + } + // prepare connect handles + ncclResult_t res; + struct { + int isMaster; + ncclConnect connect; + } *allConnects = NULL; + ncclConnect *masterConnects = NULL; + //NCCLCHECK(ncclCalloc(&masterConnects, nMasters)); + if (type == 1) { // recv side: AllGather + // all ranks must participate + //NCCLCHECK(ncclCalloc(&allConnects, nranks)); + //allConnects[rank].isMaster = isMaster; + //memcpy(&(allConnects[rank].connect), &myConnect, sizeof(struct ncclConnect)); + //NCCLCHECKGOTO(bootstrapAllGather(comm->bootstrap, allConnects, sizeof(*allConnects)), res, cleanup); + // consolidate + //int c = 0; + //for (int r = 0; r < nranks; r++) { + //if (allConnects[r].isMaster) { + //memcpy(masterConnects+c, &(allConnects[r].connect), sizeof(struct ncclConnect)); + //if (r == rank) rankInCollNet = c; + //c++; + //} + //} + } else { // send side : copy in connect info received from peer recv master + //if (isMaster) memcpy(masterConnects+rankInCollNet, &(sendrecvExchange.connect), sizeof(struct ncclConnect)); + } + // connect + if (isMaster && ret > 0) { + //NCCLCHECKGOTO(transportComm->connect(comm, masterConnects, nMasters, rankInCollNet, conn), res, cleanup); + //struct ncclPeer* devRoot = channel->devPeers+nranks; + //struct ncclConnector* devConn = (type == 1) ? &devRoot->recv : &devRoot->send; + //CUDACHECKGOTO(hipMemcpy(devConn, conn, sizeof(struct ncclConnector), hipMemcpyHostToDevice), res, cleanup); + } + // recv side sends connect info to send side + if (isMaster && type == 1) { + //sendrecvExchange.collNetRank = rankInCollNet; + //memcpy(&sendrecvExchange.connect, masterConnects+rankInCollNet, sizeof(struct ncclConnect)); + //NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, masterPeer, &sendrecvExchange, sizeof(sendrecvExchange)), res, cleanup); + INFO(NCCL_INIT, "CollNet [recv] : rank %d collNetRank %d collNetNranks %d sent connect to rank %d", rank, rankInCollNet, nMasters, masterPeer); + } + if (ret > 0) { + supported = 1; + } +cleanup: + //if (allConnects != NULL) free(allConnects); + //if (masterConnects != NULL) free(masterConnects); + return supported; +} + +static ncclResult_t checkCollNetSetup(struct ncclComm* comm, int rank, int collNetSetupFail) { + comm->collNetSupport = 1; + return ncclSuccess; +} + ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data, struct allGather3Data_t *allGather3Data, struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph) { int rank = comm->rank; int nranks = comm->nRanks; + + if (ncclParamCollNetEnable() == 1 && ncclCollNet == 0) + ncclCollNet = (ncclCollNet_t*)0x12345678; //uint64_t commHash = getHash(commId->internal, NCCL_UNIQUE_ID_BYTES); //TRACE(NCCL_INIT, "comm %p, commHash %lx, rank %d nranks %d - BEGIN", comm, commHash, rank, nranks); //NCCLCHECK(bootstrapInit(commId, rank, nranks, &comm->bootstrap)); @@ -527,6 +624,8 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t if (comm->nNodes > 1 && ncclParamCollNetEnable() == 1 && collNetSupport() && collNetGraph.nChannels) { + // Force 2 channels for CollNet + comm->collNetnChannels = collNetGraph.nChannels = 2; NCCLCHECK(ncclTopoConnectCollNet(comm, &collNetGraph, rank)); } @@ -580,11 +679,10 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t INFO(NCCL_INIT, "Connected all trees"); // Check if we can setup CollNet -#if 0 if (comm->nNodes > 1 && ncclParamCollNetEnable() == 1 && collNetSupport() && collNetGraph.nChannels) { - int logicChannels = comm->nChannels/2; + int logicChannels = comm->collNetnChannels/2; int collNetSetupFail = 0; const int recvIndex = 0; // recv GPU index is always 0 const int sendIndex = collNetGraph.pattern == NCCL_TOPO_PATTERN_TREE ? 0 : 1; // send GPU index depends on topo pattern @@ -604,7 +702,6 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t // Verify CollNet setup across ranks NCCLCHECK(checkCollNetSetup(comm, rank, collNetSetupFail)); } -#endif TRACE(NCCL_INIT, "rank %d nranks %d - CONNECTED %d RINGS AND TREES", rank, nranks, comm->nChannels); free(rings);