From 9cee2c531144a97bbb502456aa786025eea63719 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 27 Mar 2019 18:19:10 +0000 Subject: [PATCH 1/4] load program state once per agent [ROCm/clr commit: 2845b4c4b8449f0803fac09c714cebbf52404fcd] --- .../hip/hcc_detail/functional_grid_launch.hpp | 23 +-- .../include/hip/hcc_detail/hip_runtime_api.h | 33 +++-- .../include/hip/hcc_detail/program_state.hpp | 135 ++++++++++-------- projects/clr/hipamd/src/hip_module.cpp | 15 +- 4 files changed, 102 insertions(+), 104 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp index 0e541001bf..4aeb052364 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp @@ -155,30 +155,17 @@ void hipLaunchKernelGGLImpl( std::uint32_t sharedMemBytes, hipStream_t stream, void** kernarg) { - auto it0 = functions().find(function_address); - if (it0 == functions().cend()) { + auto agent = target_agent(stream); + auto it = functions(agent).find(function_address); + + if (it == functions(agent).cend()) { hip_throw(std::runtime_error{ "No device code available for function: " + name(function_address)}); } - auto agent = target_agent(stream); - - const auto it1 = std::find_if( - it0->second.cbegin(), - it0->second.cend(), - [=](const std::pair& x) { - return x.first == agent; - }); - - if (it1 == it0->second.cend()) { - hip_throw(std::runtime_error{ - "No code available for function: " + name(function_address) + - ", for agent: " + name(agent)}); - } - - hipModuleLaunchKernel(it1->second, numBlocks.x, numBlocks.y, numBlocks.z, + hipModuleLaunchKernel(it->second, numBlocks.x, numBlocks.y, numBlocks.z, dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, stream, nullptr, kernarg); } diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index ba561472e9..2b1d484178 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -2659,30 +2659,33 @@ inline __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::vector> agent_globals; + static std::unordered_map>> agent_globals; static std::once_flag f; + auto agent = this_agent(); std::call_once(f, []() { - for (auto&& agent_executables : executables()) { - std::vector tmp0; - for (auto&& executable : agent_executables.second) { - auto tmp1 = read_agent_globals(agent_executables.first, - executable); - - tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()), - make_move_iterator(tmp1.end())); - } - agent_globals.emplace(agent_executables.first, move(tmp0)); + for (auto&& agent : hip_impl::all_hsa_agents()) { + agent_globals[agent].second.clear(); } }); - const auto it = agent_globals.find(this_agent()); + std::call_once(agent_globals[agent].first, [](hsa_agent_t agent) { + std::vector tmp0; + for (auto&& executable : executables(agent)) { + auto tmp1 = read_agent_globals(agent, executable); + tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()), + make_move_iterator(tmp1.end())); + } + agent_globals[agent].second = move(tmp0); + }, agent); + + const auto it = agent_globals.find(agent); if (it == agent_globals.cend()) return hipErrorNotInitialized; - std::tie(*dptr, *bytes) = read_global_description(it->second.cbegin(), - it->second.cend(), name); + std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(), + it->second.second.cend(), name); return *dptr ? hipSuccess : hipErrorNotFound; } diff --git a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp index f49ed44930..fea900d26f 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp @@ -83,6 +83,9 @@ inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) { } namespace hip_impl { + +std::vector all_hsa_agents(); + class Kernel_descriptor { std::uint64_t kernel_object_{}; amd_kernel_code_t const* kernel_header_{nullptr}; @@ -376,47 +379,50 @@ hsa_executable_t load_executable(const std::string& file, return executable; } -std::vector all_hsa_agents(); - inline __attribute__((visibility("hidden"))) -const std::unordered_map< - hsa_agent_t, std::vector>& executables() { - static std::unordered_map> r; +const std::vector& executables(hsa_agent_t agent) { + static std::unordered_map>> r; static std::once_flag f; + // Create placeholder for each agent in the map. std::call_once(f, []() { for (auto&& agent : hip_impl::all_hsa_agents()) { - hsa_agent_iterate_isas(agent, [](hsa_isa_t x, void* pa) { - const auto it = code_object_blobs().find(x); - - if (it == code_object_blobs().cend()) return HSA_STATUS_SUCCESS; - - hsa_agent_t a = *static_cast(pa); - - for (auto&& blob : it->second) { - hsa_executable_t tmp = {}; - - hsa_executable_create_alt( - HSA_PROFILE_FULL, - HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, - nullptr, - &tmp); - - // TODO: this is massively inefficient and only meant for - // illustration. - std::string blob_to_str{blob.cbegin(), blob.cend()}; - tmp = load_executable(blob_to_str, tmp, a); - - if (tmp.handle) r[a].push_back(tmp); - } - - return HSA_STATUS_SUCCESS; - }, &agent); + r[agent].second.clear(); } }); - return r; + std::call_once(r[agent].first, [](hsa_agent_t agent) { + hsa_agent_iterate_isas(agent, [](hsa_isa_t x, void* pa) { + const auto it = code_object_blobs().find(x); + + if (it == code_object_blobs().cend()) return HSA_STATUS_SUCCESS; + + hsa_agent_t a = *static_cast(pa); + + for (auto&& blob : it->second) { + hsa_executable_t tmp = {}; + + hsa_executable_create_alt( + HSA_PROFILE_FULL, + HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, + nullptr, + &tmp); + + // TODO: this is massively inefficient and only meant for + // illustration. + std::string blob_to_str{blob.cbegin(), blob.cend()}; + tmp = load_executable(blob_to_str, tmp, a); + + if (tmp.handle) r[a].second.push_back(tmp); + } + + return HSA_STATUS_SUCCESS; + }, &agent); + }, agent); + + return r[agent].second; } inline @@ -477,55 +483,66 @@ const std::unordered_map& function_names() { inline __attribute__((visibility("hidden"))) const std::unordered_map< - std::string, std::vector>& kernels() { - static std::unordered_map< - std::string, std::vector> r; + std::string, std::vector>& kernels(hsa_agent_t agent) { + static std::unordered_map>>> r; static std::once_flag f; + // Create placeholder for each agent in the map. std::call_once(f, []() { + for (auto&& agent : hip_impl::all_hsa_agents()) { + r[agent].second.clear(); + } + }); + + std::call_once(r[agent].first, [](hsa_agent_t agent) { static const auto copy_kernels = []( - hsa_executable_t, hsa_agent_t, hsa_executable_symbol_t x, void*) { - if (type(x) == HSA_SYMBOL_KIND_KERNEL) r[name(x)].push_back(x); + 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); return HSA_STATUS_SUCCESS; }; - for (auto&& agent_executables : executables()) { - for (auto&& executable : agent_executables.second) { - hsa_executable_iterate_agent_symbols( - executable, agent_executables.first, copy_kernels, nullptr); - } + for (auto&& executable : executables(agent)) { + hsa_executable_iterate_agent_symbols( + executable, agent, copy_kernels, nullptr); } - }); + }, agent); - return r; + return r[agent].second; } inline __attribute__((visibility("hidden"))) const std::unordered_map< std::uintptr_t, - std::vector>>& functions() { - static std::unordered_map< - std::uintptr_t, - std::vector>> r; + Kernel_descriptor>& functions(hsa_agent_t agent) { + static std::unordered_map>> r; static std::once_flag f; + // Create placeholder for each agent in the map. std::call_once(f, []() { - for (auto&& function : function_names()) { - const auto it = kernels().find(function.second); - - if (it == kernels().cend()) continue; - - for (auto&& kernel_symbol : it->second) { - r[function.first].emplace_back( - agent(kernel_symbol), - Kernel_descriptor{kernel_object(kernel_symbol), it->first}); - } + for (auto&& agent : hip_impl::all_hsa_agents()) { + r[agent].second.clear(); } }); - return r; + std::call_once(r[agent].first, [](hsa_agent_t agent) { + for (auto&& function : function_names()) { + const auto it = kernels(agent).find(function.second); + + if (it == kernels(agent).cend()) continue; + + for (auto&& kernel_symbol : it->second) { + r[agent].second.emplace( + function.first, + Kernel_descriptor{kernel_object(kernel_symbol), it->first}); + } + } + }, agent); + + return r[agent].second; } inline diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 994f211bb4..6401e13727 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -490,21 +490,12 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) if (!attr) return hipErrorInvalidValue; if (!func) return hipErrorInvalidDeviceFunction; - const auto it0 = functions().find(reinterpret_cast(func)); - - if (it0 == functions().cend()) return hipErrorInvalidDeviceFunction; - auto agent = this_agent(); - const auto it1 = find_if( - it0->second.cbegin(), - it0->second.cend(), - [=](const pair& x) { - return x.first == agent; - }); + const auto it = functions(agent).find(reinterpret_cast(func)); - if (it1 == it0->second.cend()) return hipErrorInvalidDeviceFunction; + if (it == functions(agent).cend()) return hipErrorInvalidDeviceFunction; - const auto header = static_cast(it1->second)->_header; + const auto header = static_cast(it->second)->_header; if (!header) throw runtime_error{"Ill-formed Kernel_descriptor."}; From 73bb9a74bb4e426b65a5bd1935a7203b3a963800 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 28 Mar 2019 02:21:45 +0530 Subject: [PATCH 2/4] Handle D2D in memcpy2D [ROCm/clr commit: 50d623981e25918caec6b78143af2ebd5616d9d9] --- projects/clr/hipamd/src/hip_memory.cpp | 22 ++++++++++++++-------- 1 file changed, 14 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index e3823c504e..8196be91ff 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1586,21 +1586,24 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue; hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - int isLocked = 0; + int isLockedOrD2D = 0; void *pinnedPtr=NULL; void *actualSrc = (void*)src; void *actualDest = dst; if(kind == hipMemcpyHostToDevice ) { if(getLockedPointer((void*)src, spitch, &pinnedPtr) == hipSuccess ){ - isLocked = 1; + isLockedOrD2D = 1; actualSrc = pinnedPtr; } } else if(kind == hipMemcpyDeviceToHost) { if(getLockedPointer((void*)dst, dpitch, &pinnedPtr) == hipSuccess ){ - isLocked = 1; + isLockedOrD2D = 1; actualDest = pinnedPtr; } + } else if(kind == hipMemcpyDeviceToDevice) { + isLockedOrD2D = 1; } + hc::completion_future marker; hipError_t e = hipSuccess; @@ -1608,7 +1611,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false); } else { try { - if(!isLocked) { + if(!isLockedOrD2D) { for (int i = 0; i < height; ++i) stream->locked_copySync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind); @@ -1639,27 +1642,30 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp HIP_INIT_SPECIAL_API(hipMemcpy2DAsync, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream); if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue); hipError_t e = hipSuccess; - int isLocked = 0; + int isLockedOrD2D = 0; void *pinnedPtr=NULL; void *actualSrc = (void*)src; void *actualDest = dst; stream = ihipSyncAndResolveStream(stream); if(kind == hipMemcpyHostToDevice ) { if(getLockedPointer((void*)src, spitch, &pinnedPtr) == hipSuccess ){ - isLocked = 1; + isLockedOrD2D = 1; actualSrc = pinnedPtr; } } else if(kind == hipMemcpyDeviceToHost) { if(getLockedPointer((void*)dst, dpitch, &pinnedPtr) == hipSuccess ){ - isLocked = 1; + isLockedOrD2D = 1; actualDest = pinnedPtr; } + } else if(kind == hipMemcpyDeviceToDevice) { + isLockedOrD2D = 1; } + if((width == dpitch) && (width == spitch)) { hip_internal::memcpyAsync(dst, src, width*height, kind, stream); } else { try { - if(!isLocked){ + if(!isLockedOrD2D){ for (int i = 0; i < height; ++i) e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind, stream); From 5233d41c6c72d3734bd0ab1eef84bf2296403cb6 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 27 Mar 2019 21:40:27 +0000 Subject: [PATCH 3/4] improve program state commentary 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. [ROCm/clr commit: f5e4fff6cc9513b93880a3cb9b365f8e24b3b62a] --- .../hip/hcc_detail/functional_grid_launch.hpp | 3 +- .../include/hip/hcc_detail/hip_runtime_api.h | 23 +++++---- .../include/hip/hcc_detail/program_state.hpp | 50 ++++++++++++------- 3 files changed, 47 insertions(+), 29 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp index 4aeb052364..42326cbdc1 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/clr/hipamd/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/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 2b1d484178..7d9f08ac69 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/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/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp index fea900d26f..f05f41d3f5 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp +++ b/projects/clr/hipamd/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}); } From cfe930f9d67ba4e11daa7fb751e3f30ab9922c9a Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Thu, 28 Mar 2019 22:45:04 -0500 Subject: [PATCH 4/4] Make hipModuleGetGlobal be in HIP runtime so it can be discovered at runtime (#981) * Make hipModuleGetGlobal be in HIP runtime so it can be discovered at runtime In HIP PR #929, quite a few HIP public APIs were made as inline functions with hidden visibility. It was necessary to support applications with shared libraries with GPU kernels launched via hipLaunchKernelGGL(), after HIP runtime is initialized. In empirical tests, the implementation has been proved to be a bit too excessive, especially for hipModuleGetGlobal(). The function is used by another type of client applications which relies on the existence of this function within HIP runtime so global symbols from HSA code objects loaded dynamically at runtime can be retrieved programmtically. This commit moves hipModuleGetGlobal() back to src/hip_module.cpp, and makes it visible and not inline, to fulfill requirements for applications aforementioned. It does not change the behavior of applications depending on hipLaunchKernelGGL(). * Add HIP_INIT_API into the implementation of hipModuleGetGlobal Address review comments. * Fix failing HIP unit tests [ROCm/clr commit: 04915cea2f3ee33a37ac545c340bb98fdb395408] --- .../include/hip/hcc_detail/hip_runtime_api.h | 25 +++++++------------ projects/clr/hipamd/src/hip_module.cpp | 10 ++++++++ 2 files changed, 19 insertions(+), 16 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 7d9f08ac69..cee4e81054 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1419,9 +1419,14 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, hipMemcpyKind kind, hipStream_t stream __dparm(0)); #else -__attribute__((visibility("hidden"))) hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*); +namespace hip_impl { +inline +__attribute__((visibility("hidden"))) +hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, + const char* name); +} // Namespace hip_impl. /** * @brief Copies the memory address of symbol @p symbolName to @p devPtr @@ -1439,7 +1444,7 @@ hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { //HIP_INIT_API(hipGetSymbolAddress, devPtr, symbolName); hip_impl::hip_init(); size_t size = 0; - return hipModuleGetGlobal(devPtr, &size, 0, (const char*)symbolName); + return hip_impl::read_agent_global_from_process(devPtr, &size, (const char*)symbolName); } @@ -1459,7 +1464,7 @@ hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) { // HIP_INIT_API(hipGetSymbolSize, size, symbolName); hip_impl::hip_init(); void* devPtr = nullptr; - return hipModuleGetGlobal(&devPtr, size, 0, (const char*)symbolName); + return hip_impl::read_agent_global_from_process(&devPtr, size, (const char*)symbolName); } #if defined(__cplusplus) @@ -2710,20 +2715,8 @@ extern "C" { * * @returns hipSuccess, hipErrorInvalidValue, hipErrorNotInitialized */ -inline -__attribute__((visibility("hidden"))) hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, - hipModule_t hmod, const char* name) { - if (!dptr || !bytes) return hipErrorInvalidValue; - - if (!name) return hipErrorNotInitialized; - - const auto r = hmod ? - hip_impl::read_agent_global_from_module(dptr, bytes, hmod, name) : - hip_impl::read_agent_global_from_process(dptr, bytes, name); - - return r; -} + hipModule_t hmod, const char* name); #endif // __HIP_VDI__ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name); diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 6401e13727..e81204d86b 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -282,6 +282,16 @@ 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::read_agent_global_from_module(dptr, bytes, hmod, name); +} + namespace hip_impl { hsa_executable_t executable_for(hipModule_t hmod) { return hmod->executable;