From 8c101d4169f446d8b754a7631b60b77124380199 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 7 Nov 2021 22:02:06 +0300 Subject: [PATCH 1/4] fix-swdev-310166(01) [HIP offline build] Remove x86 related lib from linking --- src/hip/hip_build_utils.cpp | 35 ++++++++++++++++++++++++++++++++++- 1 file changed, 34 insertions(+), 1 deletion(-) diff --git a/src/hip/hip_build_utils.cpp b/src/hip/hip_build_utils.cpp index f22a2347a7..2dab986ab6 100644 --- a/src/hip/hip_build_utils.cpp +++ b/src/hip/hip_build_utils.cpp @@ -84,6 +84,33 @@ inline const std::string& GetCoV3Option(const bool enable) else return no_option; } + +inline void InplaceRemoveAllTokensThatInclude(std::string& str, const std::string& search) +{ + for(std::size_t pos = 0;;) + { + if((pos = str.find(search, pos)) == std::string::npos) + break; + + auto beg = str.find_last_of(" \t\n\r", pos); + beg = (beg == std::string::npos) ? 0 : beg + 1; + + const auto end = str.find_first_of(" \t\n\r", pos); + + auto len = std::string::npos; + if(end != std::string::npos) + { + const auto next = str.find_first_not_of(" \t\n\r", end); + if(next != std::string::npos) + len = next - beg; + } + str.erase(beg, len); + if(len == std::string::npos) + break; + pos = beg; + } +} + } // namespace static boost::filesystem::path HipBuildImpl(boost::optional& tmp_dir, @@ -136,7 +163,13 @@ static boost::filesystem::path HipBuildImpl(boost::optional& tmp_dir, } params += " -Wno-unused-command-line-argument -I. "; - params += MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS); + + { + std::string hip_compiler_flags = MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS); + InplaceRemoveAllTokensThatInclude(hip_compiler_flags, "builtins-x86_64"); + params += hip_compiler_flags; + } + if(IsHccCompiler()) { env += std::string("KMOPTLLC=\"-mattr=+enable-ds128 "); From c9272dcc47255592771a90eed6296032db6e32ce Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 7 Nov 2021 22:03:50 +0300 Subject: [PATCH 2/4] fix-swdev-310166(02) [future ROCm] Disable -Wreserved-identifier for declarations of llvm intrinsics in HIP kernels. --- .../include/utility/static_kernel_amd_buffer_addressing.hpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/src/kernels/static_composable_kernel/include/utility/static_kernel_amd_buffer_addressing.hpp b/src/kernels/static_composable_kernel/include/utility/static_kernel_amd_buffer_addressing.hpp index 0bc3f1804d..2fd048530a 100644 --- a/src/kernels/static_composable_kernel/include/utility/static_kernel_amd_buffer_addressing.hpp +++ b/src/kernels/static_composable_kernel/include/utility/static_kernel_amd_buffer_addressing.hpp @@ -15,6 +15,9 @@ union BufferAddressConfig int32_t range[4]; }; +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wreserved-identifier" + __device__ float __llvm_amdgcn_buffer_load_f32(int32x4_t rsrc, index_t vindex, index_t offset, @@ -792,5 +795,7 @@ __device__ void amd_buffer_atomic_add(const float* p_src, } #endif // CK_USE_AMD_BUFFER_ATOMIC_FADD +#pragma clang diagnostic pop // "-Wreserved-identifier" + } // namespace ck #endif From 1f2741e43614fbcdda5e37bc2695cf6c7965de0c Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 7 Nov 2021 22:55:46 +0300 Subject: [PATCH 3/4] fix-swdev-310166(03) [COMGR] Remove x86 builtin lib from linking --- src/comgr.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/comgr.cpp b/src/comgr.cpp index 25119d817e..f00801905d 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -273,7 +273,7 @@ static void RemoveCommonOptionsUnwanted(OptionList& list) || (option == "-hc") || (option == "-x hip") || (option == "-xhip") || (option == "--hip-link") - || (option == "-lclang_rt.builtins-x86_64") + || (option.find("builtins-x86_64") != std::string::npos) || miopen::StartsWith(option, "-mllvm -amdgpu-early-inline-all") || miopen::StartsWith(option, "-mllvm -amdgpu-function-calls") || miopen::StartsWith(option, "--hip-device-lib-path="); // clang-format on From 5776cb81a58f4968bc04f79394656b51fe13f156 Mon Sep 17 00:00:00 2001 From: Artem Tamazov Date: Mon, 8 Nov 2021 21:37:45 +0300 Subject: [PATCH 4/4] fix-swdev-310166(04) Review comments --- src/comgr.cpp | 18 ++++++++++++------ src/hip/hip_build_utils.cpp | 36 ++++++++++++++++++++++++------------ 2 files changed, 36 insertions(+), 18 deletions(-) diff --git a/src/comgr.cpp b/src/comgr.cpp index f00801905d..f650906794 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -266,19 +266,25 @@ static bool IsLinkerOption(const std::string& option) static void RemoveCommonOptionsUnwanted(OptionList& list) { - list.erase(remove_if(list.begin(), - list.end(), - [&](const auto& option) { // clang-format off + list.erase( + remove_if( + list.begin(), + list.end(), + [&](const auto& option) { // clang-format off return miopen::StartsWith(option, "-mcpu=") || (option == "-hc") || (option == "-x hip") || (option == "-xhip") || (option == "--hip-link") - || (option.find("builtins-x86_64") != std::string::npos) + // The following matches current "-lclang_rt.builtins-x86_64" (4.5) as weel as + // upcoming ".../libclang_rt.builtins-x86_64.a" and even future things like + // "...x86_64.../libclang_rt.builtins.a" etc. + || ((option.find("clang_rt.builtins") != std::string::npos) + && (option.find("x86_64") != std::string::npos)) || miopen::StartsWith(option, "-mllvm -amdgpu-early-inline-all") || miopen::StartsWith(option, "-mllvm -amdgpu-function-calls") || miopen::StartsWith(option, "--hip-device-lib-path="); // clang-format on - }), - list.end()); + }), + list.end()); } static void RemoveCompilerOptionsUnwanted(OptionList& list) diff --git a/src/hip/hip_build_utils.cpp b/src/hip/hip_build_utils.cpp index 2dab986ab6..f0a73afdb5 100644 --- a/src/hip/hip_build_utils.cpp +++ b/src/hip/hip_build_utils.cpp @@ -85,29 +85,41 @@ inline const std::string& GetCoV3Option(const bool enable) return no_option; } -inline void InplaceRemoveAllTokensThatInclude(std::string& str, const std::string& search) +inline void InplaceRemoveAllTokensThatInclude(std::string& str, + const std::string& search_1, + const std::string& search_2) { for(std::size_t pos = 0;;) { - if((pos = str.find(search, pos)) == std::string::npos) + if((pos = str.find(search_1, pos)) == std::string::npos) break; - auto beg = str.find_last_of(" \t\n\r", pos); - beg = (beg == std::string::npos) ? 0 : beg + 1; - + auto beg = str.find_last_of(" \t\n\r", pos); + beg = (beg == std::string::npos) ? 0 : beg + 1; const auto end = str.find_first_of(" \t\n\r", pos); - auto len = std::string::npos; + auto len = std::string::npos; + auto next = std::string::npos; if(end != std::string::npos) { - const auto next = str.find_first_not_of(" \t\n\r", end); + next = str.find_first_not_of(" \t\n\r", end); if(next != std::string::npos) len = next - beg; } - str.erase(beg, len); - if(len == std::string::npos) - break; - pos = beg; + + const auto tok = + end == std::string::npos ? std::string{str, beg} : std::string{str, beg, end - beg}; + if(tok.find(search_2) != std::string::npos) + { + str.erase(beg, len); + if(len == std::string::npos) + break; + pos = beg; + } + else + { + pos = next; + } } } @@ -166,7 +178,7 @@ static boost::filesystem::path HipBuildImpl(boost::optional& tmp_dir, { std::string hip_compiler_flags = MIOPEN_STRINGIZE(HIP_COMPILER_FLAGS); - InplaceRemoveAllTokensThatInclude(hip_compiler_flags, "builtins-x86_64"); + InplaceRemoveAllTokensThatInclude(hip_compiler_flags, "clang_rt.builtins", "x86_64"); params += hip_compiler_flags; }