* 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: 04915cea2f]
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: f5e4fff6cc]
For code objects with global symbols of length 0, ROCR runtime would
ignore them even though they exist in the symbol table. Therefore the
result from read_agent_globals() can't be trusted entirely.
As a workaround to tame applications which depend on the existence of
global symbols with length 0, always return hipSuccess here.
This behavior shall be reverted once ROCR runtime has been fixed to
address SWDEV-173477
[ROCm/clr commit: cf7ad0f184]
* adding prof primitives generator
* minor change, renaming
* minor cosmetic changes, comments correcting and dead code removing
* minor changes and renaming
* minor chane, fixing comments
[ROCm/clr commit: 1229750546]
* 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
[ROCm/clr commit: fa9495841b]
- `immarg`, immediate argument, is enabled on all AMDGPU intrinsics.
Revise device functions using these intrinsics with immediate
arguments.
[ROCm/clr commit: b74b4500c4]
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
[ROCm/clr commit: fcde9fe3df]
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*
[ROCm/clr commit: f8aca3159c]
* 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.
[ROCm/clr commit: ed48847237]
Add 2 extra memset functions which fills memory with integer-typed data
Also change the parameters of ihipMemset to better explain the semantic
[ROCm/clr commit: 392271f4db]
hipMalloc and hipHostMalloc are C++ function templates by default. In some
applications they need to discover the decltype of underlying C functions. In
this commit, a new macro __HIP_DISABLE_CPP_FUNCTIONS__ is introduced to
suppress the C++ function templates.
Existing HIP applications don't need to be changed. hipMalloc and hipHostMalloc
will behave just the same.
For new applications which depends on identifying decltypes of underlying C
functions, __HIP_DISABLE_CPP_FUNCTIONS__ shall be defined prior to including
HIP headers.
[ROCm/clr commit: f4e1f74e9e]
Wen-Heng (Jack) Chung
Jeff Daily
Maneesh Gupta
Michael LIAO
Nico Trost
eshcherb
Siu Chi Chan
Yaxun Sam Liu
Evgeny
Aaron Enye Shi
Michael Liao
Steven Noonan
Alex Voicu
Rahul Garg