AI/rocm-systems
2
0 Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

Use parallel init of LDS and adjust P2P channels for gfx94x (#943)

Bladeren bron 
* Use parallel init of LDS and adjust P2P channels for gfx94x

* Move another init to parallel

* Fix NCCL_NCHANNELS_PER_PEER setting
This commit is contained in:
Wenkai Du authored and GitHub committed 2023-11-03 16:06:49 -07:00
1 bovenliggende d50bab28bf
commit dbcba2923b
2 files changed
+25 -22
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files

No files matched your search

src/graph/paths.cc
+4 -9
Bestand weergeven
@@ -761,12 +761,8 @@ static ncclResult_t ncclTopoGetNchannels(struct ncclTopoSystem* system, int g /*
// Local rank
path = system->nodes[GPU].nodes[peer].paths[GPU]+g;
if (path->type == PATH_NVL) {
if (ncclParamNChannelsPerPeer() == -2) {
float nvlBw = ncclTopoXGMISpeed(system->nodes[GPU].nodes[g].gpu.gcn);
*nChannels = 2*std::max(1, (int)(path->bw / nvlBw));
} else {
*nChannels = ncclParamNChannelsPerPeer();
}
float nvlBw = ncclTopoXGMISpeed(system->nodes[GPU].nodes[g].gpu.gcn);
*nChannels = (IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ? 4 : 2)*std::max(1, (int)(path->bw / nvlBw));
} else {
*nChannels = 2;
}
@@ -816,10 +812,9 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) {
// Adjust P2P channels on Rome
comm->p2pnChannelsPerPeer = 2;
comm->p2pnChannels = 2;
}
else {
} else {
// Round to next pow2 nChannelsPerPeer and nChannels
comm->p2pnChannelsPerPeer = nextPow2(minChannels);
comm->p2pnChannelsPerPeer = (ncclParamNChannelsPerPeer() == -2 ? nextPow2(minChannels) : ncclParamNChannelsPerPeer());
comm->p2pnChannels = nextPow2(comm->p2pnChannels);
}