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

Add another Rome model (#553)

Browse Source 
* Add another Rome model

* Add option to force enable intranet on single node

* Limit p2p channels to number of ranks

* Refine p2p channels handling

[ROCm/rccl commit: ef499c4810]
This commit is contained in:
Wenkai Du
2022-05-31 11:31:30 -07:00
committed by GitHub
parent a03ab8e752
commit 5becf1669f
7 changed files with 160 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/graph/paths.cc
+7 -1
View File
@@ -576,6 +576,8 @@ ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclPeer
return ncclSuccess;
}
RCCL_PARAM(EnableIntranet, "ENABLE_INTRANET", 0);
ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm* comm) {
int *domains;
int64_t *ids;
@@ -669,6 +671,10 @@ ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm*
}
}
}
if (rcclParamEnableIntranet()) {
remove = 0;
system->type |= RCCL_TOPO_FORCE_INTRA;
}
comm->localRanks = system->nodes[GPU].count;
if (system->nodes[GPU].count == comm->nRanks && remove) {
for (int n=system->nodes[NET].count-1; n>=0; n--)
@@ -741,7 +747,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) {
else {
// Round to next pow2 nChannelsPerPeer and nChannels
comm->p2pnChannelsPerPeer = nextPow2(minChannels);
comm->p2pnChannels = nextPow2(comm->p2pnChannels);
comm->p2pnChannels = std::min(comm->topo->nodes[GPU].count == comm->topo->nRanks ? 2*comm->nRanks : comm->nRanks, nextPow2(comm->p2pnChannels));
}
// Init channels that weren't used so far
projects/rccl/src/graph/rome_models.cc
+2 -2
View File
@@ -573,8 +573,8 @@ static struct rcclRomeModel rome_model_76 = {
.connMatrix = { 0, 4, 1, 0, 0, 0, 2, 0, 4, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 4, 2, 0, 0, 0, 0, 1, 4, 0, 0, 0, 0, 1, 0, 0, 2, 0, 0, 4, 1, 0, 0, 1, 0, 0, 4, 0, 0, 1, 2, 0, 0, 0, 1, 0, 0, 4, 0, 0, 0, 1, 0, 1, 4, 0, },
.gdrLevel = { PATH_PXB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PXB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PXB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PXB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PXB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PXB, PATH_PXB, PATH_PHB, PATH_PHB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PXB, PATH_SYS, PATH_SYS, PATH_SYS, PATH_SYS, PATH_PHB, PATH_PHB, PATH_PXB, PATH_PXB, },
.pattern = "4444",
.ringBase = "N4 4 5 7 6 0 1 3 2 N2|N2 2 3 1 0 6 7 5 4 N4|N7 7 6 0 1 5 4 2 3 N3|N1 1 0 6 7 3 2 4 5 N5|N3 3 2 4 5 7 6 0 1 N1|N5 5 4 2 3 1 0 6 7 N7|N0 0 1 5 4 2 3 7 6 N6|N6 6 7 3 2 4 5 1 0 N0",
.options = "tuning=3",
.ringBase = "N0 0 1 3 2 4 5 7 6 N6|N2 2 3 1 0 6 7 5 4 N4|N5 5 4 2 3 7 6 0 1 N1|N1 1 0 6 7 3 2 4 5 N5|N4 4 5 7 6 0 1 3 2 N2|N2 2 3 1 0 6 7 5 4 N4|N0 0 1 5 4 2 3 7 6 N6|N3 3 2 4 5 1 0 6 7 N7|N4 4 5 7 6 0 1 3 2 N2|N6 6 7 5 4 2 3 1 0 N0|N7 7 6 0 1 5 4 2 3 N3|N6 6 7 3 2 4 5 1 0 N0|N3 3 2 0 1 5 4 6 7 N7|N1 1 0 2 3 7 6 4 5 N5|N5 5 4 6 7 3 2 0 1 N1|N7 7 6 4 5 1 0 2 3 N3",
.options = "disableNumaMatching=1,tuning=3",
.treeBase = "",
};
projects/rccl/src/graph/topo.h
+1
View File
@@ -88,6 +88,7 @@ struct ncclTopoLinkList {
#define RCCL_TOPO_4P2H_ROME 2
#define RCCL_TOPO_GDR_ALL 4
#define RCCL_TOPO_16P1H 8
#define RCCL_TOPO_FORCE_INTRA 16
struct ncclTopoNode {
int type;