From f22df90e5c6c7c589510b74d33e55bfc91cb4cd3 Mon Sep 17 00:00:00 2001 From: Nilesh M Negi Date: Thu, 26 Oct 2023 12:09:15 -0500 Subject: [PATCH] remove gcnArch support (#920) Signed-off-by: nileshnegi --- CMakeLists.txt | 11 ---------- src/graph/topo.cc | 20 +++++------------- src/graph/xml.cc | 15 ++++++++++---- src/include/archinfo.h | 2 +- src/misc/archinfo.cc | 47 ++++++++++++++++-------------------------- 5 files changed, 35 insertions(+), 60 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index efd0caba46..114ef58930 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -125,14 +125,6 @@ check_symbol_exists("hipEventDisableSystemFence" "hip/hip_runtime_api.h" HIP_EVE ### Check for hipDeviceMallocUncached support check_symbol_exists("hipDeviceMallocUncached" "hip/hip_runtime_api.h" HIP_UNCACHED_MEMORY) -message(STATUS "HIP Library version: ${hip_VERSION_MINOR}") -### Check the version of HIP to see if we can use gcnArchName instead of gcnArch (deprecated) -if(${hipcc_version_string} VERSION_LESS "5.7.31921") - set(HIP_NO_GCNARCHNAME ON) - else() - set(HIP_NO_GCNARCHNAME OFF) -endif() - ### Check for indirect function call support if(ENABLE_IFC) if(${hipcc_version_string} VERSION_GREATER_EQUAL "5.5.30201") @@ -566,9 +558,6 @@ else() target_compile_options(rccl PRIVATE --hipcc-func-supp) endif() endif() -if(HIP_NO_GCNARCHNAME) - target_compile_definitions(rccl PRIVATE HIP_NO_GCNARCHNAME) -endif() if (BUILD_BFD) if (HAVE_BFD) target_compile_definitions(rccl PRIVATE HAVE_BFD) diff --git a/src/graph/topo.cc b/src/graph/topo.cc index b7b820a1e0..588f9609b0 100644 --- a/src/graph/topo.cc +++ b/src/graph/topo.cc @@ -369,21 +369,11 @@ ncclResult_t ncclTopoAddNic(struct ncclXmlNode* xmlNic, struct ncclTopoSystem* s ncclResult_t ncclTopoAddGpu(struct ncclXmlNode* xmlGpu, struct ncclTopoSystem* system, struct ncclTopoNode* gpu) { NCCLCHECK(xmlGetAttrInt(xmlGpu, "sm", &gpu->gpu.cudaCompCap)); - NCCLCHECK(xmlGetAttr(xmlGpu, "gcn", &gpu->gpu.gcn)); - if (strcmp(gpu->gpu.gcn, "906") == 0) { - gpu->gpu.gcn = "gfx906"; - } else if (strcmp(gpu->gpu.gcn, "908") == 0) { - gpu->gpu.gcn = "gfx908"; - } else if (strcmp(gpu->gpu.gcn, "910") == 0) { - gpu->gpu.gcn = "gfx90a"; - } else if (strcmp(gpu->gpu.gcn, "940") == 0) { - gpu->gpu.gcn = "gfx940"; - } else if (strcmp(gpu->gpu.gcn, "941") == 0) { - gpu->gpu.gcn = "gfx941"; - } else if (strcmp(gpu->gpu.gcn, "942") == 0) { - gpu->gpu.gcn = "gfx942"; - } - + const char* gcnArch; + const char* gcnArchName; + NCCLCHECK(xmlGetAttr(xmlGpu, "gcn", &gcnArch)); + convertGcnArchToGcnArchName(gcnArch, gcnArchName); + gpu->gpu.gcn = strdup(gcnArchName); rcclHipDeviceArch_t arch; NCCLCHECK(xmlGetAttrInt(xmlGpu, "arch", &arch.value)); memcpy(&gpu->gpu.arch, &arch.arch, sizeof(hipDeviceArch_t)); diff --git a/src/graph/xml.cc b/src/graph/xml.cc index d11f80864a..e68aa092f0 100644 --- a/src/graph/xml.cc +++ b/src/graph/xml.cc @@ -636,15 +636,22 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev int sm; NCCLCHECK(xmlGetAttrInt(gpuNode, "sm", &sm)); - int gcn; + const char* gcn; + const char* gcnArchName; NCCLCHECK(xmlGetAttrIndex(gpuNode, "gcn", &index)); if (index == -1) { hipDeviceProp_t devProp; CUDACHECK(hipGetDeviceProperties(&devProp, 0)); - gcn = devProp.gcnArch; - NCCLCHECK(xmlSetAttrInt(gpuNode, "gcn", gcn)); + //extract only the releveant info from the gcnArchName attribute + //e.g.: convert "gfx908:sramecc+:xnack-" to "gfx908" + char gcnArchNameSubstr[6]; + GcnArchNameFormat(devProp.gcnArchName, gcnArchNameSubstr); + gcn = gcnArchNameSubstr; + NCCLCHECK(xmlSetAttr(gpuNode, "gcn", gcn)); } - NCCLCHECK(xmlGetAttrInt(gpuNode, "gcn", &gcn)); + NCCLCHECK(xmlGetAttr(gpuNode, "gcn", &gcn)); + convertGcnArchToGcnArchName(gcn, gcnArchName); + NCCLCHECK(xmlSetAttr(gpuNode, "gcn", gcnArchName)); rcclHipDeviceArch_t arch; NCCLCHECK(xmlGetAttrIndex(gpuNode, "arch", &index)); diff --git a/src/include/archinfo.h b/src/include/archinfo.h index a4a163bc3f..3fa4cb47bd 100644 --- a/src/include/archinfo.h +++ b/src/include/archinfo.h @@ -31,7 +31,7 @@ THE SOFTWARE. */ void GcnArchNameFormat(char *gcnArchName, char* out); -void GcnArchConvertToGcnArchName(int gcnArch, char* out); +void convertGcnArchToGcnArchName(const char* gcnArch, const char* gcnArchName); int GetGcnArchName(int deviceId, char* out); double GetDeviceWallClockRateInKhz(int deviceId); bool IsArchMatch(char const* arch, char const* target); diff --git a/src/misc/archinfo.cc b/src/misc/archinfo.cc index 7973cbd7be..96b3bd5f25 100644 --- a/src/misc/archinfo.cc +++ b/src/misc/archinfo.cc @@ -21,6 +21,7 @@ THE SOFTWARE. */ #include "archinfo.h" +#include "checks.h" #include #include @@ -31,44 +32,32 @@ void GcnArchNameFormat(char* gcnArchName, char* out) { strcpy(out, gcnArchNameToken); } -void GcnArchConvertToGcnArchName(int gcnArch, char* gcnArchName) { +void convertGcnArchToGcnArchName(const char* gcnArch, const char* gcnArchName) { // gcnArch is deprecated and we should instead use gcnArchName; however, some data files still have // the older gcnArch value. There's only a handful of architectures that were coded prior to deprecation, // so we handle those cases here. - //char gcnArchName[256] = {0}; // why 256? Because that's what gcnArchName gives us, so we're matching it. - gcnArchName[6] = 0; - switch (gcnArch) { - case 906: - strncpy(gcnArchName, "gfx906", 6); - break; - case 908: - strncpy(gcnArchName, "gfx908", 6); - break; - case 910: - // this is actually 90a - strncpy(gcnArchName, "gfx90a", 6); - break; - } + if (strcmp(gcnArch, "906") == 0) + gcnArchName = "gfx906"; + else if (strcmp(gcnArch, "908") == 0) + gcnArchName = "gfx908"; + else if (strcmp(gcnArch, "910") == 0) + gcnArchName = "gfx90a"; + else if (strcmp(gcnArch, "940") == 0) + gcnArchName = "gfx940"; + else if (strcmp(gcnArch, "941") == 0) + gcnArchName = "gfx941"; + else if (strcmp(gcnArch, "942") == 0) + gcnArchName = "gfx942"; + else + gcnArchName = gcnArch; } int GetGcnArchName(int deviceId, char* out) { - // this is a generic call in to get a consistent gcnArchName regardless of which version of rocm we're using. - // or which version of rocm we're using. + // this is a generic call in to get a consistent gcnArchName hipDeviceProp_t devProp; - hipError_t status = hipGetDeviceProperties(&devProp, deviceId); - if (status != hipSuccess) { - //std::cerr << "Encountered HIP error getting device properties: " - // << hipGetErrorString(status) << "\n"; - exit(-1); - } -#ifdef HIP_NO_GCNARCHNAME - // we're using a HIP version before 3.7. - GcnArchConvertToGcnArchName(devProp.gcnArch, out); - return 1; -#else + CUDACHECK(hipGetDeviceProperties(&devProp, deviceId)); GcnArchNameFormat(devProp.gcnArchName, out); return 0; -#endif } double GetDeviceWallClockRateInKhz(int deviceId) {