diff --git a/projects/rccl/src/graph/connect.cc b/projects/rccl/src/graph/connect.cc index 24b66f481c..7f52375ae8 100644 --- a/projects/rccl/src/graph/connect.cc +++ b/projects/rccl/src/graph/connect.cc @@ -240,7 +240,7 @@ int ncclMaxNchannels() { return maxNchannels; } -ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePatterns, struct ncclTopoRanks** allTopoRanks, int* rings, int gcn, int nnets) { +ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePatterns, struct ncclTopoRanks** allTopoRanks, int* rings, int nc) { // Gather data from all ranks int *ringRecv, *ringSend, *ringPrev, *ringNext, *treeToParent, *treeToChild0, *treeToChild1; int nranks = comm->nRanks; @@ -272,13 +272,6 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa memcpy(ringPrev+nChannels*nranks, ringPrev, nChannels*nranks*sizeof(int)); memcpy(ringNext+nChannels*nranks, ringNext, nChannels*nranks*sizeof(int)); - int nc = nChannels*2; - if (gcn == 908) nc = std::max(nc, 4); - if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G)) nc = nChannels*4; - if (!nnets) nnets = comm->topo->nodes[NET].count; - if (nnets && (comm->topo->type & RCCL_TOPO_4P2H_ROME)) nc = (nnets > 3 ? 2 : 4)*nnets; - int end = std::min((int)ncclMaxNchannels(), std::max(nc, ncclMinNchannels())); - // Duplication should be complete now nChannels = comm->nChannels = std::min(MAXCHANNELS,nChannels*2); @@ -286,7 +279,7 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa // We permit combining max, then min, to only use the first channels, then duplicate them. nChannels = comm->nChannels = std::min((int)ncclMaxNchannels(), nChannels); int c; - for (c=nChannels; cchannels+c, comm->channels+c-nChannels, sizeof(struct ncclChannel)); diff --git a/projects/rccl/src/include/graph.h b/projects/rccl/src/include/graph.h index ac7fcbf992..893db11814 100644 --- a/projects/rccl/src/include/graph.h +++ b/projects/rccl/src/include/graph.h @@ -97,7 +97,7 @@ ncclResult_t ncclTopoPreset(struct ncclComm* comm, struct ncclTopoRanks* topoRanks); ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePatterns, - struct ncclTopoRanks** allTopoRanks, int* rings, int gcn, int nnets); + struct ncclTopoRanks** allTopoRanks, int* rings, int nc); ncclResult_t ncclTopoConnectCollNet(struct ncclComm* comm, struct ncclTopoGraph* collNetGraph, int rank); diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 512114fd05..cefe38cad6 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -924,7 +924,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm int cudaCompCap; int fullCudaCompCap; int nChannels; - int gcn; + int nc; struct ncclGraphInfo tree; struct ncclGraphInfo ring; struct ncclGraphInfo collNet; @@ -935,10 +935,14 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm int idx; NCCLCHECK(ncclTopoIdToIndex(comm->topo, GPU, myInfo->busId, &idx)); allGather3Data[rank].cudaCompCap = comm->topo->nodes[GPU].nodes[idx].gpu.cudaCompCap; - allGather3Data[rank].gcn = comm->topo->nodes[GPU].nodes[idx].gpu.gcn; - allGather3Data[rank].nChannels = comm->nChannels = treeGraph.nChannels = ringGraph.nChannels = std::min(treeGraph.nChannels, ringGraph.nChannels); + allGather3Data[rank].nc = comm->nChannels*2; + if (comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 908) allGather3Data[rank].nc = std::max(allGather3Data[rank].nc, 4); + if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G)) + allGather3Data[rank].nc = comm->nChannels*4; + if (comm->topo->nodes[GPU].count != comm->topo->nRanks && comm->topo->nodes[NET].count && (comm->topo->type & RCCL_TOPO_4P2H_ROME)) + allGather3Data[rank].nc = (comm->topo->nodes[NET].count > 3 ? 2 : 4)*comm->topo->nodes[NET].count; allGather3Data[rank].tree.pattern = treeGraph.pattern; allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels; allGather3Data[rank].tree.speedIntra = treeGraph.speedIntra; @@ -986,10 +990,10 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm int nChannelsOrig = comm->nChannels; struct ncclTopoRanks** allTopoRanks; NCCLCHECK(ncclCalloc(&allTopoRanks, comm->nRanks)); - int gcn = allGather3Data[0].gcn; + int nc = allGather3Data[0].nc; for (int i=0; inChannels = std::min(allGather3Data[i].nChannels, comm->nChannels); treeGraph.sameChannels = std::min(allGather3Data[i].tree.sameChannels, treeGraph.sameChannels); @@ -1009,22 +1013,6 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm collNetGraph.typeInter = std::min(allGather3Data[i].collNet.typeInter, collNetGraph.typeInter); } - // count NETs used by ring - int nNets = 0; - int nets[MAXCHANNELS*2]; - // do not count NETs in case of single node, i.e comm->topo->nodes[GPU].count == comm->topo->nRanks - for (int i = 0; comm->topo->nodes[GPU].count != comm->topo->nRanks && i < ringGraph.nChannels; i++) { - for (int j = 0; j < 2; j++) { - int k; - for (k = 0; k < nNets; k++) - if (nets[k] == ringGraph.inter[i*2+j]) break; - if (k >= nNets) { - nets[nNets] = ringGraph.inter[i*2+j]; - nNets++; - } - } - } - if (comm->nChannels < nChannelsOrig) { // We started duplicating channels during Preset(), so we need to move the // duplicated channels since we have removed some. @@ -1034,7 +1022,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm int *rings; NCCLCHECK(ncclCalloc(&rings, nranks*MAXCHANNELS)); - NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, nodesTreePatterns, allTopoRanks, rings, gcn, nNets)); + NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, nodesTreePatterns, allTopoRanks, rings, nc)); if (comm->nNodes > 1 && ncclParamCollNetEnable() == 1 && collNetSupport() && collNetGraph.nChannels) { diff --git a/projects/rccl/tools/topo_expl/include/utils.h b/projects/rccl/tools/topo_expl/include/utils.h index a111b8be29..62dd89d1cb 100644 --- a/projects/rccl/tools/topo_expl/include/utils.h +++ b/projects/rccl/tools/topo_expl/include/utils.h @@ -28,7 +28,7 @@ struct allGather3Data_t{ int cudaCompCap; int fullCudaCompCap; int nChannels; - int gcn; + int nc; int alltoallDisable; struct ncclGraphInfo tree; struct ncclGraphInfo ring; diff --git a/projects/rccl/tools/topo_expl/utils.cpp b/projects/rccl/tools/topo_expl/utils.cpp index a18cd3e590..f6c0cee8e7 100644 --- a/projects/rccl/tools/topo_expl/utils.cpp +++ b/projects/rccl/tools/topo_expl/utils.cpp @@ -369,7 +369,7 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t int cudaCompCap; int fullCudaCompCap; int nChannels; - int gcn; + int nc; struct ncclGraphInfo tree; struct ncclGraphInfo ring; struct ncclGraphInfo collNet; @@ -381,10 +381,14 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t int idx; NCCLCHECK(ncclTopoIdToIndex(comm->topo, GPU, myInfo->busId, &idx)); allGather3Data[rank].cudaCompCap = comm->topo->nodes[GPU].nodes[idx].gpu.cudaCompCap; - allGather3Data[rank].gcn = comm->topo->nodes[GPU].nodes[idx].gpu.gcn; - allGather3Data[rank].nChannels = comm->nChannels = treeGraph.nChannels = ringGraph.nChannels = std::min(treeGraph.nChannels, ringGraph.nChannels); + allGather3Data[rank].nc = comm->nChannels*2; + if (comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 908) allGather3Data[rank].nc = std::max(allGather3Data[rank].nc, 4); + if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G)) + allGather3Data[rank].nc = comm->nChannels*4; + if (comm->topo->nodes[GPU].count != comm->topo->nRanks && comm->topo->nodes[NET].count && (comm->topo->type & RCCL_TOPO_4P2H_ROME)) + allGather3Data[rank].nc = (comm->topo->nodes[NET].count > 3 ? 2 : 4)*comm->topo->nodes[NET].count; allGather3Data[rank].tree.pattern = treeGraph.pattern; allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels; allGather3Data[rank].tree.speedIntra = treeGraph.speedIntra; @@ -578,10 +582,10 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t int nChannelsOrig = comm->nChannels; struct ncclTopoRanks** allTopoRanks; NCCLCHECK(ncclCalloc(&allTopoRanks, comm->nRanks)); - int gcn = allGather3Data[0].gcn; + int nc = allGather3Data[0].nc; for (int i=0; inChannels = std::min(allGather3Data[i].nChannels, comm->nChannels); treeGraph.sameChannels = std::min(allGather3Data[i].tree.sameChannels, treeGraph.sameChannels); @@ -601,22 +605,6 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t collNetGraph.typeInter = std::min(allGather3Data[i].collNet.typeInter, collNetGraph.typeInter); } - // count NETs used by ring - int nNets = 0; - int nets[MAXCHANNELS*2]; - // do not count NETs in case of single node, i.e comm->topo->nodes[GPU].count == comm->topo->nRanks - for (int i = 0; comm->topo->nodes[GPU].count != comm->topo->nRanks && i < ringGraph.nChannels; i++) { - for (int j = 0; j < 2; j++) { - int k; - for (k = 0; k < nNets; k++) - if (nets[k] == ringGraph.inter[i*2+j]) break; - if (k >= nNets) { - nets[nNets] = ringGraph.inter[i*2+j]; - nNets++; - } - } - } - if (comm->nChannels < nChannelsOrig) { // We started duplicating channels during Preset(), so we need to move the // duplicated channels since we have removed some. @@ -626,7 +614,7 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t int *rings; NCCLCHECK(ncclCalloc(&rings, nranks*MAXCHANNELS)); - NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, nodesTreePatterns, allTopoRanks, rings, gcn, nNets)); + NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, nodesTreePatterns, allTopoRanks, rings, nc)); if (comm->nNodes > 1 && ncclParamCollNetEnable() == 1 && collNetSupport() && collNetGraph.nChannels) {