From 1fe031402a044e3cc4652d11d512015b5c426810 Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Wed, 14 Apr 2021 08:29:00 -0700 Subject: [PATCH] Add gfx90a target (#344) * Add gfx90a target * Support gfx90a topology Co-authored-by: Eiden Yoshida --- CMakeLists.txt | 31 +++++++----- cmake/Dependencies.cmake | 1 + src/graph/rome_models.cc | 12 +++++ src/graph/topo.cc | 3 +- src/graph/xml.cc | 2 + tools/scripts/topo_val.sh | 2 +- tools/topo_expl/models/topo_8p_90a.xml | 67 ++++++++++++++++++++++++++ tools/topo_expl/topo_expl.cpp | 1 + 8 files changed, 106 insertions(+), 13 deletions(-) create mode 100644 tools/topo_expl/models/topo_8p_90a.xml diff --git a/CMakeLists.txt b/CMakeLists.txt index 2656388a7d..84be06df7d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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 diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index 2eff262a74..b4ea4bd913 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -105,3 +105,4 @@ include(ROCMCreatePackage) include(ROCMInstallTargets) include(ROCMPackageConfigHelpers) include(ROCMInstallSymlinks) +include(ROCMCheckTargetIds OPTIONAL) diff --git a/src/graph/rome_models.cc b/src/graph/rome_models.cc index bb88a7b8b6..37e7354368 100755 --- a/src/graph/rome_models.cc +++ b/src/graph/rome_models.cc @@ -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 : diff --git a/src/graph/topo.cc b/src/graph/topo.cc index 61a38c1c84..d7ce17fd6f 100644 --- a/src/graph/topo.cc +++ b/src/graph/topo.cc @@ -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; snSubs; s++) { diff --git a/src/graph/xml.cc b/src/graph/xml.cc index 1e61332fda..e1a0b431d8 100644 --- a/src/graph/xml.cc +++ b/src/graph/xml.cc @@ -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. diff --git a/tools/scripts/topo_val.sh b/tools/scripts/topo_val.sh index 5d55a33fc5..5bd8e8efa0 100755 --- a/tools/scripts/topo_val.sh +++ b/tools/scripts/topo_val.sh @@ -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 diff --git a/tools/topo_expl/models/topo_8p_90a.xml b/tools/topo_expl/models/topo_8p_90a.xml new file mode 100644 index 0000000000..e81097f86a --- /dev/null +++ b/tools/topo_expl/models/topo_8p_90a.xml @@ -0,0 +1,67 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/tools/topo_expl/topo_expl.cpp b/tools/topo_expl/topo_expl.cpp index e4c284e41c..2c671e8a7b 100644 --- a/tools/topo_expl/topo_expl.cpp +++ b/tools/topo_expl/topo_expl.cpp @@ -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[])