diff --git a/src/comgr.cpp b/src/comgr.cpp index 25119d817e..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 == "-lclang_rt.builtins-x86_64") + // 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 f22a2347a7..f0a73afdb5 100644 --- a/src/hip/hip_build_utils.cpp +++ b/src/hip/hip_build_utils.cpp @@ -84,6 +84,45 @@ inline const std::string& GetCoV3Option(const bool enable) else return no_option; } + +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_1, 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; + auto next = std::string::npos; + if(end != std::string::npos) + { + next = str.find_first_not_of(" \t\n\r", end); + if(next != std::string::npos) + len = next - 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; + } + } +} + } // namespace static boost::filesystem::path HipBuildImpl(boost::optional& tmp_dir, @@ -136,7 +175,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, "clang_rt.builtins", "x86_64"); + params += hip_compiler_flags; + } + if(IsHccCompiler()) { env += std::string("KMOPTLLC=\"-mattr=+enable-ds128 "); 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