Fix HIP builds for the upcoming HIP compiler (SWDEV-310166)

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)

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<TmpDir>& tmp_dir,
@@ -136,7 +175,13 @@ static boost::filesystem::path HipBuildImpl(boost::optional<TmpDir>& 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 ");

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<float, 4>(const float* p_src,
 }
 #endif // CK_USE_AMD_BUFFER_ATOMIC_FADD
 
+#pragma clang diagnostic pop // "-Wreserved-identifier"
+
 } // namespace ck
 #endif