diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index d870963101..17a42e2a16 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2590,6 +2590,7 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func); +#if 0 struct Agent_global { Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {} @@ -2629,6 +2630,7 @@ struct Agent_global { hipDeviceptr_t address; uint32_t byte_cnt; }; +#endif #if !__HIP_VDI__ #if defined(__cplusplus) @@ -2636,134 +2638,35 @@ struct Agent_global { #endif namespace hip_impl { -hsa_executable_t executable_for(hipModule_t); -const char* hash_for(hipModule_t); + class agent_globals_impl; + class agent_globals { + public: + agent_globals(); + ~agent_globals(); + agent_globals(const agent_globals&) = delete; -template -std::pair read_global_description( - ForwardIterator f, ForwardIterator l, const char* name) { - const auto it = std::find_if(f, l, [=](const Agent_global& x) { - return strcmp(x.name, name) == 0; - }); + hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name); + hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, + const char* name); + private: + agent_globals_impl* impl; + }; - return it == l ? - std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt); -} - -std::vector read_agent_globals(hsa_agent_t agent, - hsa_executable_t executable); -hsa_agent_t this_agent(); - - -class agent_globals_impl { -private: - std::pair< - std::mutex, - std::unordered_map< - std::string, std::vector>> globals_from_module; - - std::unordered_map< - hsa_agent_t, - std::pair< - std::once_flag, - std::vector>> globals_from_process; - -public: - - hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, - hipModule_t hmod, const char* name) { - // the key of the map would the hash of code object associated with the - // hipModule_t instance - std::string key(hash_for(hmod)); - - if (globals_from_module.second.count(key) == 0) { - std::lock_guard lck{globals_from_module.first}; - - if (globals_from_module.second.count(key) == 0) { - globals_from_module.second.emplace( - key, read_agent_globals(this_agent(), executable_for(hmod))); - } - } - - const auto it0 = globals_from_module.second.find(key); - if (it0 == globals_from_module.second.cend()) { - hip_throw( - std::runtime_error{"agent_globals data structure corrupted."}); - } - - std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), - it0->second.cend(), name); - - return *dptr ? hipSuccess : hipErrorNotFound; + inline + __attribute__((visibility("hidden"))) + agent_globals& get_agent_globals() { + static agent_globals ag; + return ag; } + extern "C" + inline + __attribute__((visibility("hidden"))) hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, - const char* name) { - - auto agent = this_agent(); - - std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) { - std::vector tmp0; - for (auto&& executable : hip_impl::get_program_state().executables(aa)) { - auto tmp1 = read_agent_globals(aa, executable); - tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()), - make_move_iterator(tmp1.end())); - } - globals_from_process[aa].second = move(move(tmp0)); - }, agent); - - const auto it = globals_from_process.find(agent); - - if (it == globals_from_process.cend()) return hipErrorNotInitialized; - - std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(), - it->second.second.cend(), name); - - return *dptr ? hipSuccess : hipErrorNotFound; + const char* name) { + return get_agent_globals().read_agent_global_from_process(dptr, bytes, name); } - -}; - -class agent_globals { -public: - agent_globals() : impl(new agent_globals_impl()) { - if (!impl) - hip_throw( - std::runtime_error{"Error when constructing agent global data structures."}); - } - ~agent_globals() { delete impl; } - - hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, - hipModule_t hmod, const char* name) { - return impl->read_agent_global_from_module(dptr, bytes, hmod, name); - } - - hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, - const char* name) { - return impl->read_agent_global_from_process(dptr, bytes, name); - } - -private: - agent_globals_impl* impl; -}; - -inline -__attribute__((visibility("hidden"))) -agent_globals& get_agent_globals() { - static agent_globals ag; - return ag; -} - - -extern "C" -inline -__attribute__((visibility("hidden"))) -hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, - const char* name) { - return get_agent_globals().read_agent_global_from_process(dptr, bytes, name); -} - - } // Namespace hip_impl. #if defined(__cplusplus) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index e8a8801e98..2b41988ac5 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -289,16 +289,6 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0)); } -hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, - hipModule_t hmod, const char* name) { - HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name); - if (!dptr || !bytes || !hmod) return hipErrorInvalidValue; - - if (!name) return hipErrorNotInitialized; - - return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name); -} - namespace hip_impl { hsa_executable_t executable_for(hipModule_t hmod) { return hmod->executable; @@ -323,10 +313,159 @@ namespace hip_impl { return currentDevice->_hsaAgent; } + + struct Agent_global { + Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {} + Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) + : name(nullptr), address(address), byte_cnt(byte_cnt) { + if (name) + this->name = strdup(name); + } + + Agent_global& operator=(Agent_global&& t) { + if (this == &t) return *this; + + if (name) free(name); + name = t.name; + address = t.address; + byte_cnt = t.byte_cnt; + + t.name = nullptr; + t.address = nullptr; + t.byte_cnt = 0; + + return *this; + } + + Agent_global(Agent_global&& t) + : name(nullptr), address(nullptr), byte_cnt(0) { + *this = std::move(t); + } + + // not needed, delete them to prevent bugs + Agent_global(const Agent_global&) = delete; + Agent_global& operator=(Agent_global& t) = delete; + + ~Agent_global() { if (name) free(name); } + + char* name; + hipDeviceptr_t address; + uint32_t byte_cnt; + }; + + template + std::pair read_global_description( + ForwardIterator f, ForwardIterator l, const char* name) { + const auto it = std::find_if(f, l, [=](const Agent_global& x) { + return strcmp(x.name, name) == 0; + }); + + return it == l ? + std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt); + } + + std::vector read_agent_globals(hsa_agent_t agent, + hsa_executable_t executable); + class agent_globals_impl { + private: + std::pair< + std::mutex, + std::unordered_map< + std::string, std::vector>> globals_from_module; + + std::unordered_map< + hsa_agent_t, + std::pair< + std::once_flag, + std::vector>> globals_from_process; + + public: + + hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name) { + // the key of the map would the hash of code object associated with the + // hipModule_t instance + std::string key(hash_for(hmod)); + + if (globals_from_module.second.count(key) == 0) { + std::lock_guard lck{globals_from_module.first}; + + if (globals_from_module.second.count(key) == 0) { + globals_from_module.second.emplace( + key, read_agent_globals(this_agent(), executable_for(hmod))); + } + } + + const auto it0 = globals_from_module.second.find(key); + if (it0 == globals_from_module.second.cend()) { + hip_throw( + std::runtime_error{"agent_globals data structure corrupted."}); + } + + std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), + it0->second.cend(), name); + + return *dptr ? hipSuccess : hipErrorNotFound; + } + + hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, + const char* name) { + + auto agent = this_agent(); + + std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) { + std::vector tmp0; + for (auto&& executable : hip_impl::get_program_state().executables(aa)) { + auto tmp1 = read_agent_globals(aa, executable); + tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()), + make_move_iterator(tmp1.end())); + } + globals_from_process[aa].second = move(move(tmp0)); + }, agent); + + const auto it = globals_from_process.find(agent); + + if (it == globals_from_process.cend()) return hipErrorNotInitialized; + + std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(), + it->second.second.cend(), name); + + return *dptr ? hipSuccess : hipErrorNotFound; + } + + }; + + agent_globals::agent_globals() : impl(new agent_globals_impl()) { + if (!impl) + hip_throw( + std::runtime_error{"Error when constructing agent global data structures."}); + } + agent_globals::~agent_globals() { delete impl; } + + hipError_t agent_globals::read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name) { + return impl->read_agent_global_from_module(dptr, bytes, hmod, name); + } + + hipError_t agent_globals::read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, + const char* name) { + return impl->read_agent_global_from_process(dptr, bytes, name); + } + } // Namespace hip_impl. +hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name) { + HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name); + if (!dptr || !bytes || !hmod) return hipErrorInvalidValue; + + if (!name) return hipErrorNotInitialized; + + return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name); +} + namespace { -inline void track(const Agent_global& x, hsa_agent_t agent) { +inline void track(const hip_impl::Agent_global& x, hsa_agent_t agent) { tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name, x.address, x.byte_cnt); @@ -347,7 +486,7 @@ inline void track(const Agent_global& x, hsa_agent_t agent) { } -template > +template > inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t agent, hsa_executable_symbol_t x, void* out) { using namespace hip_impl; @@ -358,7 +497,7 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t); if (t == HSA_SYMBOL_KIND_VARIABLE) { - Agent_global tmp(name(x).c_str(), address(x), size(x)); + hip_impl::Agent_global tmp(name(x).c_str(), address(x), size(x)); static_cast(out)->push_back(std::move(tmp)); track(static_cast(out)->back(),agent);