AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

improve program state commentary

Sfoglia il codice sorgente 
Disambiguate calling many varibles "agent".
More detail in exception message.
Create and discard map placeholders; no need to call std::vector::clear() on map value.
This commit is contained in:
Jeff Daily
2019-03-27 21:40:27 +00:00
parent c9117de8eb
commit 7aada87cbd
3 ha cambiato i file con 47 aggiunte e 29 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
include/hip/hcc_detail/functional_grid_launch.hpp
+2 -1
Vedi File
@@ -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,
include/hip/hcc_detail/hip_runtime_api.h
+14 -9
Vedi File
@@ -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<hsa_agent_t, std::pair<std::once_flag,
std::vector<Agent_global>>> agent_globals;
std::vector<Agent_global>>> 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<Agent_global> 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);
include/hip/hcc_detail/program_state.hpp
+31 -19
Vedi File
@@ -388,13 +388,17 @@ const std::vector<hsa_executable_t>& 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<hsa_executable_t>& 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});
}