Hook into the creaky lazy-reinit machinery. Try to minimise race-risk.

include/hip/hcc_detail/functional_grid_launch.hpp

@@ -103,18 +103,23 @@ inline std::vector<std::uint8_t> make_kernarg(
 
     if (sizeof...(Formals) == 0) return {};
 
-    const auto it = function_names().find(
-        reinterpret_cast<std::uintptr_t>(kernel));
-
+    auto it = function_names().find(reinterpret_cast<std::uintptr_t>(kernel));
     if (it == function_names().cend()) {
-        throw std::runtime_error{"Undefined __global__ function."};
+        it =
+            function_names(true).find(reinterpret_cast<std::uintptr_t>(kernel));
+        if (it == function_names().cend()) {
+            throw std::runtime_error{"Undefined __global__ function."};
+        }
     }
 
-    const auto it1 = kernargs().find(it->second);
-
+    auto it1 = kernargs().find(it->second);
     if (it1 == kernargs().end()) {
-        throw std::runtime_error{
-            "Missing metadata for __global__ function: " + it->second};
+        it1 = kernargs(true).find(it->second);
+
+        if (it1 == kernargs().end()) {
+            throw std::runtime_error{
+                "Missing metadata for __global__ function: " + it->second};
+        }
     }
 
     std::tuple<Formals...> to_formals{std::move(actuals)};

include/hip/hcc_detail/program_state.hpp

@@ -99,8 +99,9 @@ const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kern
 functions(bool rebuild = false);
 const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
 std::unordered_map<std::string, void*>& globals(bool rebuild = false);
-std::unordered_map<
-    std::string, std::vector<std::pair<std::size_t, std::size_t>>>& kernargs();
+const std::unordered_map<
+    std::string, std::vector<std::pair<std::size_t, std::size_t>>>&
+    kernargs(bool rebuild = false);
 
 hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
                                  hsa_agent_t agent);