AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge remote-tracking branch 'rccl/develop' into 2.19.4

Browse Source
This commit is contained in:
Wenkai Du
2024-02-09 11:31:03 -06:00
parent e64324a64a 5669b0d7b6
commit d999d9ad21
13 changed files with 108 additions and 120 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/graph/paths.cc
+3 -4
View File
@@ -863,7 +863,8 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) {
} else {
// Round to next pow2 nChannelsPerPeer and nChannels
comm->p2pnChannelsPerPeer = (ncclParamNChannelsPerPeer() == -2 ? nextPow2(minChannels) : ncclParamNChannelsPerPeer());
comm->p2pnChannelsPerPeer *= 2;
// Doubling P2P channels per peer on single node
if (comm->topo->nodes[GPU].count == comm->topo->nRanks && IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94")) comm->p2pnChannelsPerPeer *= 2;
comm->p2pnChannels = nextPow2(comm->p2pnChannels);
}
@@ -874,9 +875,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) {
// fill the whole space of nChannels. To do so we mirror the bits in the
// nChannels space.
for (int c=0; c<comm->p2pnChannels; c++) {
int mirror = 0;
for (int b=1, mb=(comm->p2pnChannels>>1); b<comm->p2pnChannels; b<<=1, mb>>=1) if (c & b) mirror |= mb;
comm->p2pChannels[c] = mirror;
comm->p2pChannels[c] = mirrorBits(c, comm->p2pnChannels);
}
return ncclSuccess;
}
src/graph/search.cc
-8
View File
@@ -467,14 +467,6 @@ ncclResult_t ncclTopoSelectNets(struct ncclTopoSystem* system, int typeInter, in
for (int n=0; n<system->nodes[NET].count; n++) {
if (paths[n].type == t) localNets[localNetCount++] = n;
}
if (localNetCount == 0) continue;
// Shuffle by gpu NVML device number so that GPUs on the same PCI switch
// with multiple NICs don't use the same one as first choice.
for (int r=0; r<system->nodes[GPU].nodes[g].gpu.dev % localNetCount; r++) {
int net0 = localNets[0];
for (int i=0; i<localNetCount-1; i++) localNets[i] = localNets[i+1];
localNets[localNetCount-1] = net0;
}
// Append NICs to list
for (int i=0; i<localNetCount; i++) {
int n = localNets[i];
src/graph/topo.cc
+39 -98
View File
@@ -792,126 +792,67 @@ ncclResult_t ncclTopoGetSystem(struct ncclComm* comm, struct ncclTopoSystem** sy
return ncclSuccess;
}
static ncclResult_t getLocalNetMask(struct ncclTopoSystem* system, int g, uint64_t* localNetMask, int* type) {
ncclResult_t ncclTopoGetLocal(struct ncclTopoSystem* system, int type, int index, int resultType, int** locals, int* localCount, int* pathType) {
int minType = PATH_DIS;
float maxBw = 0;
int count = 0;
int* nets;
NCCLCHECK(ncclCalloc(&nets, system->nodes[NET].count));
for (int n=0; n<system->nodes[NET].count; n++) {
struct ncclTopoLinkList* path = system->nodes[NET].nodes[n].paths[GPU]+g;
if (path->bw > maxBw || (path->bw == maxBw && path->type < minType)) {
maxBw = path->bw;
minType = path->type;
if (type) *type = minType;
NCCLCHECK(ncclCalloc(locals, system->nodes[resultType].count));
struct ncclTopoLinkList* paths = system->nodes[type].nodes[index].paths[resultType];
for (int i=0; i<system->nodes[resultType].count; i++) {
if (paths[i].bw > maxBw || (paths[i].bw == maxBw && paths[i].type < minType)) {
maxBw = paths[i].bw;
minType = paths[i].type;
if (pathType) *pathType = minType;
count = 0;
}
if (path->bw == maxBw && path->type == minType) nets[count++] = system->nodes[NET].nodes[n].id;
if (paths[i].bw == maxBw && paths[i].type == minType) (*locals)[count++] = i;
}
*localNetMask = 0ULL;
for (int n=0; n<count; n++) {
if (nets[n] >= 64) return ncclInternalError;
*localNetMask |= 1ULL<<nets[n];
}
free(nets);
*localCount = count;
return ncclSuccess;
}
ncclResult_t ncclTopoGetLocalNet(struct ncclTopoSystem* system, int rank, int channelId, int* id) {
uint64_t* localNetMasks;
int ngpus = system->nodes[GPU].count;
NCCLCHECK(ncclCalloc(&localNetMasks, ngpus));
// Fill localNetMasks for all GPUs.
for (int g=0; g<ngpus; g++) {
NCCLCHECK(getLocalNetMask(system, g, localNetMasks+g, NULL));
}
// Find GPUs which have the same mask as rank, i.e. share the same local Nets.
int gpu;
NCCLCHECK(ncclTopoRankToIndex(system, rank, &gpu));
int netLocalGpus = 0, netLocalGpu = 0;
for (int g=0; g<ngpus; g++) {
if (localNetMasks[g] == localNetMasks[gpu]) {
if (g == gpu) netLocalGpu = netLocalGpus;
netLocalGpus++;
int* localNets;
int localNetCount;
NCCLCHECK(ncclTopoGetLocal(system, GPU, gpu, NET, &localNets, &localNetCount, NULL));
int* localGpus;
int localGpuCount;
NCCLCHECK(ncclTopoGetLocal(system, NET, localNets[0], GPU, &localGpus, &localGpuCount, NULL));
int net = 0;
for (int i = 0; i < localGpuCount; i++) {
if (gpu == localGpus[i]) {
net = i;
break;
}
}
uint64_t localNetMask = localNetMasks[gpu];
free(localNetMasks);
if (localNetMask == 0) return ncclInternalError;
// Round robin on GPUs and channels
int gIndex = 0, cId = 0, n = 0;
while (1) {
if (1ULL << n & localNetMask) {
if (gIndex == netLocalGpu && cId == channelId) {
*id = n;
return ncclSuccess;
}
gIndex++;
if (gIndex == netLocalGpus) {
gIndex = 0;
cId++;
}
}
n = (n+1) % 64;
if (isPow2(localNetCount)) net = mirrorBits(net, localNetCount);
if (localNetCount == 0) {
*id = -1;
} else {
net += channelId%(DIVUP(localNetCount,localGpuCount));
*id = system->nodes[NET].nodes[localNets[net%localNetCount]].id;
}
free(localNets);
free(localGpus);
return ncclSuccess;
}
ncclResult_t ncclTopoGetLocalGpu(struct ncclTopoSystem* system, int net, int* gpuIndex) {
int ngpus = system->nodes[GPU].count;
int* gpus;
NCCLCHECK(ncclCalloc(&gpus, ngpus));
// Find localNetMask which includes net with the most local GPUs.
int netLocalGpus = 0, minType = PATH_DIS;
uint64_t localNetMask = 0ULL;
for (int g=0; g<ngpus; g++) {
int type = PATH_DIS;
uint64_t mask;
NCCLCHECK(getLocalNetMask(system, g, &mask, &type));
if ((1ULL<<net) & mask) {
if (type < minType) {
localNetMask = mask;
netLocalGpus = 0;
minType = type;
}
if (type == minType) {
if (localNetMask && mask != localNetMask) {
WARN("Gpus %d and %d both have a type of %d with net %d yet have different netMasks of %lx and %lx\n", g, gpus[netLocalGpus-1], minType, net, mask, localNetMask);
free(gpus);
return ncclInternalError;
}
gpus[netLocalGpus] = g;
netLocalGpus++;
}
}
}
if (localNetMask == 0ULL) {
*gpuIndex = -1;
free(gpus);
return ncclSuccess;
}
// Round robin on GPUs and channels
int gIndex = 0, cId = 0, n = 0;
while (1) {
if (1ULL << n & localNetMask) {
if (n == net) {
*gpuIndex = gpus[gIndex];
free(gpus);
for (int c=0; c<MAXCHANNELS; c++) {
for (int g=0; g<system->nodes[GPU].count; g++) {
struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
int id;
NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &id));
if (net == id) {
*gpuIndex = g;
return ncclSuccess;
}
gIndex++;
if (gIndex == netLocalGpus) {
gIndex = 0;
cId++;
}
}
n = (n+1) % 64;
}
*gpuIndex = -1;
return ncclSuccess;
}
/****************************/
src/graph/topo.h
+9
View File
@@ -242,4 +242,13 @@ static float ncclTopoXGMISpeed(const char* gcn) {
#define ncclGetKernelIndex(p_comm) (0)
#endif
// Mirror bits
static bool isPow2(int val) {
return (val & (val-1)) == 0;
}
static int mirrorBits(int val, int pow2) {
int mirror = 0;
for (int b=1, mb=(pow2>>1); b<pow2; b<<=1, mb>>=1) if (val & b) mirror |= mb;
return mirror;
}
#endif
src/graph/xml.cc
+12 -1
View File
@@ -602,7 +602,18 @@ ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml*
}
}
pciNode->parent = parent;
parent->subs[parent->nSubs++] = pciNode;
// Keep PCI sub devices ordered by PCI Bus ID (Issue #820)
int subIndex = parent->nSubs;
const char* newBusId;
NCCLCHECK(xmlGetAttrStr(pciNode, "busid", &newBusId));
for (int s=0; s<parent->nSubs; s++) {
const char* busId;
NCCLCHECK(xmlGetAttrStr(parent->subs[s], "busid", &busId));
if (strcmp(newBusId, busId) < 0) { subIndex = s; break; }
}
for (int s = parent->nSubs; s > subIndex; s--) parent->subs[s] = parent->subs[s-1];
parent->subs[subIndex] = pciNode;
parent->nSubs++;
}
if (strcmp(parent->name, "pci") == 0) {
NCCLCHECK(ncclTopoGetXmlFromSys(parent, xml));