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

topo_expl: 2.19 update

Browse Source
This commit is contained in:
Wenkai Du
2024-01-31 15:47:32 -06:00
parent 1a134b283b
commit d1575a1622
4 changed files with 480 additions and 62 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tools/topo_expl/models/topo_4p_940.xml
+58
View File
@@ -0,0 +1,58 @@
<system version="2">
<cpu numaid="2" affinity="000000ff,ffff0000,00000000,000000ff,ffff0000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:82:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="0" sm="94" gcn="940" arch="38911" rank="0" gdr="1">
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0000:81:00.0" class="0x020000" vendor="0x17db" device="0x0501" subsystem_vendor="0x17db" subsystem_device="0x0000" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="cxi0" dev="0" speed="200000000" port="1" latency="0.000000" guid="0x0" maxconn="128" gdr="1"/>
</nic>
</pci>
</cpu>
<cpu numaid="3" affinity="ffffff00,00000000,00000000,ffffff00,00000000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:c2:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="1" sm="94" gcn="940" arch="38911" rank="1" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0000:c1:00.0" class="0x020000" vendor="0x17db" device="0x0501" subsystem_vendor="0x17db" subsystem_device="0x0000" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="cxi1" dev="1" speed="200000000" port="1" latency="0.000000" guid="0x1" maxconn="128" gdr="1"/>
</nic>
</pci>
</cpu>
<cpu numaid="0" affinity="00000000,00000000,00ffffff,00000000,00000000,00ffffff" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:02:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="2" sm="94" gcn="940" arch="38911" rank="2" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0000:01:00.0" class="0x020000" vendor="0x17db" device="0x0501" subsystem_vendor="0x17db" subsystem_device="0x0000" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="cxi2" dev="2" speed="200000000" port="1" latency="0.000000" guid="0x2" maxconn="128" gdr="1"/>
</nic>
</pci>
</cpu>
<cpu numaid="1" affinity="00000000,0000ffff,ff000000,00000000,0000ffff,ff000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:42:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="3" sm="94" gcn="940" arch="38911" rank="3" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0000:41:00.0" class="0x020000" vendor="0x17db" device="0x0501" subsystem_vendor="0x17db" subsystem_device="0x0000" link_speed="16.0 GT/s PCIe" link_width="16">
<nic>
<net name="cxi3" dev="3" speed="200000000" port="1" latency="0.000000" guid="0x3" maxconn="128" gdr="1"/>
</nic>
</pci>
</cpu>
</system>
tools/topo_expl/models/topo_8p_940.xml
+182
View File
@@ -0,0 +1,182 @@
<system version="2">
<cpu numaid="0" affinity="00000000,000000ff,ffffffff,ffff0000,00000000,00ffffff,ffffffff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="143">
<pci busid="0000:05:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:08:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:0a:00.0" class="0x060400" vendor="0x1022" device="0x1500" subsystem_vendor="0x1022" subsystem_device="0x1500" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:0c:00.0" class="0x120000" vendor="0x1002" device="0x74a1" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="0" sm="94" gcn="gfx942" arch="38911" rank="0" gdr="1">
<xgmi target="0000:15:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:1e:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:69:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:9f:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:af:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:bf:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:df:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:07:00.0" class="0x020700" vendor="0x15b3" device="0x1021" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_0" dev="0" speed="400000" port="1" latency="0.000000" guid="0xf01cbc0103dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:0e:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:11:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:13:00.0" class="0x060400" vendor="0x1022" device="0x1500" subsystem_vendor="0x1022" subsystem_device="0x1500" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:15:00.0" class="0x120000" vendor="0x1002" device="0x74a1" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="1" sm="94" gcn="gfx942" arch="38911" rank="1" gdr="1">
<xgmi target="0000:0c:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:1e:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:69:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:9f:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:af:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:bf:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:df:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:10:00.0" class="0x020700" vendor="0x15b3" device="0x1021" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_1" dev="1" speed="400000" port="1" latency="0.000000" guid="0xf11cbc0203dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:17:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:1a:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:1c:00.0" class="0x060400" vendor="0x1022" device="0x1500" subsystem_vendor="0x1022" subsystem_device="0x1500" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:1e:00.0" class="0x120000" vendor="0x1002" device="0x74a1" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="2" sm="94" gcn="gfx942" arch="38911" rank="2" gdr="1">
<xgmi target="0000:0c:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:15:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:69:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:9f:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:af:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:bf:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:df:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:19:00.0" class="0x020700" vendor="0x15b3" device="0x1021" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_2" dev="2" speed="400000" port="1" latency="0.000000" guid="0xf21cbc0303dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:62:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:65:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:67:00.0" class="0x060400" vendor="0x1022" device="0x1500" subsystem_vendor="0x1022" subsystem_device="0x1500" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:69:00.0" class="0x120000" vendor="0x1002" device="0x74a1" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="3" sm="94" gcn="gfx942" arch="38911" rank="3" gdr="1">
<xgmi target="0000:0c:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:15:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:1e:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:9f:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:af:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:bf:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:df:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:64:00.0" class="0x020700" vendor="0x15b3" device="0x1021" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_3" dev="3" speed="400000" port="1" latency="0.000000" guid="0xf31cbc0403dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
</cpu>
<cpu numaid="1" affinity="ffffffff,ffffff00,00000000,0000ffff,ffffffff,ff000000,00000000" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="143">
<pci busid="0000:98:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:9b:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:9d:00.0" class="0x060400" vendor="0x1022" device="0x1500" subsystem_vendor="0x1022" subsystem_device="0x1500" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:9f:00.0" class="0x120000" vendor="0x1002" device="0x74a1" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="4" sm="94" gcn="gfx942" arch="38911" rank="4" gdr="1">
<xgmi target="0000:0c:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:15:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:1e:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:69:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:af:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:bf:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:df:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:9a:00.0" class="0x020700" vendor="0x15b3" device="0x1021" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_4" dev="4" speed="400000" port="1" latency="0.000000" guid="0xf41cbc0503dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:a8:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:ab:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:ad:00.0" class="0x060400" vendor="0x1022" device="0x1500" subsystem_vendor="0x1022" subsystem_device="0x1500" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:af:00.0" class="0x120000" vendor="0x1002" device="0x74a1" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="5" sm="94" gcn="gfx942" arch="38911" rank="5" gdr="1">
<xgmi target="0000:0c:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:15:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:1e:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:69:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:9f:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:bf:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:df:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:aa:00.0" class="0x020700" vendor="0x15b3" device="0x1021" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_5" dev="5" speed="400000" port="1" latency="0.000000" guid="0xf51cbc0603dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:b8:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:bb:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:bd:00.0" class="0x060400" vendor="0x1022" device="0x1500" subsystem_vendor="0x1022" subsystem_device="0x1500" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:bf:00.0" class="0x120000" vendor="0x1002" device="0x74a1" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="6" sm="94" gcn="gfx942" arch="38911" rank="6" gdr="1">
<xgmi target="0000:0c:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:15:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:1e:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:69:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:9f:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:af:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:df:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:ba:00.0" class="0x020700" vendor="0x15b3" device="0x1021" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_6" dev="6" speed="400000" port="1" latency="0.000000" guid="0xf61cbc0703dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
<pci busid="0000:d8:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:db:00.0" class="0x060400" vendor="0x15b3" device="0x1979" subsystem_vendor="0x0000" subsystem_device="0x0000" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:dd:00.0" class="0x060400" vendor="0x1022" device="0x1500" subsystem_vendor="0x1022" subsystem_device="0x1500" link_speed="32.0 GT/s PCIe" link_width="16">
<pci busid="0000:df:00.0" class="0x120000" vendor="0x1002" device="0x74a1" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="7" sm="94" gcn="gfx942" arch="38911" rank="7" gdr="1">
<xgmi target="0000:0c:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:15:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:1e:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:69:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:9f:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:af:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:bf:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
</pci>
</pci>
<pci busid="0000:da:00.0" class="0x020700" vendor="0x15b3" device="0x1021" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_7" dev="7" speed="400000" port="1" latency="0.000000" guid="0xf71cbc0803dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</pci>
</cpu>
</system>
tools/topo_expl/models/topo_8p_940vm.xml
+134
View File
@@ -0,0 +1,134 @@
<system version="2">
<cpu numaid="0" affinity="00000000,0000ffff,ffffffff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="143">
<pci busid="0021:00:00.0" class="0x120000" vendor="0x1002" device="0x74b5" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="1" sm="94" gcn="gfx942" arch="38911" rank="0" gdr="1">
<xgmi target="0011:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0031:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0041:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0051:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0061:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0071:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0081:00:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0011:00:00.0" class="0x120000" vendor="0x1002" device="0x74b5" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="0" sm="94" gcn="gfx942" arch="38911" rank="1" gdr="1">
<xgmi target="0021:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0031:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0041:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0051:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0061:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0071:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0081:00:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0031:00:00.0" class="0x120000" vendor="0x1002" device="0x74b5" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="2" sm="94" gcn="gfx942" arch="38911" rank="2" gdr="1">
<xgmi target="0011:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0021:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0041:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0051:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0061:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0071:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0081:00:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0041:00:00.0" class="0x120000" vendor="0x1002" device="0x74b5" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="3" sm="94" gcn="gfx942" arch="38911" rank="5" gdr="1">
<xgmi target="0011:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0021:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0031:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0051:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0061:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0071:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0081:00:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0020:00:00.0" class="0x020700" vendor="0x15b3" device="0x101e" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_0" dev="0" speed="400000" port="1" latency="0.000000" guid="0xd972b70203dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
<pci busid="0030:00:00.0" class="0x020700" vendor="0x15b3" device="0x101e" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_1" dev="1" speed="400000" port="1" latency="0.000000" guid="0xda72b70303dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
<pci busid="0010:00:00.0" class="0x020700" vendor="0x15b3" device="0x101e" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_2" dev="2" speed="400000" port="1" latency="0.000000" guid="0xd872b70103dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
<pci busid="0040:00:00.0" class="0x020700" vendor="0x15b3" device="0x101e" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_3" dev="3" speed="400000" port="1" latency="0.000000" guid="0xdb72b70403dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</cpu>
<cpu numaid="1" affinity="ffffffff,ffff0000,00000000" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="143">
<pci busid="0071:00:00.0" class="0x120000" vendor="0x1002" device="0x74b5" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="6" sm="94" gcn="gfx942" arch="38911" rank="3" gdr="1">
<xgmi target="0011:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0021:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0031:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0041:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0051:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0061:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0081:00:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0061:00:00.0" class="0x120000" vendor="0x1002" device="0x74b5" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="5" sm="94" gcn="gfx942" arch="38911" rank="4" gdr="1">
<xgmi target="0011:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0021:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0031:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0041:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0051:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0071:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0081:00:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0051:00:00.0" class="0x120000" vendor="0x1002" device="0x74b5" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="4" sm="94" gcn="gfx942" arch="38911" rank="6" gdr="1">
<xgmi target="0011:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0021:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0031:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0041:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0061:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0071:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0081:00:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0081:00:00.0" class="0x120000" vendor="0x1002" device="0x74b5" subsystem_vendor="0x1002" subsystem_device="0x74a1" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="7" sm="94" gcn="gfx942" arch="38911" rank="7" gdr="1">
<xgmi target="0011:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0021:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0031:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0041:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0051:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0061:00:00.0" count="1" tclass="0x120000"/>
<xgmi target="0071:00:00.0" count="1" tclass="0x120000"/>
</gpu>
</pci>
<pci busid="0080:00:00.0" class="0x020700" vendor="0x15b3" device="0x101e" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_4" dev="4" speed="400000" port="1" latency="0.000000" guid="0xdf72b70803dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
<pci busid="0050:00:00.0" class="0x020700" vendor="0x15b3" device="0x101e" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_5" dev="5" speed="400000" port="1" latency="0.000000" guid="0xdc72b70503dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
<pci busid="0060:00:00.0" class="0x020700" vendor="0x15b3" device="0x101e" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_6" dev="6" speed="400000" port="1" latency="0.000000" guid="0xdd72b70603dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
<pci busid="0070:00:00.0" class="0x020700" vendor="0x15b3" device="0x101e" subsystem_vendor="0x15b3" subsystem_device="0x0101" link_speed="32.0 GT/s PCIe" link_width="16">
<nic>
<net name="mlx5_7" dev="7" speed="400000" port="1" latency="0.000000" guid="0xde72b70703dd6838" maxconn="131072" gdr="1"/>
</nic>
</pci>
</cpu>
</system>
tools/topo_expl/utils.cpp
+106 -62
View File
@@ -272,9 +272,11 @@ ncclResult_t initChannel(struct ncclComm* comm, int channelId) {
/* channel->devPeers is not shared, so just free it when calling commFree() */
NCCLCHECK(ncclCudaCallocAsync(&channel->devPeers, nPeers, sharedRes->deviceStream.cudaStream));
ncclCommPushCudaFree(comm, channel->devPeers);
NCCLCHECK(ncclCalloc(&channel->devPeersHostPtr, nPeers));
for (int r = 0; r < nRanks; r++) {
uintptr_t addr = (uintptr_t)(comm->sharedRes->devPeers[channelId] + comm->topParentRanks[r]);
NCCLCHECK(ncclCudaMemcpyAsync((uintptr_t*)(channel->devPeers + r), (uintptr_t*)&addr, 1, sharedRes->deviceStream.cudaStream));
channel->devPeersHostPtr[r] = (struct ncclDevChannelPeer*)addr;
}
}
#endif
@@ -282,9 +284,9 @@ ncclResult_t initChannel(struct ncclComm* comm, int channelId) {
//NCCLCHECK(ncclCudaCallocAsync(&channel->devRingUserRanks, nRanks, sharedRes->deviceStream.cudaStream));
//ncclCommPushCudaFree(comm, channel->devRingUserRanks);
/* guarantee addr has been copied into channel->devPeers */
//NCCLCHECK(ncclStrongStreamSynchronize(&sharedRes->deviceStream));
//NCCLCHECK(ncclStrongStreamRelease(ncclCudaGraphNone(), &sharedRes->deviceStream));
//CUDACHECK(hipEventRecord(sharedRes->deviceStream.scratchEvent, sharedRes->deviceStream.cudaStream));
//CUDACHECK(hipStreamWaitEvent(sharedRes->deviceStream.cudaStream, sharedRes->deviceStream.scratchEvent, 0));
return ncclSuccess;
}
@@ -332,7 +334,7 @@ static ncclResult_t setupChannel(struct ncclComm* comm, int channelId, int rank,
}
template <int type>
static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclConnect* connect, int channelId, int peer, int connIndex, int* transportType) {
static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph* graph, struct ncclConnect* connect, int channelId, int peer, int connIndex, int* transportType, bool* needsProxy) {
struct ncclPeerInfo* myInfo = comm->peerInfo+comm->rank;
struct ncclPeerInfo* peerInfo = comm->peerInfo+peer;
struct ncclConnector* connector = (type == 1) ? comm->channels[channelId].peers[peer]->send + connIndex :
@@ -357,6 +359,7 @@ static ncclResult_t selectTransport(struct ncclComm* comm, struct ncclTopoGraph*
connector->transportComm = transportComm;
NCCLCHECK(transportComm->setup(comm, graph, myInfo, peerInfo, connect, connector, channelId, connIndex));
if (transportType) *transportType = t;
if (needsProxy) *needsProxy = (transportComm->proxyProgress != NULL);
return ncclSuccess;
}
}
@@ -390,14 +393,24 @@ void dumpData(struct ncclConnect* data, int ndata) {
}
}
NCCL_PARAM(ConnectRoundMaxPeers, "CONNECT_ROUND_MAX_PEERS", 128);
NCCL_PARAM(ReportConnectProgress, "REPORT_CONNECT_PROGRESS", 0);
#include <sys/time.h>
ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, int connIndex, int* highestTransportType/*=NULL*/, bool* needsProxy/*=NULL*/) {
// Stream used during transport setup; need for P2P pre-connect + CUDA Graph
ncclResult_t ret = ncclSuccess;
int highestType = TRANSPORT_P2P; // track highest transport type
struct ncclConnect** data = (ncclConnect**) malloc(sizeof(ncclConnect*) * comm->nRanks); // Store intermediate send/recvData structs for connect
struct ncclConnect** recvData = (ncclConnect**) malloc(sizeof(ncclConnect*) * comm->nRanks); // Points to entries inside data for given recv connection within a channel
struct ncclConnect** sendData = (ncclConnect**) malloc(sizeof(ncclConnect*) * comm->nRanks); // Points to entries inside data for given send connection within a channel
bool needsProxyResult = false;
struct ncclConnect** data; // Store intermediate send/recvData structs for connect
struct ncclConnect** recvData; // Points to entries inside data for given recv connection within a channel
struct ncclConnect** sendData; // Points to entries inside data for given send connection within a channel
int done = 0;
int maxPeers = ncclParamConnectRoundMaxPeers();
data = (ncclConnect**) malloc(sizeof(ncclConnect*) * maxPeers); // Store intermediate send/recvData structs for connect
recvData = (ncclConnect**) malloc(sizeof(ncclConnect*) * maxPeers); // Points to entries inside data for given recv connection within a channel
sendData = (ncclConnect**) malloc(sizeof(ncclConnect*) * maxPeers); // Points to entries inside data for given send connection within a channel
//NCCLCHECKGOTO(ncclStrongStreamAcquireUncaptured(&comm->sharedRes->hostStream), ret, fail);
// First time initialization
for (int i=1; i<comm->nRanks; i++) {
@@ -412,24 +425,27 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
// The first N entries contain recvData, connection information for recv connections
// The next M entries contain sendData, connection information for send connections
// It's not guaranteed that each entry of data has the same number of total or send/recv specific connections
data[i] = (ncclConnect*) malloc(sizeof(ncclConnect) * 2*MAXCHANNELS);
recvData[i] = data[i];
int p = i-(done+1);
if (recvMask || sendMask) data[p] = (ncclConnect*) malloc(sizeof(ncclConnect) * 2*MAXCHANNELS);
recvData[p] = data[p];
int sendChannels = 0, recvChannels = 0;
int type;
bool proxy;
TIME_START(0);
for (int c=0; c<MAXCHANNELS; c++) {
if (recvMask & (1UL<<c)) {
NCCLCHECKGOTO(selectTransport<0>(comm, graph, recvData[i]+recvChannels++, c, recvPeer, connIndex, &type), ret, fail);
NCCLCHECKGOTO(selectTransport<0>(comm, graph, recvData[p]+recvChannels++, c, recvPeer, connIndex, &type, &proxy), ret, fail);
if (type > highestType) highestType = type;
}
}
TIME_STOP(0);
TIME_START(1);
sendData[i] = recvData[i]+recvChannels;
sendData[p] = recvData[p]+recvChannels;
for (int c=0; c<MAXCHANNELS; c++) {
if (sendMask & (1UL<<c)) {
NCCLCHECKGOTO(selectTransport<1>(comm, graph, sendData[i]+sendChannels++, c, sendPeer, connIndex, &type), ret, fail);
NCCLCHECKGOTO(selectTransport<1>(comm, graph, sendData[p]+sendChannels++, c, sendPeer, connIndex, &type, &proxy), ret, fail);
if (type > highestType) highestType = type;
needsProxyResult |= proxy;
}
}
TIME_STOP(1);
@@ -439,8 +455,8 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
if (recvChannels+sendChannels) {
//NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, recvPeer, bootstrapTag, data[i], sizeof(struct ncclConnect)*(recvChannels+sendChannels)), ret, fail);
//NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, recvPeer, bootstrapTag, data[i], sizeof(struct ncclConnect)*(recvChannels+sendChannels)), ret, fail);
sendData[i] = data[i];
recvData[i] = data[i]+sendChannels;
sendData[p] = data[p];
recvData[p] = data[p]+sendChannels;
}
} else {
//if (recvChannels) NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, recvPeer, bootstrapTag, recvData[i], sizeof(struct ncclConnect)*recvChannels), ret, fail);
@@ -449,70 +465,97 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
//if (recvChannels) NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, recvPeer, bootstrapTag, recvData[i], sizeof(struct ncclConnect)*recvChannels), ret, fail);
}
TIME_STOP(2);
}
// Loop until all channels with all ranks have been connected
bool allChannelsConnected;
allChannelsConnected = false;
while (!allChannelsConnected) {
allChannelsConnected = true;
for (int i=1; i<comm->nRanks; i++) {
int recvPeer = (comm->rank - i + comm->nRanks) % comm->nRanks;
int sendPeer = (comm->rank + i) % comm->nRanks;
uint64_t recvMask = comm->connectRecv[recvPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)];
uint64_t sendMask = comm->connectSend[sendPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)];
if (i-done == maxPeers || i == comm->nRanks-1) {
// Loop until all channels with all ranks have been connected
bool allChannelsConnected;
allChannelsConnected = false;
while (!allChannelsConnected) {
allChannelsConnected = true;
for (int j=done+1; j<=i; j++) {
int recvPeer = (comm->rank - j + comm->nRanks) % comm->nRanks;
int sendPeer = (comm->rank + j) % comm->nRanks;
uint64_t recvMask = comm->connectRecv[recvPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)];
uint64_t sendMask = comm->connectSend[sendPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)];
int sendDataOffset = 0;
int recvDataOffset = 0;
for (int c=0; c<MAXCHANNELS; c++) {
TIME_START(3);
if (sendMask & (1UL<<c)) {
struct ncclConnector* conn = comm->channels[c].peers[sendPeer]->send + connIndex;
// This connector hasn't completed connection yet
if (conn->connected == 0) {
//NCCLCHECKGOTO(conn->transportComm->connect(comm, sendData[i] + sendDataOffset++, 1, comm->rank, conn), ret, fail);
if (ret == ncclSuccess) {
struct ncclDevChannelPeer* addr;
conn->connected = 1;
/* comm->channels[c].devPeers[sendPeer]->send[connIndex] is a device memory access. */
//CUDACHECKGOTO(cudaMemcpyAsync(&addr, &comm->channels[c].devPeers[sendPeer], sizeof(struct ncclDevChannelPeer*), cudaMemcpyDeviceToHost, comm->sharedRes->hostStream.cudaStream), ret, fail);
//CUDACHECKGOTO(cudaMemcpyAsync(&addr->send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), ret, fail);
} else if (ret == ncclInProgress) {
allChannelsConnected = false;
int p = j-(done+1);
int sendDataOffset = 0;
int recvDataOffset = 0;
for (int c=0; c<MAXCHANNELS; c++) {
TIME_START(3);
if (sendMask & (1UL<<c)) {
struct ncclConnector* conn = comm->channels[c].peers[sendPeer]->send + connIndex;
// This connector hasn't completed connection yet
if (conn->connected == 0) {
//NCCLCHECKGOTO(conn->transportComm->connect(comm, sendData[p] + sendDataOffset++, 1, comm->rank, conn), ret, fail);
if (ret == ncclSuccess) {
conn->connected = 1;
/* comm->channels[c].devPeers[sendPeer]->send[connIndex] is a device memory access. */
//CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeersHostPtr[sendPeer]->send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), ret, fail);
} else if (ret == ncclInProgress) {
allChannelsConnected = false;
}
}
}
}
TIME_STOP(3);
TIME_STOP(3);
// Start with recv channels
TIME_START(4);
if (recvMask & (1UL<<c)) {
struct ncclConnector* conn = comm->channels[c].peers[recvPeer]->recv + connIndex;
// This connector hasn't completed connection yet
if (conn->connected == 0) {
//NCCLCHECKGOTO(conn->transportComm->connect(comm, recvData[i] + recvDataOffset++, 1, comm->rank, conn), ret, fail);
if (ret == ncclSuccess) {
struct ncclDevChannelPeer* addr;
conn->connected = 1;
/* comm->channels[c].devPeers[recvPeer]->recv[connIndex] is a device memory access. */
//CUDACHECKGOTO(cudaMemcpyAsync(&addr, &comm->channels[c].devPeers[recvPeer], sizeof(struct ncclDevChannelPeer*), cudaMemcpyDeviceToHost, comm->sharedRes->hostStream.cudaStream), ret, fail);
//CUDACHECKGOTO(cudaMemcpyAsync(&addr->recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), ret, fail);
} else if (ret == ncclInProgress) {
allChannelsConnected = false;
// Start with recv channels
TIME_START(4);
if (recvMask & (1UL<<c)) {
struct ncclConnector* conn = comm->channels[c].peers[recvPeer]->recv + connIndex;
// This connector hasn't completed connection yet
if (conn->connected == 0) {
//NCCLCHECKGOTO(conn->transportComm->connect(comm, recvData[p] + recvDataOffset++, 1, comm->rank, conn), ret, fail);
if (ret == ncclSuccess) {
conn->connected = 1;
/* comm->channels[c].devPeers[recvPeer]->recv[connIndex] is a device memory access. */
//CUDACHECKGOTO(cudaMemcpyAsync(&comm->channels[c].devPeersHostPtr[recvPeer]->recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), cudaMemcpyHostToDevice, comm->sharedRes->hostStream.cudaStream), ret, fail);
} else if (ret == ncclInProgress) {
allChannelsConnected = false;
}
}
}
TIME_STOP(4);
}
TIME_STOP(4);
if (sendMask || recvMask) {
free(data[p]);
data[p] = NULL;
}
}
}
done = i;
}
}
// Clear all connect masks and free each connectInfo array
for (int i=1; i<comm->nRanks; i++) {
/* We need to sync ranks here since some ranks might run too fast after connection setup
* and start to destroy the connection after returning from this function; however, the
* others might still be trying to connect and import the buffer. No sync can lead to invalid
* shmem/cuda buffer. In addition, we also clear all connect masks and free each connectInfo array */
for (int i = 1; i < comm->nRanks; i++) {
int bootstrapTag = (i << 8) + (graph ? graph->id + 1 : 0);
int recvPeer = (comm->rank - i + comm->nRanks) % comm->nRanks;
int sendPeer = (comm->rank + i) % comm->nRanks;
int flag = 0;
/*if (recvPeer != sendPeer) {
if (comm->connectSend[sendPeer] != 0UL)
NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, sendPeer, bootstrapTag, &flag, sizeof(int)), ret, fail);
if (comm->connectRecv[recvPeer] != 0UL)
NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, recvPeer, bootstrapTag, &flag, sizeof(int)), ret, fail);
if (comm->connectSend[sendPeer] != 0UL)
NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, sendPeer, bootstrapTag, &flag, sizeof(int)), ret, fail);
if (comm->connectRecv[recvPeer] != 0UL)
NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, recvPeer, bootstrapTag, &flag, sizeof(int)), ret, fail);
} else {
if (comm->connectSend[sendPeer] != 0UL || comm->connectRecv[recvPeer] != 0UL) {
NCCLCHECKGOTO(bootstrapSend(comm->bootstrap, sendPeer, bootstrapTag, &flag, sizeof(int)), ret, fail);
NCCLCHECKGOTO(bootstrapRecv(comm->bootstrap, sendPeer, bootstrapTag, &flag, sizeof(int)), ret, fail);
}
}
comm->connectRecv[recvPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)] = comm->connectSend[sendPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)] = 0UL;
free(data[i]);
*/
}
free(data);
@@ -520,6 +563,7 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
free(recvData);
if (highestTransportType != NULL) *highestTransportType = highestType;
if (needsProxy != NULL) *needsProxy = needsProxyResult;
TIME_PRINT("P2P Setup/Connect");
exit:
//NCCLCHECK(ncclStrongStreamWaitStream(ncclCudaGraphNone(), &comm->sharedRes->deviceStream, &comm->sharedRes->hostStream));