diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index 4aeb052364..42326cbdc1 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -162,7 +162,8 @@ void hipLaunchKernelGGLImpl( if (it == functions(agent).cend()) { hip_throw(std::runtime_error{ "No device code available for function: " + - name(function_address)}); + name(function_address) + + ", for agent: " + name(agent)}); } hipModuleLaunchKernel(it->second, numBlocks.x, numBlocks.y, numBlocks.z, diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 2b1d484178..7d9f08ac69 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2660,29 +2660,34 @@ __attribute__((visibility("hidden"))) hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, const char* name) { static std::unordered_map>> agent_globals; + std::vector>> globals; static std::once_flag f; auto agent = this_agent(); + // Create placeholder for each agent in the map. std::call_once(f, []() { - for (auto&& agent : hip_impl::all_hsa_agents()) { - agent_globals[agent].second.clear(); + for (auto&& x : hip_impl::all_hsa_agents()) { + (void)globals[x]; } }); - std::call_once(agent_globals[agent].first, [](hsa_agent_t agent) { + if (globals.find(agent) == globals.cend()) { + hip_throw(std::runtime_error{"invalid agent"}); + } + + std::call_once(globals[agent].first, [](hsa_agent_t aa) { std::vector tmp0; - for (auto&& executable : executables(agent)) { - auto tmp1 = read_agent_globals(agent, executable); + for (auto&& executable : executables(aa)) { + auto tmp1 = read_agent_globals(aa, executable); tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()), make_move_iterator(tmp1.end())); } - agent_globals[agent].second = move(tmp0); + globals[aa].second = move(tmp0); }, agent); - const auto it = agent_globals.find(agent); + const auto it = globals.find(agent); - if (it == agent_globals.cend()) return hipErrorNotInitialized; + if (it == globals.cend()) return hipErrorNotInitialized; std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(), it->second.second.cend(), name); diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index fea900d26f..f05f41d3f5 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -388,13 +388,17 @@ const std::vector& executables(hsa_agent_t agent) { // Create placeholder for each agent in the map. std::call_once(f, []() { - for (auto&& agent : hip_impl::all_hsa_agents()) { - r[agent].second.clear(); + for (auto&& x : hip_impl::all_hsa_agents()) { + (void)r[x]; } }); - std::call_once(r[agent].first, [](hsa_agent_t agent) { - hsa_agent_iterate_isas(agent, [](hsa_isa_t x, void* pa) { + if (r.find(agent) == r.cend()) { + hip_throw(std::runtime_error{"invalid agent"}); + } + + std::call_once(r[agent].first, [](hsa_agent_t aa) { + hsa_agent_iterate_isas(aa, [](hsa_isa_t x, void* pa) { const auto it = code_object_blobs().find(x); if (it == code_object_blobs().cend()) return HSA_STATUS_SUCCESS; @@ -419,7 +423,7 @@ const std::vector& executables(hsa_agent_t agent) { } return HSA_STATUS_SUCCESS; - }, &agent); + }, &aa); }, agent); return r[agent].second; @@ -490,22 +494,26 @@ const std::unordered_map< // Create placeholder for each agent in the map. std::call_once(f, []() { - for (auto&& agent : hip_impl::all_hsa_agents()) { - r[agent].second.clear(); + for (auto&& x : hip_impl::all_hsa_agents()) { + (void)r[x]; } }); - std::call_once(r[agent].first, [](hsa_agent_t agent) { + if (r.find(agent) == r.cend()) { + hip_throw(std::runtime_error{"invalid agent"}); + } + + std::call_once(r[agent].first, [](hsa_agent_t aa) { static const auto copy_kernels = []( - hsa_executable_t, hsa_agent_t agent, hsa_executable_symbol_t x, void*) { - if (type(x) == HSA_SYMBOL_KIND_KERNEL) r[agent].second[name(x)].push_back(x); + hsa_executable_t, hsa_agent_t a, hsa_executable_symbol_t x, void*) { + if (type(x) == HSA_SYMBOL_KIND_KERNEL) r[a].second[name(x)].push_back(x); return HSA_STATUS_SUCCESS; }; - for (auto&& executable : executables(agent)) { + for (auto&& executable : executables(aa)) { hsa_executable_iterate_agent_symbols( - executable, agent, copy_kernels, nullptr); + executable, aa, copy_kernels, nullptr); } }, agent); @@ -523,19 +531,23 @@ const std::unordered_map< // Create placeholder for each agent in the map. std::call_once(f, []() { - for (auto&& agent : hip_impl::all_hsa_agents()) { - r[agent].second.clear(); + for (auto&& x : hip_impl::all_hsa_agents()) { + (void)r[x]; } }); - std::call_once(r[agent].first, [](hsa_agent_t agent) { - for (auto&& function : function_names()) { - const auto it = kernels(agent).find(function.second); + if (r.find(agent) == r.cend()) { + hip_throw(std::runtime_error{"invalid agent"}); + } - if (it == kernels(agent).cend()) continue; + std::call_once(r[agent].first, [](hsa_agent_t aa) { + for (auto&& function : function_names()) { + const auto it = kernels(aa).find(function.second); + + if (it == kernels(aa).cend()) continue; for (auto&& kernel_symbol : it->second) { - r[agent].second.emplace( + r[aa].second.emplace( function.first, Kernel_descriptor{kernel_object(kernel_symbol), it->first}); }