AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Increase minimal channels for gfx908 (#259)

Sfoglia il codice sorgente
This commit is contained in:
Wenkai Du
2020-08-26 11:40:11 -07:00
committato da GitHub
parent b0919dc46c
commit c5cbece6d0
8 ha cambiato i file con 107 aggiunte e 7 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
tools/topo_expl/include/utils.h
+1
Vedi File
@@ -25,6 +25,7 @@ struct allGather3Data_t{
int cudaCompCap;
int fullCudaCompCap;
int nChannels;
int gcn;
struct ncclGraphInfo tree;
struct ncclGraphInfo ring;
struct ncclGraphInfo collNet;
tools/topo_expl/models/topo_4p3l_ia.xml
+85
Vedi File
@@ -0,0 +1,85 @@
<system version="2">
<cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
<gpu dev="0" sm="98" gcn="908" arch="38911" rank="0" gdr="0">
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
<gpu dev="1" sm="98" gcn="908" arch="38911" rank="1" gdr="0">
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
<gpu dev="2" sm="98" gcn="908" arch="38911" rank="2" gdr="0">
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
<gpu dev="3" sm="98" gcn="908" arch="38911" rank="3" gdr="0">
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:1a:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
<nic>
<net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
<gpu dev="4" sm="98" gcn="908" arch="38911" rank="4" gdr="0">
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
<gpu dev="5" sm="98" gcn="908" arch="38911" rank="5" gdr="0">
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
<gpu dev="6" sm="98" gcn="908" arch="38911" rank="6" gdr="0">
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
<pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
<gpu dev="7" sm="98" gcn="908" arch="38911" rank="7" gdr="0">
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
</pci>
</cpu>
</system>
tools/topo_expl/topo_expl.cpp
+2
Vedi File
@@ -103,6 +103,8 @@ NodeModelDesc model_descs[] = {
{4, "topo_8p_ts1_n4.xml", "4 nodes 8 VEGA20 TS1 NPS=4"},
{1, "topo_8p_ts1_n4_1.xml", "single node 8 VEGA20 TS1 NPS=4 Alt. Model"},
{4, "topo_8p_ts1_n4_1.xml", "4 nodes 8 VEGA20 TS1 NPS=4 Alt. Model"},
{1, "topo_4p3l_ia.xml", "single node 8 gfx908"},
{4, "topo_4p3l_ia.xml", "4 nodes 8 gfx908"},
};
int main(int argc,char* argv[])
tools/topo_expl/utils.cpp
+7 -2
Vedi File
@@ -216,7 +216,10 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t
}
// AllGather3 - begin
allGather3Data[rank].cudaCompCap = ncclCudaCompCap();
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 = std::min(treeGraph.nChannels, ringGraph.nChannels);
allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels;
allGather3Data[rank].tree.speedIntra = treeGraph.speedIntra;
@@ -397,8 +400,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;
for (int i=0; i<nranks; i++) {
allTopoRanks[i] = &allGather3Data[i].topoRanks;
gcn = std::min(allGather3Data[i].gcn, gcn);
// Make sure we align all ranks so that the tuning is consistent across ranks
treeGraph.nChannels = ringGraph.nChannels = comm->nChannels = std::min(allGather3Data[i].nChannels, comm->nChannels);
treeGraph.sameChannels = std::min(allGather3Data[i].tree.sameChannels, treeGraph.sameChannels);
@@ -424,7 +429,7 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
int *rings;
NCCLCHECK(ncclCalloc(&rings, nranks*MAXCHANNELS));
NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, allTopoRanks, rings));
NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, allTopoRanks, rings, gcn));
if (comm->nNodes > 1 &&
ncclParamCollNetEnable() == 1 &&
collNetSupport()) {