Merge remote-tracking branch 'nccl/v2.19' into develop
[ROCm/rccl commit: 81ddf9de89]
This commit is contained in:
BertanDogancay
@@ -52,10 +52,10 @@ ncclResult_t ncclTopoSearchInit(struct ncclTopoSystem* system) {
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
static ncclResult_t findRevLink(struct ncclTopoNode* node1, struct ncclTopoNode* node2, struct ncclTopoLink** revLink) {
|
||||
static ncclResult_t findRevLink(struct ncclTopoNode* node1, struct ncclTopoNode* node2, int type, struct ncclTopoLink** revLink) {
|
||||
for (int l=0; l<node2->nlinks; l++) {
|
||||
struct ncclTopoLink* link = node2->links+l;
|
||||
if (link->remNode == node1) {
|
||||
if (link->remNode == node1 && link->type == type) {
|
||||
*revLink = link;
|
||||
return ncclSuccess;
|
||||
}
|
||||
@@ -88,11 +88,11 @@ static ncclResult_t followPath(struct ncclTopoLinkList* path, struct ncclTopoNod
|
||||
float fwBw = link->type == LINK_PCI ? pciBw : bw;
|
||||
float revBw = 0;
|
||||
if (link->remNode->type == GPU && link->remNode->gpu.cudaCompCap < 80 && start->type != GPU) {
|
||||
if (revLink == NULL) NCCLCHECK(findRevLink(node, link->remNode, &revLink));
|
||||
if (revLink == NULL) NCCLCHECK(findRevLink(node, link->remNode, link->type, &revLink));
|
||||
revBw += fwBw/8;
|
||||
}
|
||||
if (link->remNode->type == CPU && link->type == LINK_NVL) {
|
||||
if (revLink == NULL) NCCLCHECK(findRevLink(node, link->remNode, &revLink));
|
||||
if (link->remNode->type == CPU && link->remNode->cpu.arch == NCCL_TOPO_CPU_ARCH_POWER && link->type == LINK_NVL) {
|
||||
if (revLink == NULL) NCCLCHECK(findRevLink(node, link->remNode, link->type, &revLink));
|
||||
revBw += fwBw;
|
||||
}
|
||||
if (link->bw < fwBw || (revBw && revLink->bw < revBw)) { *steps = step; return ncclSuccess; }
|
||||
@@ -263,6 +263,32 @@ ncclResult_t ncclTopoSearchNextGpuSort(struct ncclTopoSystem* system, struct ncc
|
||||
} else {
|
||||
for (int i=0; i<count; i++) next[i] = scores[i].g;
|
||||
}
|
||||
|
||||
if (system->nodes[NVS].count) {
|
||||
// NVSwitches prefer when we talk to a limited set of peers. Try to use neighbors first.
|
||||
int index = gpu-system->nodes[GPU].nodes;
|
||||
int i;
|
||||
int prevGpu = (index-1+ngpus)%ngpus;
|
||||
int nextGpu = (index+1)%ngpus;
|
||||
int firstGpus[2];
|
||||
int firstGpuCount = 0;
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_RING) {
|
||||
firstGpus[0] = nextGpu; firstGpus[1] = prevGpu; firstGpuCount = 2;
|
||||
} else if (graph->pattern == NCCL_TOPO_PATTERN_SPLIT_TREE ||
|
||||
graph->pattern == NCCL_TOPO_PATTERN_BALANCED_TREE) {
|
||||
firstGpus[0] = prevGpu; firstGpus[1] = nextGpu; firstGpuCount = 2;
|
||||
} else {
|
||||
firstGpus[0] = nextGpu; firstGpuCount = 1;
|
||||
}
|
||||
for (int g=0; g<firstGpuCount; g++) {
|
||||
for (i=0; i<count && next[i] != firstGpus[g]; i++);
|
||||
if (i<count) {
|
||||
for (; i>0; i--) next[i] = next[i-1];
|
||||
next[0] = firstGpus[g];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
*countPtr = count;
|
||||
return ncclSuccess;
|
||||
}
|
||||
@@ -270,7 +296,7 @@ ncclResult_t ncclTopoSearchNextGpuSort(struct ncclTopoSystem* system, struct ncc
|
||||
ncclResult_t ncclTopoSearchRec(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, struct ncclTopoGraph* saveGraph, int* time);
|
||||
|
||||
// Try to keep all searchs within one second
|
||||
#define NCCL_SEARCH_GLOBAL_TIMEOUT (5ULL<<16)
|
||||
#define NCCL_SEARCH_GLOBAL_TIMEOUT (1ULL<<19)
|
||||
#define NCCL_SEARCH_TIMEOUT (1<<14)
|
||||
#define NCCL_SEARCH_TIMEOUT_TREE (1<<14)
|
||||
#define NCCL_SEARCH_TIMEOUT_SAMECHANNELS (1<<8)
|
||||
@@ -374,6 +400,7 @@ ncclResult_t ncclTopoCompareGraphs(struct ncclTopoSystem* system, struct ncclTop
|
||||
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) { // NVLS channels correspond to GPUs pulling from NVLS. So the more the better.
|
||||
if (graph->nChannels > refGraph->nChannels && graph->nChannels <= system->nodes[GPU].count) *copy = 1;
|
||||
if (graph->nChannels*graph->bwInter > refGraph->nChannels*refGraph->bwInter) *copy = 1;
|
||||
return ncclSuccess;
|
||||
}
|
||||
// 2. Try to get better bandwidth
|
||||
@@ -394,24 +421,43 @@ ncclResult_t ncclTopoCompareGraphs(struct ncclTopoSystem* system, struct ncclTop
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
// Build a list of the best NETs to try.
|
||||
// Build a sorted list of the NETs to try.
|
||||
//
|
||||
// "gpu" can be set to -1 to build a list suitable for all GPUs (search start) or to a given gpu
|
||||
// index when trying to get back to the NIC.
|
||||
//
|
||||
// The list is built the following way:
|
||||
// 1. Select NETs starting with those close to GPU(s), based on paths[n].type.
|
||||
// 2. For each GPU, once that list of NICs with a given distance is prepared, shuffle the list
|
||||
// based on the GPU NVML index so that e.g. GPU 1 chooses NIC 1 first instead of NIC 0 which
|
||||
// might have been choosen by GPU 0 (case with multiple independent communicators per node)
|
||||
// 3. Then add the NETs to the final list if they were not already added by another closer GPU.
|
||||
// 2. add other NETs satisfying typeInter but not already in the list.
|
||||
|
||||
ncclResult_t ncclTopoSelectNets(struct ncclTopoSystem* system, int typeInter, int gpu, int* nets, int* netCountRet) {
|
||||
int netCount = 0;
|
||||
int localNetCount;
|
||||
int* localNets;
|
||||
NCCLCHECK(ncclCalloc(&localNets, system->nodes[NET].count));
|
||||
NCCLCHECK(ncclCalloc(&localNets, MAXCHANNELS));
|
||||
|
||||
// First add the preferred NICs
|
||||
for (int g=0; g<system->nodes[GPU].count; g++) {
|
||||
if (gpu != -1 && gpu != g) continue;
|
||||
localNetCount = 0;
|
||||
struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
|
||||
for (int c = 0; c<MAXCHANNELS; c++) {
|
||||
int netId;
|
||||
NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &netId));
|
||||
NCCLCHECK(ncclTopoIdToIndex(system, NET, netId, localNets+localNetCount));
|
||||
if (localNetCount > 0 && localNets[localNetCount] == localNets[0]) break;
|
||||
localNetCount++;
|
||||
}
|
||||
// Append NICs to list
|
||||
for (int i=0; i<localNetCount; i++) {
|
||||
int n = localNets[i];
|
||||
int found = 0;
|
||||
while (nets[found] != n && found<netCount) found++;
|
||||
if (found == netCount) nets[netCount++] = n;
|
||||
}
|
||||
}
|
||||
|
||||
// Then add others satisfying typeInter
|
||||
for (int t=0; t <= typeInter; t++) {
|
||||
for (int g=0; g<system->nodes[GPU].count; g++) {
|
||||
if (gpu != -1 && gpu != g) continue;
|
||||
@@ -473,11 +519,11 @@ ncclResult_t ncclTopoSearchRecGpu(struct ncclTopoSystem* system, struct ncclTopo
|
||||
int startNetIndex;
|
||||
NCCLCHECK(getNetIndex(system, graph->inter[graph->nChannels*2], &startNetIndex));
|
||||
struct ncclTopoNode* startNet = system->nodes[NET].nodes+startNetIndex;
|
||||
int netcount;
|
||||
int netCount;
|
||||
int* nets;
|
||||
NCCLCHECK(ncclCalloc(&nets, system->nodes[NET].count));
|
||||
NCCLCHECK(ncclTopoSelectNets(system, graph->typeInter, g, nets, &netcount));
|
||||
for (int i=0; i<netcount; i++) {
|
||||
NCCLCHECK(ncclTopoSelectNets(system, graph->typeInter, g, nets, &netCount));
|
||||
for (int i=0; i<netCount; i++) {
|
||||
int n = nets[i];
|
||||
struct ncclTopoNode* net = system->nodes[NET].nodes+n;
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_TREE && net->id != startNet->id) continue; // Trees are symmetric
|
||||
@@ -545,12 +591,12 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
|
||||
const int bw = graph->bwInter;
|
||||
int* nets;
|
||||
NCCLCHECK(ncclCalloc(&nets, system->nodes[NET].count));
|
||||
int netcount;
|
||||
NCCLCHECK(ncclTopoSelectNets(system, graph->typeInter, -1, nets, &netcount));
|
||||
for (int i=0; i<netcount; i++) {
|
||||
int n = nets[i];
|
||||
int netCount;
|
||||
NCCLCHECK(ncclTopoSelectNets(system, graph->typeInter, -1, nets, &netCount));
|
||||
for (int i=0; i<netCount; i++) {
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS && i>0) continue;
|
||||
int n = nets[(graph->nChannels+i)%netCount];
|
||||
struct ncclTopoNode* net = system->nodes[NET].nodes+n;
|
||||
struct ncclTopoNode* gpu;
|
||||
if (graph->collNet && net->net.collSupport == 0) continue;
|
||||
if (net->net.bw < bw) continue;
|
||||
|
||||
@@ -564,12 +610,12 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
|
||||
}
|
||||
}
|
||||
|
||||
// NVLS needs to balance on all NICs
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) {
|
||||
if (graph->nChannels < netcount) {
|
||||
// NVLS search only tries to find NIC:GPU combinations to compute the heads.
|
||||
if (graph->nChannels < netCount) {
|
||||
int gpu;
|
||||
NCCLCHECK(ncclTopoGetLocalGpu(system, nets[graph->nChannels], &gpu));
|
||||
if (gpu != -1) NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, -1, -1, gpu));
|
||||
NCCLCHECK(ncclTopoGetLocalGpu(system, net->id, &gpu));
|
||||
if (gpu != -1) NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, NET, n, gpu));
|
||||
}
|
||||
} else {
|
||||
if (graph->nChannels > 0) {
|
||||
@@ -579,7 +625,7 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
|
||||
NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, FORCED_ORDER_REPLAY, time, NET, n, g));
|
||||
}
|
||||
if (graph->nChannels == 0 || graph->sameChannels == 0) {
|
||||
if (graph->nChannels == 0) {
|
||||
if (graph->nChannels == 0 && system->nodes[NVS].count == 0) {
|
||||
// Always try the PCI order first to set a reference, but don't count in the timeout nor let it run for long
|
||||
struct ncclTopoLinkList* paths = net->paths[GPU];
|
||||
int f = 0, f_gdr = 0;
|
||||
@@ -613,18 +659,10 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
|
||||
}
|
||||
}
|
||||
if (maxBw >= bw) {
|
||||
// In the first loop, avoid using GPUs in both directions between channels (one channel
|
||||
// sending from that GPU and one channel receiving to that GPU), since that usually leads
|
||||
// to lower BW.
|
||||
for (int tryGpuBidir=0; tryGpuBidir<2; tryGpuBidir++) {
|
||||
for (int g=0; g<system->nodes[GPU].count; g++) {
|
||||
if (paths[g].bw == maxBw && paths[g].count == minHops) {
|
||||
gpu = system->nodes[GPU].nodes+g;
|
||||
int gpuUsed = gpuPciBw(gpu) > 0 ? 0 : 1;
|
||||
if (tryGpuBidir == gpuUsed) {
|
||||
NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, NET, n, g));
|
||||
}
|
||||
}
|
||||
for (int i=0; i<system->nodes[GPU].count; i++) {
|
||||
int g = (graph->nChannels+i)%system->nodes[GPU].count;
|
||||
if (paths[g].bw == maxBw && paths[g].count == minHops) {
|
||||
NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, NET, n, g));
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -840,6 +878,21 @@ ncclResult_t ncclTopoGetXmlFromGraphs(int ngraphs, struct ncclTopoGraph** graphs
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclResult_t ncclTopoDupChannels(struct ncclTopoGraph* graph, int ccMin, int ngpus) {
|
||||
if (graph->nChannels == 0) return ncclSuccess;
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) return ncclSuccess;
|
||||
if (graph->bwIntra < 25.0) return ncclSuccess;
|
||||
if (ccMin > 80 && graph->bwIntra < 50.0 && graph->nChannels > 4) return ncclSuccess;
|
||||
|
||||
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));
|
||||
graph->bwIntra /= DIVUP(dupChannels, graph->nChannels);
|
||||
graph->bwInter /= DIVUP(dupChannels, graph->nChannels);
|
||||
graph->nChannels = dupChannels;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
|
||||
float speedArrayIntra[] = { 48.0, 24.0, 20.0, 18.0, 15.0, 12.0, 10.0, 9.0, 7.0, 6.0, 5.0, 4.0, 3.0, 2.4, 1.2, 0.24, 0.12 };
|
||||
float speedArrayInter[] = { 48.0, 24.0, 20.0, 18.0, 15.0, 12.0, 10.0, 9.0, 7.0, 6.0, 5.0, 4.0, 3.0, 2.4, 1.2, 0.24, 0.12 };
|
||||
@@ -851,8 +904,8 @@ float speedArrayInter[] = { 48.0, 30.0, 28.0, 24.0, 20.0, 18.0, 15.0, 12.0, 10.0
|
||||
#define NSPEEDSINTRA (sizeof(speedArrayIntra)/sizeof(float))
|
||||
#define NSPEEDSINTER (sizeof(speedArrayInter)/sizeof(float))
|
||||
|
||||
float sm90SpeedArrayIntra[] = { 60.0, 40.0, 30.0, 24.0, 20.0, 15.0, 12.0, 6.0, 3.0 };
|
||||
float sm90SpeedArrayInter[] = { 48.0, 45.0, 42.0, 40.0, 30.0, 24.0, 20.0, 17.5, 15.0, 12.0, 6.0, 3.0, 2.4, 1.2, 0.24, 0.12 };
|
||||
float sm90SpeedArrayIntra[] = { 60.0, 50.0, 40.0, 30.0, 24.0, 20.0, 15.0, 12.0, 6.0, 3.0 };
|
||||
float sm90SpeedArrayInter[] = { 48.0, 45.0, 42.0, 40.0, 30.0, 24.0, 22.0, 20.0, 17.5, 15.0, 12.0, 6.0, 3.0, 2.4, 1.2, 0.24, 0.12 };
|
||||
#define NSPEEDSINTRA_SM90 (sizeof(sm90SpeedArrayIntra)/sizeof(float))
|
||||
#define NSPEEDSINTER_SM90 (sizeof(sm90SpeedArrayInter)/sizeof(float))
|
||||
#endif
|
||||
@@ -861,14 +914,13 @@ RCCL_PARAM(ModelMatchingDisable, "MODEL_MATCHING_DISABLE", 0);
|
||||
|
||||
ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph) {
|
||||
int ngpus = system->nodes[GPU].count;
|
||||
graph->crossNic = ncclParamCrossNic();
|
||||
int crossNic = (system->nodes[NET].count > 1) && graph->crossNic &&
|
||||
int crossNic = (system->nodes[NET].count > 1) &&
|
||||
(graph->pattern == NCCL_TOPO_PATTERN_RING ||
|
||||
graph->pattern == NCCL_TOPO_PATTERN_BALANCED_TREE ||
|
||||
graph->pattern == NCCL_TOPO_PATTERN_SPLIT_TREE) ? 1 : 0;
|
||||
graph->pattern == NCCL_TOPO_PATTERN_BALANCED_TREE ||
|
||||
graph->pattern == NCCL_TOPO_PATTERN_SPLIT_TREE) ? ncclParamCrossNic() : 0;
|
||||
graph->crossNic = crossNic == 1 ? 1 : 0;
|
||||
graph->bwIntra = graph->bwInter = 0;
|
||||
graph->latencyInter = 0;
|
||||
if (graph->crossNic == 2) graph->crossNic = 0;
|
||||
graph->typeIntra = ngpus == 1 ? PATH_LOC : PATH_NVL;
|
||||
graph->typeInter = PATH_PIX;
|
||||
graph->nChannels = 0;
|
||||
@@ -877,7 +929,10 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph
|
||||
int trySameChannels = graph->pattern == NCCL_TOPO_PATTERN_NVLS ? 0 : 1;
|
||||
graph->sameChannels = trySameChannels;
|
||||
|
||||
char* str = getenv("NCCL_GRAPH_FILE");
|
||||
int cpuArch, cpuVendor, cpuModel;
|
||||
NCCLCHECK(ncclTopoCpuType(system, &cpuArch, &cpuVendor, &cpuModel));
|
||||
|
||||
const char* str = ncclGetEnv("NCCL_GRAPH_FILE");
|
||||
if (str) {
|
||||
INFO(NCCL_ENV, "NCCL_GRAPH_FILE set by environment to %s", str);
|
||||
struct ncclXml* xml;
|
||||
@@ -930,6 +985,8 @@ ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph
|
||||
int ccMin;
|
||||
NCCLCHECK(ncclTopoGetCompCap(system, &ccMin, NULL));
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS && (system->nodes[NVS].count == 0 || ccMin < 90)) return ncclSuccess;
|
||||
// NVLS search must have ngpus heads at most.
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) graph->maxChannels = system->nodes[GPU].count;
|
||||
|
||||
if (ngpus == 1) if (graph->pattern != NCCL_TOPO_PATTERN_RING) graph->pattern = NCCL_TOPO_PATTERN_TREE;
|
||||
|
||||
@@ -968,7 +1025,7 @@ search:
|
||||
|
||||
NCCLCHECK(ncclTopoSearchRec(system, &tmpGraph, graph, &time));
|
||||
#if 0
|
||||
printf("Pattern %d, crossNic %d, Bw %g/%g, type %d/%d, channels %d-%d sameChannels %d -> nChannels %dx%g/%g %s\n", tmpGraph.pattern, tmpGraph.crossNic, tmpGraph.bwInter, tmpGraph.bwIntra, tmpGraph.typeInter, tmpGraph.typeIntra, tmpGraph.minChannels, tmpGraph.maxChannels, tmpGraph.sameChannels, graph->nChannels, graph->bwInter, graph->bwIntra, time == 0 ? "TIMEOUT" : time == -1 ? "PERFECT" : "");
|
||||
printf("Id %d Pattern %d, crossNic %d, Bw %g/%g, type %d/%d, channels %d-%d sameChannels %d -> nChannels %dx%g/%g %s\n", tmpGraph.id, tmpGraph.pattern, tmpGraph.crossNic, tmpGraph.bwInter, tmpGraph.bwIntra, tmpGraph.typeInter, tmpGraph.typeIntra, tmpGraph.minChannels, tmpGraph.maxChannels, tmpGraph.sameChannels, graph->nChannels, graph->bwInter, graph->bwIntra, time == 0 ? "TIMEOUT" : time == -1 ? "PERFECT" : "");
|
||||
for (int c=0; c<graph->nChannels; c++) {
|
||||
printf("%2d : ", c);
|
||||
for (int g=0; g<ngpus; g++) {
|
||||
@@ -985,8 +1042,9 @@ search:
|
||||
if (pass == 1) {
|
||||
// First pass, we don't have a solution yet ; try other options
|
||||
|
||||
// Try having different channels
|
||||
if (tmpGraph.sameChannels == 1) {
|
||||
// Try having different channels (except when going through AMD CPUs)
|
||||
if (tmpGraph.sameChannels == 1 &&
|
||||
!(cpuArch == NCCL_TOPO_CPU_ARCH_X86 && cpuVendor == NCCL_TOPO_CPU_VENDOR_AMD && tmpGraph.typeIntra == PATH_SYS)) {
|
||||
tmpGraph.sameChannels = 0;
|
||||
goto search;
|
||||
}
|
||||
@@ -1016,12 +1074,12 @@ search:
|
||||
}
|
||||
tmpGraph.typeInter = PATH_PIX;
|
||||
|
||||
if (crossNic && tmpGraph.crossNic == 0) {
|
||||
if (crossNic == 2 && tmpGraph.crossNic == 0) {
|
||||
// Try again with crossNic if permitted
|
||||
tmpGraph.crossNic = crossNic;
|
||||
tmpGraph.crossNic = 1;
|
||||
goto search;
|
||||
}
|
||||
tmpGraph.crossNic = 0;
|
||||
tmpGraph.crossNic = crossNic == 1 ? 1 : 0;
|
||||
|
||||
// Decrease bw until we find a solution
|
||||
if ((speedIndex < nspeeds-1) && (graph->nChannels == 0 || (speedArray[speedIndex+1]/graph->bwInter > .49))) {
|
||||
@@ -1038,6 +1096,7 @@ done:
|
||||
// We have a solution. Start from that solution and move to pass 2.
|
||||
if (pass == 1) {
|
||||
time = -1;
|
||||
NCCLCHECK(ncclTopoDupChannels(graph, ccMin, ngpus));
|
||||
memcpy(&tmpGraph, graph, sizeof(tmpGraph));
|
||||
speedIndex = 0;
|
||||
while (speedArray[speedIndex] > graph->bwInter && speedIndex < nspeeds-1) speedIndex++;
|
||||
@@ -1046,13 +1105,22 @@ done:
|
||||
pass = 2;
|
||||
}
|
||||
|
||||
// 3. See if we can increase bwIntra for trees (2 nodes or collnet)
|
||||
if (pass == 2) {
|
||||
if (time != 0 && graph->pattern != NCCL_TOPO_PATTERN_RING &&
|
||||
tmpGraph.bwIntra == graph->bwIntra && tmpGraph.bwIntra < tmpGraph.bwInter*2 &&
|
||||
speedIndex > 0) {
|
||||
tmpGraph.bwIntra = speedArray[--speedIndex];
|
||||
goto search;
|
||||
// See if we can increase bw
|
||||
if (time != 0 && speedIndex > 0) {
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_RING) {
|
||||
// increase bw for Ring
|
||||
tmpGraph.bwIntra = tmpGraph.bwInter = speedArray[--speedIndex];
|
||||
goto search;
|
||||
} else if (graph->pattern == NCCL_TOPO_PATTERN_NVLS && tmpGraph.bwInter == graph->bwInter && tmpGraph.bwInter < tmpGraph.bwIntra*2) {
|
||||
tmpGraph.minChannels = tmpGraph.maxChannels = graph->nChannels;
|
||||
tmpGraph.bwInter = speedArray[--speedIndex];
|
||||
goto search;
|
||||
} else if (tmpGraph.bwIntra == graph->bwIntra && tmpGraph.bwIntra < tmpGraph.bwInter*2) {
|
||||
// increase bwIntra for trees (2 nodes or collnet)
|
||||
tmpGraph.bwIntra = speedArray[--speedIndex];
|
||||
goto search;
|
||||
}
|
||||
}
|
||||
time = -1;
|
||||
memcpy(&tmpGraph, graph, sizeof(tmpGraph));
|
||||
@@ -1066,18 +1134,6 @@ done:
|
||||
graph->typeIntra = graph->typeInter = PATH_SYS;
|
||||
graph->nChannels = 1;
|
||||
}
|
||||
|
||||
if (graph->nChannels == 0) return ncclSuccess;
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) return ncclSuccess;
|
||||
if (graph->bwIntra < 25.0) return ncclSuccess;
|
||||
if (ccMin > 80 && graph->bwIntra < 50.0 && graph->nChannels > 4) return ncclSuccess;
|
||||
|
||||
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));
|
||||
graph->bwIntra /= DIVUP(dupChannels, graph->nChannels);
|
||||
graph->bwInter /= DIVUP(dupChannels, graph->nChannels);
|
||||
graph->nChannels = dupChannels;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
@@ -1117,7 +1173,7 @@ ncclResult_t ncclTopoPrintGraph(struct ncclTopoSystem* system, struct ncclTopoGr
|
||||
}
|
||||
|
||||
ncclResult_t ncclTopoDumpGraphs(struct ncclTopoSystem* system, int ngraphs, struct ncclTopoGraph** graphs) {
|
||||
char* str = getenv("NCCL_GRAPH_DUMP_FILE");
|
||||
const char* str = ncclGetEnv("NCCL_GRAPH_DUMP_FILE");
|
||||
if (str) {
|
||||
INFO(NCCL_ENV, "NCCL_GRAPH_DUMP_FILE set by environment to %s", str);
|
||||
struct ncclXml* xml;
|
||||
|
||||
Reference in New Issue
Block a user