__hipRegisterFunction is called during by .init functions during program initialization.
It calls hipModuleGetFunction to locate kernel symbol in code objects. hipModuleGetFunction
assumes current device when locating kernel symbols. This works for HCC but not for hip-clang,
since hip-clang needs to locate kernel symbols for different devices without switching
between devices.
This patch introduces a new hsa agent parameter to ihipModuleGetFunction, which allows
__hipRegisterFunction to choose the correct hsa agent when locating kernel symbols. By
default it uses this_agent(), therefore this patch has no impact on HCC.
* 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
* reimplement HIP_INIT as a function, expose it as hip_impl::hip_init()
so that it could be called from hipLaunchKernelGGL and other inlined
HIP functions
* Don't call hip_init from ihipPreLaunchKernel
Issue: Header uses std::vector<Agent_global> agent_globals which is created by hip_module.cpp
- Move iterator fails to copy Agent_global from library source into header version
- Due to different versions of std::string name in struct Agent_global
Fix: Change Agent_global to use char* name instead of std::string name
Issue: mismatch undefined symbols in different user env
- Binary expects modified return value std::string&
- Fails to match libhip_hcc.so: return value is std::string& but doesn't match modified C++ env
Fix: Change return value to char*, create new key std::string in header from char*
* Initial attempt to switch over to internally linked state.
* Add missing CMake update.
* hipLaunchKernelGGLImpl must be inline as well. Ensure internal linkage.
* Ensure global retrieval uses internally linked state.
* Hide HC in the implementation. Minimise ADL woes.
* Strange software exists, and must be catered to.
* Use a less spammy mechanism for ensuring internal linkage / non-export.
* Remove leftover internal detail.
The implementation is copied from HCC runtime.
For hcc it has no effect since apps can find them in either hcc runtime or HIP
runtime.
hip-clang needs it in HIP/HCC runtime so that HIP/HCC and HIP/VDI runtime are
swappable.
A hash calculated via FNV-1a algorithm is introduced in ihipModule_t, the
internal of hipModule_t. The hash is used by HIP module APIs such as
- read_agent_global_from_module
to determine whether the agent-scope globals for a module have been iterated.
This commit fixes one issue that applications which load / unload modules
frequently would occasionally fail. After deep investigation of the issue it
turns out the old implementation in read_agent_global_from_module uses
hipModule_t as the key, which is not robust enough, as hipModule_t instances
are allocated dynamically so there are cases that one memory address may be
used by multiple hipModule_t instances. The real solution is to introduce a
uniquely identifiable hash for the code object associated with the HIP module.
And that's the rationale behind this commit.
Yaxun Sam Liu
Wen-Heng (Jack) Chung
Maneesh Gupta
Rahul Garg
Jeff Daily
Siu Chi Chan
Evgeny
Aaron Enye Shi
Alex Voicu
Wilkin Chau
lmoriche
Laurent Morichetti
Rahul Garg
Qianfeng Zhang