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

Add gfx90a target (#344)

Browse Source 
* Add gfx90a target

* Support gfx90a topology

Co-authored-by: Eiden Yoshida <eiden.yoshida@amd.com>
This commit is contained in:
Wenkai Du
2021-04-14 08:29:00 -07:00
committed by GitHub
parent 3f18540f50
commit 1fe031402a
8 changed files with 106 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+20 -11
View File
@@ -11,27 +11,36 @@ set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "")
project(rccl CXX)
include(cmake/Dependencies.cmake)
# Detect compiler support for target ID
# This section is deprecated. Please use rocm_check_target_ids for future use.
if( CMAKE_CXX_COMPILER MATCHES ".*/hipcc$" )
execute_process(COMMAND ${CMAKE_CXX_COMPILER} "--help"
OUTPUT_VARIABLE CXX_OUTPUT
OUTPUT_STRIP_TRAILING_WHITESPACE
ERROR_STRIP_TRAILING_WHITESPACE)
string(REGEX MATCH ".mcode\-object\-version" TARGET_ID_SUPPORT ${CXX_OUTPUT})
execute_process(COMMAND ${CMAKE_CXX_COMPILER} "--help"
OUTPUT_VARIABLE CXX_OUTPUT
OUTPUT_STRIP_TRAILING_WHITESPACE
ERROR_STRIP_TRAILING_WHITESPACE)
string(REGEX MATCH ".mcode\-object\-version" TARGET_ID_SUPPORT ${CXX_OUTPUT})
endif()
# Use target ID syntax if supported for AMDGPU_TARGETS
if(TARGET_ID_SUPPORT)
set(AMDGPU_TARGETS gfx803;gfx900;gfx906:xnack-;gfx908:xnack- CACHE STRING "List of specific machine types for library to target")
#Set the AMDGPU_TARGETS with backward compatiblity
if(COMMAND rocm_check_target_ids)
rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx1030"
)
else()
set(AMDGPU_TARGETS gfx803;gfx900;gfx906;gfx908 CACHE STRING "List of specific machine types for library to target")
# Use target ID syntax if supported for AMDGPU_TARGETS
if(TARGET_ID_SUPPORT)
set(DEFAULT_AMDGPU_TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx1030")
else()
set(DEFAULT_AMDGPU_TARGETS "gfx803;gfx900;gfx906;gfx908")
endif()
endif()
set(AMDGPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "List of specific machine types for library to target")
option(BUILD_TESTS "Build test programs" OFF)
option(INSTALL_DEPENDENCIES "Force install dependencies" OFF)
include(cmake/Dependencies.cmake)
# parse version from Makefile NCCL_MAJOR, NCCL_MINOR, NCCL_PATCH must exist
# NCCL_SUFFIX is optional NCCL_VERSION formatting is ((X) * 1000 + (Y) * 100 +
# (Z)) so we must first detect one or two digits first
cmake/Dependencies.cmake
+1
View File
@@ -105,3 +105,4 @@ include(ROCMCreatePackage)
include(ROCMInstallTargets)
include(ROCMPackageConfigHelpers)
include(ROCMInstallSymlinks)
include(ROCMCheckTargetIds OPTIONAL)
src/graph/rome_models.cc
+12
View File
@@ -264,6 +264,17 @@ static struct rcclRomeModel rome_model_49 = {
.ringBase = "N0 0 1 2 3 4 5 6 7 N3|N3 7 6 5 4 3 2 1 0 N0|N1 2 3 0 1 6 7 4 5 N2|N2 5 4 7 6 1 0 3 2 N1",
};
static struct rcclRomeModel rome_model_52 = {
.nGpus = 8, .nCpus = 1, .nNics = 0, .nLinks = 3,
.gpuIds = { 0xc1000, 0xc5000, 0xc9000, 0xcd000, 0xd1000, 0xd5000, 0xd9000, 0xdd000, },
.nicIds = { },
.gpuNuma = { 0, 0, 0, 0, 0, 0, 0, 0, },
.nicNuma = { },
.connMatrix = { 0, 1, 1, 0, 0, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, 0, 1, 0, 0, 1, 1, 0, 0, 0, 0, 1, 1, 0, 0, 0, 0, 1, 0, 0, 1, 0, 0, 1, 1, 0, 0, 1, 0, 0, 1, 0, 0, 1, 1, 0, 0, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 1, 1, 0, },
.pattern = "80",
.ringBase = "0 1 3 2 4 5 7 6|6 7 5 4 2 3 1 0|0 1 5 4 6 7 3 2|2 3 7 6 4 5 1 0",
};
static struct rcclRomeModel romeTopoModels[] = {
rome_model_22,
rome_model_25,
@@ -285,6 +296,7 @@ static struct rcclRomeModel romeTopoModels[] = {
rome_model_46,
rome_model_48,
rome_model_49,
rome_model_52,
};
/* Parse user defined rings. Format is like :
src/graph/topo.cc
+2 -1
View File
@@ -415,7 +415,8 @@ ncclResult_t ncclTopoAddCpu(struct ncclXmlNode* xmlCpu, struct ncclTopoSystem* s
int familyId, modelId;
NCCLCHECK(xmlGetAttrInt(xmlCpu, "familyid", &familyId));
NCCLCHECK(xmlGetAttrInt(xmlCpu, "modelid", &modelId));
cpu->cpu.model = (familyId == 143 && modelId >= 49) ? NCCL_TOPO_CPU_TYPE_ROME : NCCL_TOPO_CPU_TYPE_ZEN;
// Treat "Milan" also as "Rome"
cpu->cpu.model = ((familyId == 143 && modelId >= 49) || familyId == 175) ? NCCL_TOPO_CPU_TYPE_ROME : NCCL_TOPO_CPU_TYPE_ZEN;
}
}
for (int s=0; s<xmlCpu->nSubs; s++) {
src/graph/xml.cc
+2
View File
@@ -519,6 +519,8 @@ ncclResult_t ncclTopoGetXmlFromSys(struct ncclXmlNode* pciNode, struct ncclXml*
// Save that for later in case next step is a CPU
char numaIdStr[MAX_STR_LEN];
NCCLCHECK(ncclTopoGetStrFromSys(path, "numa_node", numaIdStr));
// Workaround kernel bug for now
if (strcmp(numaIdStr, "-1") == 0) strcpy(numaIdStr, "0");
// Go up one level in the PCI tree. Rewind two "/" and follow the upper PCI
// switch, or stop if we reach a CPU root complex.
tools/scripts/topo_val.sh
+1 -1
View File
@@ -21,7 +21,7 @@
DIR="$(cd -P "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
for i in {0..51}
for i in {0..52}
do
if [[ $i -eq 50 ]] || [[ $i -eq 51 ]]
then
tools/topo_expl/models/topo_8p_90a.xml
+67
View File
@@ -0,0 +1,67 @@
<system version="2">
<cpu numaid="0" affinity="ffffffff,ffffffff,ffffffff,ffffffff" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="48">
<pci busid="0000:c1:00.0" class="0x038000" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="0" sm="90" gcn="910" arch="38911" rank="0" gdr="0">
<xgmi target="0000:c5:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:c9:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d9:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
<pci busid="0000:c5:00.0" class="0x038000" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="1" sm="90" gcn="910" arch="38911" rank="1" gdr="0">
<xgmi target="0000:c1:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:cd:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d5:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
<pci busid="0000:c9:00.0" class="0x038000" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="2" sm="90" gcn="910" arch="38911" rank="2" gdr="0">
<xgmi target="0000:c1:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:cd:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d1:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
<pci busid="0000:cd:00.0" class="0x038000" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="3" sm="90" gcn="910" arch="38911" rank="3" gdr="0">
<xgmi target="0000:c5:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:c9:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:dd:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
<pci busid="0000:d1:00.0" class="0x038000" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="4" sm="90" gcn="910" arch="38911" rank="4" gdr="0">
<xgmi target="0000:c9:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d5:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d9:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
<pci busid="0000:d5:00.0" class="0x038000" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="5" sm="90" gcn="910" arch="38911" rank="5" gdr="0">
<xgmi target="0000:c5:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d1:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:dd:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
<pci busid="0000:d9:00.0" class="0x038000" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="6" sm="90" gcn="910" arch="38911" rank="6" gdr="0">
<xgmi target="0000:c1:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d1:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:dd:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
<pci busid="0000:dd:00.0" class="0x038000" link_speed="16.0 GT/s PCIe" link_width="16">
<gpu dev="7" sm="90" gcn="910" arch="38911" rank="7" gdr="0">
<xgmi target="0000:cd:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d5:00.0" count="1" tclass="0x038000"/>
<xgmi target="0000:d9:00.0" count="1" tclass="0x038000"/>
</gpu>
</pci>
<pci busid="0000:91:00.0" class="0x060400" link_speed="16.0 GT/s PCIe" link_width="4">
<pci busid="0000:94:00.0" class="0x020000" link_speed="2.5 GT/s PCIe" link_width="1">
<nic>
<net name="eth0" dev="0" speed="1000" port="0" guid="0x0" maxconn="65536" gdr="0"/>
</nic>
</pci>
</pci>
</cpu>
</system>
tools/topo_expl/topo_expl.cpp
+1
View File
@@ -121,6 +121,7 @@ NodeModelDesc model_descs[] = {
{4, "topo_8p_rome_4nics.xml", "4 nodes node 8 gfx908 Rome 4 NICs"},
{4, "topo_collnet_n1.xml", "4 nodes collnet 1 NICs"},
{4, "topo_collnet_n4.xml", "4 nodes collnet 4 NICs"},
{1, "topo_8p_90a.xml", "single node gfx90a"},
};
int main(int argc,char* argv[])