AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Add another Rome model (#553)

Просмотр исходного кода 
* Add another Rome model

* Add option to force enable intranet on single node

* Limit p2p channels to number of ranks

* Refine p2p channels handling
Этот коммит содержится в:
Wenkai Du
2022-05-31 11:31:30 -07:00
коммит произвёл GitHub
родитель a0a686e74c
Коммит ef499c4810
7 изменённых файлов: 160 добавлений и 5 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
src/graph/paths.cc
+7 -1
Просмотреть файл
@@ -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
src/graph/rome_models.cc
+2 -2
Просмотреть файл
@@ -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 = "",
};
src/graph/topo.h
+1
Просмотреть файл
@@ -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;
src/init.cc
+1 -1
Просмотреть файл
@@ -817,7 +817,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
if ((comm->topo->type & RCCL_TOPO_4P2H_ROME) && (comm->topo->type & RCCL_TOPO_GDR_ALL)) {
if (rcclParamP2pNetDisable() == 0) {
comm->p2pNet = 1;
if (!(comm->topo->type & RCCL_TOPO_FORCE_INTRA)) comm->p2pNet = 1;
INFO(NCCL_INIT, "RCCL enabled same node P2P over network");
}
else
tools/scripts/topo_val.sh
+1 -1
Просмотреть файл
@@ -21,7 +21,7 @@
DIR="$(cd -P "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
for i in {0..76}
for i in {0..78}
do
if [[ $i -eq 50 ]] || [[ $i -eq 51 ]]
then
tools/topo_expl/models/topo_8p1h_5.xml
+146
Просмотреть файл
@@ -0,0 +1,146 @@
<system version="2">
<cpu numaid="0" affinity="00000000,00000000,ffffffff,ffffffff,00000000,00000000,ffffffff,ffffffff" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="1">
<pci busid="0000:22:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:2e:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:30:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x1022" subsystem_device="0x14c7" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:32:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="0" sm="90" gcn="910" arch="38911" rank="0" gdr="1">
<xgmi target="0000:35:00.0" count="4" tclass="0x038000"/>
<xgmi target="0000:11:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:8e:00.0" count="2" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:33:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x1022" subsystem_device="0x14c7" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:35:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="1" sm="90" gcn="910" arch="38911" rank="1" gdr="1">
<xgmi target="0000:32:00.0" count="4" tclass="0x038000"/>
<xgmi target="0000:14:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:b3:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:24:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:26:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x15b3" subsystem_device="0x0007" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_1" dev="0" speed="200000" port="1" latency="0.000000" guid="0x203d2b0003fd7010" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:2b:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:2d:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x15b3" subsystem_device="0x0007" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_2" dev="1" speed="200000" port="1" latency="0.000000" guid="0x503d2b0003fd7010" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
</pci>
<pci busid="0000:01:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:0d:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:0f:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x1022" subsystem_device="0x14c7" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:11:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="2" sm="90" gcn="910" arch="38911" rank="2" gdr="1">
<xgmi target="0000:32:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:14:00.0" count="4" tclass="0x038000"/>
<xgmi target="0000:ae:00.0" count="2" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:12:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x1022" subsystem_device="0x14c7" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:14:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="3" sm="90" gcn="910" arch="38911" rank="3" gdr="1">
<xgmi target="0000:35:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:11:00.0" count="4" tclass="0x038000"/>
<xgmi target="0000:93:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:03:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:05:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x15b3" subsystem_device="0x0007" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_3" dev="2" speed="200000" port="1" latency="0.000000" guid="0x143d2b0003fd7010" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:0a:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:0c:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x15b3" subsystem_device="0x0007" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_4" dev="3" speed="200000" port="1" latency="0.000000" guid="0x243d2b0003fd7010" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
</pci>
</cpu>
<cpu numaid="1" affinity="ffffffff,ffffffff,00000000,00000000,ffffffff,ffffffff,00000000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="1">
<pci busid="0000:a1:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:a9:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:ac:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x1022" subsystem_device="0x14c7" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:ae:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="4" sm="90" gcn="910" arch="38911" rank="4" gdr="1">
<xgmi target="0000:11:00.0" count="2" tclass="0x038000"/>
<xgmi target="0000:b3:00.0" count="4" tclass="0x038000"/>
<xgmi target="0000:8e:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:ab:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x15b3" subsystem_device="0x0007" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_6" dev="4" speed="200000" port="1" latency="0.000000" guid="0x883a2b0003fd7010" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:af:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:b1:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x1022" subsystem_device="0x14c7" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:b3:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="5" sm="90" gcn="910" arch="38911" rank="5" gdr="1">
<xgmi target="0000:35:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:ae:00.0" count="4" tclass="0x038000"/>
<xgmi target="0000:93:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:b4:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x15b3" subsystem_device="0x0007" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_7" dev="5" speed="200000" port="1" latency="0.000000" guid="0xf83c2b0003fd7010" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
</pci>
<pci busid="0000:81:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:89:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:8c:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x1022" subsystem_device="0x14c7" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:8e:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="6" sm="90" gcn="910" arch="38911" rank="6" gdr="1">
<xgmi target="0000:32:00.0" count="2" tclass="0x038000"/>
<xgmi target="0000:ae:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:93:00.0" count="4" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:8b:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x15b3" subsystem_device="0x0007" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_8" dev="6" speed="200000" port="1" latency="0.000000" guid="0x2c3d2b0003fd7010" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:8f:00.0" class="0x060400" vendor="0x1000" device="0xc010" subsystem_vendor="0x1000" subsystem_device="0xa096" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:91:00.0" class="0x060400" vendor="0x1022" device="0x14c7" subsystem_vendor="0x1022" subsystem_device="0x14c7" link_speed="16.0 GT/s PCIe" link_width="16">
<pci busid="0000:93:00.0" class="0x038000" vendor="0x1002" device="0x740c" subsystem_vendor="0x1002" subsystem_device="0x0b0c" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="7" sm="90" gcn="910" arch="38911" rank="7" gdr="1">
<xgmi target="0000:14:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:b3:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:8e:00.0" count="4" tclass="0x038000"/>
</gpu>
</pci>
</pci>
<pci busid="0000:94:00.0" class="0x020000" vendor="0x15b3" device="0x101b" subsystem_vendor="0x15b3" subsystem_device="0x0007" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_9" dev="7" speed="200000" port="1" latency="0.000000" guid="0x283d2b0003fd7010" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
</pci>
</cpu>
</system>
tools/topo_expl/topo_expl.cpp
+2
Просмотреть файл
@@ -147,6 +147,8 @@ NodeModelDesc model_descs[] = {
{4, "topo_8p1h_3.xml", "4 nodes 8P1H Alt."},
{1, "topo_8p1h_4.xml", "Single node 8P1H Alt."},
{2, "topo_8p1h_4.xml", "2 nodes 8P1H Alt."},
{1, "topo_8p1h_5.xml", "Single node 8P1H Alt."},
{2, "topo_8p1h_5.xml", "2 nodes 8P1H Alt."},
};
int main(int argc,char* argv[])