From d87dc7c2e81df5de3e3add32b08ce08d626c2a3b Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Thu, 25 Mar 2021 20:59:32 -0700 Subject: [PATCH] collnet: support multiple NICs (#335) --- src/enqueue.cc | 3 +- src/graph/connect.cc | 29 +++-- src/graph/search.cc | 36 +++++- src/graph/tuning.cc | 8 +- src/init.cc | 8 +- tools/scripts/topo_val.sh | 4 +- tools/topo_expl/include/utils.h | 2 + tools/topo_expl/models/topo_collnet_n4.xml | 126 +++++++++++++++++++++ tools/topo_expl/topo_expl.cpp | 14 +-- tools/topo_expl/utils.cpp | 15 ++- 10 files changed, 209 insertions(+), 36 deletions(-) create mode 100644 tools/topo_expl/models/topo_collnet_n4.xml diff --git a/src/enqueue.cc b/src/enqueue.cc index 92b99b94d0..b5365a492f 100644 --- a/src/enqueue.cc +++ b/src/enqueue.cc @@ -283,10 +283,11 @@ static ncclResult_t getAlgoInfo(struct ncclInfo* info) { info->algorithm = -1; info->protocol = -1; int nAlgos = NCCL_NUM_ALGORITHMS; + #define SHARP_COLL_SAT_THRESHOLD 16384 // Check collNet support int collNetTypeSupport = 0; - if (info->comm->collNetSupport) + if (info->comm->collNetSupport && info->nBytes < SHARP_COLL_SAT_THRESHOLD*comm->collNetnChannels/2) NCCLCHECK(collNetReduceSupport(info->datatype, info->op, &collNetTypeSupport)); if (collNetTypeSupport != 1) nAlgos--; for (int a=0; arank; int localRanks = comm->localRanks; - int nChannels = comm->nChannels; - for (int c=0; cnChannels; c++) { struct ncclChannel* channel = comm->channels+c; channel->ring.prev = channel->ring.next = -1; channel->tree.up = -1; for (int i=0; itree.down[i] = -1; - channel->collTree.up = -1; - for (int i=0; icollTree.down[i] = -1; int* ringIntra = ringGraph->intra+c*localRanks; int* treeIntra = treeGraph->intra+c*localRanks; - int* collNetIntra = collNetGraph->intra+c*localRanks; for (int i=0; itree.up = i == 0 ? -1 : treeIntra[i-1]; channel->tree.down[0] = i == localRanks-1 ? -1 : treeIntra[i+1]; } + } + topoRanks->ringPrev[c] = channel->ring.prev; + topoRanks->ringNext[c] = channel->ring.next; + } + // Duplicate channels rings/trees + struct ncclChannel* channel0 = comm->channels; + struct ncclChannel* channel1 = channel0+comm->nChannels; + memcpy(channel1, channel0, comm->nChannels*sizeof(struct ncclChannel)); + // Setup collnet tree + for (int c=0; ccollNetnChannels; c++) { + struct ncclChannel* channel = comm->channels+c; + channel->collTree.up = -1; + for (int i=0; icollTree.down[i] = -1; + + int* collNetIntra = collNetGraph->intra+c*localRanks; + + for (int i=0; icollTree.down[0] = collNetIntra[next]; } } - topoRanks->ringPrev[c] = channel->ring.prev; - topoRanks->ringNext[c] = channel->ring.next; } - // Duplicate channels rings/trees - struct ncclChannel* channel0 = comm->channels; - struct ncclChannel* channel1 = channel0+nChannels; - memcpy(channel1, channel0, nChannels*sizeof(struct ncclChannel)); return ncclSuccess; } diff --git a/src/graph/search.cc b/src/graph/search.cc index d6eb13fc2e..b4720e66fd 100644 --- a/src/graph/search.cc +++ b/src/graph/search.cc @@ -1153,9 +1153,32 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph 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)); + struct ncclTopoGraph tmpGraph; + memcpy(&tmpGraph, graph, sizeof(struct ncclTopoGraph)); + int nets[MAXCHANNELS], n = 0; + for (int i = 0; i < tmpGraph.nChannels; i++) { + int j; + for (j = 0; j < n; j++) { + if (nets[j] == tmpGraph.inter[i*2]) + break; + } + if (j >= n) + nets[n++] = tmpGraph.inter[i*2]; + } + for (int i = 0; i < n; i++) { + int j; + for (j = 0; j < tmpGraph.nChannels; j++) { + if (nets[i] == tmpGraph.inter[j*2]) + break; + } + if (j < tmpGraph.nChannels) { + memcpy(graph->intra+i*ngpus, &tmpGraph.intra[j*ngpus], ngpus*sizeof(int)); + memcpy(graph->inter+i*2, &tmpGraph.inter[j*2], 2*sizeof(int)); + } + } + memcpy(graph->intra+n*ngpus, graph->intra, ngpus*sizeof(int)*n); + memcpy(graph->inter+n*2, graph->inter, 2*sizeof(int)*n); + graph->nChannels = n; } if (graph->nChannels) return ncclSuccess; @@ -1286,7 +1309,12 @@ done: graph->nChannels = 1; } - if (graph->speedIntra >= 25.0) { + if (graph->nChannels && graph->collNet) { + // duplicate collnet channels + memcpy(graph->intra+graph->nChannels*ngpus, graph->intra, ngpus*sizeof(int)*graph->nChannels); + memcpy(graph->inter+graph->nChannels*2, graph->inter, 2*sizeof(int)*graph->nChannels); + } + else if (graph->speedIntra >= 25.0) { int dupChannels = std::min(graph->nChannels*2, graph->maxChannels); memcpy(graph->intra+graph->nChannels*ngpus, graph->intra, (dupChannels-graph->nChannels)*ngpus*sizeof(int)); memcpy(graph->inter+graph->nChannels*2,graph->inter, (dupChannels-graph->nChannels)*2*sizeof(int)); diff --git a/src/graph/tuning.cc b/src/graph/tuning.cc index 7982c41ee1..88f19e0e90 100644 --- a/src/graph/tuning.cc +++ b/src/graph/tuning.cc @@ -63,11 +63,11 @@ static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 39.0 // Tree/Simple is the latency a 256kB chunk, which is ~ base lat + 256k/12GB/s (+ 256k/12GB/s for the network). static const float hwLat [3][NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { /* NVLINK */ - { /* Tree (LL/LL128/Simple)*/ { 2.5, 2.5, 5.5 }, /* Ring (LL/LL128/Simple)*/ { 2.5, 2.5, 5 }, /* CollNet (LL/LL128/Simple)*/ { 1.2, 1.2, 3.8 } }, + { /* Tree (LL/LL128/Simple)*/ { 2.5, 2.5, 5.5 }, /* Ring (LL/LL128/Simple)*/ { 2.5, 2.5, 5 }, /* CollNet (LL/LL128/Simple)*/ { 1.1, 1.1, 3.3 } }, /* PCI */ - { /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 1.3, 1.3, 1.9 }, /* CollNet (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 } }, + { /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 1.3, 1.3, 1.9 }, /* CollNet (LL/LL128/Simple)*/ { 1.1, 1.1, 1.7 } }, /* NET */ - { /* Tree (LL/LL128/Simple)*/ { 28.0, 28.0, 66.0 }, /* Ring (LL/LL128/Simple)*/ { 8.5, 8.5, 19.0 }, /* CollNet (LL/LL128/Simple)*/ { 9.8, 9.8, 19.5 } } + { /* Tree (LL/LL128/Simple)*/ { 28.0, 28.0, 66.0 }, /* Ring (LL/LL128/Simple)*/ { 8.5, 8.5, 19.0 }, /* CollNet (LL/LL128/Simple)*/ { 6.5, 6.5, 14.5 } } }; // LL128 max BW (per channel) for the different collectives @@ -148,7 +148,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_LL128) busBw = std::min(busBw * (nNodes == 1 ? 7.0/9.0 : 0.915 /*120.0/128.0*/), ll128MaxBwPerCh[coll]*graphs[a]->nChannels); #endif if (a == NCCL_ALGO_COLLNET) busBw *= .9; - if (a == NCCL_ALGO_COLLNET && p == NCCL_PROTO_LL) busBw *= 1.0/6.0; // Take into account that GDR read is disabled on both sides + if (a == NCCL_ALGO_COLLNET && p == NCCL_PROTO_LL) busBw *= 1.0/2.0; // Take into account that GDR read is disabled on both sides if (a == NCCL_ALGO_COLLNET && p == NCCL_PROTO_LL128) busBw = 0; // CollNet does not support LL128 // Convert bus BW to algorithm BW diff --git a/src/init.cc b/src/init.cc index deacdba573..0805935207 100644 --- a/src/init.cc +++ b/src/init.cc @@ -958,6 +958,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm allGather3Data[rank].collNet.typeIntra = collNetGraph.typeIntra; allGather3Data[rank].collNet.typeInter = collNetGraph.typeInter; + // CollNet channels are already duplicated + comm->collNetnChannels = 2*collNetGraph.nChannels; NCCLCHECK(ncclTopoPreset(comm, &treeGraph, &ringGraph, &collNetGraph, &allGather3Data[rank].topoRanks)); NCCLCHECK(bootstrapAllGather(comm->bootstrap, allGather3Data, sizeof(*allGather3Data))); @@ -1035,9 +1037,9 @@ 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)); + } else { + comm->collNetnChannels = 0; } free(allTopoRanks); @@ -1094,6 +1096,8 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm if (comm->nNodes > 1 && ncclParamCollNetEnable() == 1 && collNetSupport() && collNetGraph.nChannels) { + for (int c=comm->nChannels; ccollNetnChannels; c++) + NCCLCHECK(initChannel(comm, c));; int logicChannels = comm->collNetnChannels/2; int collNetSetupFail = 0; const int recvIndex = 0; // recv GPU index is always 0 diff --git a/tools/scripts/topo_val.sh b/tools/scripts/topo_val.sh index 4d27a33db6..5d55a33fc5 100755 --- a/tools/scripts/topo_val.sh +++ b/tools/scripts/topo_val.sh @@ -21,9 +21,9 @@ DIR="$(cd -P "$(dirname "${BASH_SOURCE[0]}")" && pwd)" -for i in {0..50} +for i in {0..51} do - if [[ $i -eq 50 ]] + if [[ $i -eq 50 ]] || [[ $i -eq 51 ]] then NCCL_COLLNET_ENABLE=1 $DIR/../topo_expl/topo_expl -m $i > "topo_m$i.log" else diff --git a/tools/topo_expl/include/utils.h b/tools/topo_expl/include/utils.h index be426f7f44..a111b8be29 100644 --- a/tools/topo_expl/include/utils.h +++ b/tools/topo_expl/include/utils.h @@ -36,6 +36,8 @@ struct allGather3Data_t{ struct ncclTopoRanks topoRanks; }; +void initCollNet(); + ncclResult_t ncclTopoGetSystem(const char* xmlTopoFile, struct ncclTopoSystem** system); ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem** topoSystem); diff --git a/tools/topo_expl/models/topo_collnet_n4.xml b/tools/topo_expl/models/topo_collnet_n4.xml new file mode 100644 index 0000000000..a453ce03e6 --- /dev/null +++ b/tools/topo_expl/models/topo_collnet_n4.xml @@ -0,0 +1,126 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp index d2029f8140..e4c284e41c 100644 --- a/tools/topo_expl/topo_expl.cpp +++ b/tools/topo_expl/topo_expl.cpp @@ -120,6 +120,7 @@ NodeModelDesc model_descs[] = { {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"}, + {4, "topo_collnet_n4.xml", "4 nodes collnet 4 NICs"}, }; int main(int argc,char* argv[]) @@ -148,6 +149,8 @@ int main(int argc,char* argv[]) NetworkModel network; NodeModel* node; + initCollNet(); + NodeModelDesc *desc = &model_descs[model_id]; for (int i=0; inum_nodes; i++) { node = new NodeModel(desc->filename); @@ -190,14 +193,9 @@ int main(int argc,char* argv[]) } struct ncclTopoGraph *treeGraph, *ringGraph, *collNetGraph; - treeGraph = (struct ncclTopoGraph *)malloc(sizeof(struct ncclTopoGraph)*nranks); - ringGraph = (struct ncclTopoGraph *)malloc(sizeof(struct ncclTopoGraph)*nranks); - collNetGraph = (struct ncclTopoGraph *)malloc(sizeof(struct ncclTopoGraph)*nranks); - if (!treeGraph || !ringGraph || !collNetGraph) { - printf("Failed to allocate memory for graphs\n"); - return -1; - } - + NCCLCHECK(ncclCalloc(&treeGraph, nranks)); + NCCLCHECK(ncclCalloc(&ringGraph, nranks)); + NCCLCHECK(ncclCalloc(&collNetGraph, nranks)); for (int i = 0; i < nranks; i++) { node_model = network.GetNode(i); assert(node_model!=0); diff --git a/tools/topo_expl/utils.cpp b/tools/topo_expl/utils.cpp index a69cd6c4ca..3ea1e8a13c 100644 --- a/tools/topo_expl/utils.cpp +++ b/tools/topo_expl/utils.cpp @@ -231,13 +231,16 @@ static ncclResult_t checkCollNetSetup(struct ncclComm* comm, int rank, int collN return ncclSuccess; } +void initCollNet() { + if (ncclParamCollNetEnable() == 1 && ncclCollNet == 0) + ncclCollNet = (ncclCollNet_t*)0x12345678; +} + 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)); @@ -401,6 +404,8 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t allGather3Data[rank].collNet.typeIntra = collNetGraph.typeIntra; allGather3Data[rank].collNet.typeInter = collNetGraph.typeInter; + // CollNet channels are already duplicated + comm->collNetnChannels = 2*collNetGraph.nChannels; NCCLCHECK(ncclTopoPreset(comm, &treeGraph, &ringGraph, &collNetGraph, &allGather3Data[rank].topoRanks)); return ncclSuccess; @@ -624,9 +629,9 @@ 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)); + } else { + comm->collNetnChannels = 0; } free(allTopoRanks); @@ -682,6 +687,8 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t if (comm->nNodes > 1 && ncclParamCollNetEnable() == 1 && collNetSupport() && collNetGraph.nChannels) { + for (int c=comm->nChannels; ccollNetnChannels; c++) + NCCLCHECK(initChannel(comm, c));; int logicChannels = comm->collNetnChannels/2; int collNetSetupFail = 0; const int recvIndex = 0; // recv GPU index is always 0