* Add ersatz for NVRTC.
* Fix extraneous paren and use correct namespace.
* Use lowerCamelCase (yuck, yuck) consistently.
* Link against FS when building hiprtc lib.
* Correctly mark Manipulators. Fix dual compile.
* Add unit tests. Extend HIT to accept linker options.
* Make sure the HIPRTC library is installed.
* Better logging. Try to auto-detect the target.
* Stop specifying the target explicitly.
* Add missing flavour of `hipModuleLaunchKernel`.
* Program was already destroyed.
* Don't use `--genco`. Fix mangled name trimming.
* Fix HIPRTC breakage due to upstream noise.
* [dtests] Replace RUN -> TEST in hiprtc tests
Change-Id: Ie499e92dfe4e5c94634b1c2b76cf52d241bcfea3
* [hit] Set HIP_PATH to HIP_ROOT_DIR for all tests
Change-Id: Ib0ad1f99bc71c03e363e055dd508a7a4a210680a
[ROCm/clr commit: a538eb705a]
* Revert "Revert "Use COMgr to read Kernel Args Metadata (#1006)""
This reverts commit f8d108a815.
* Revert "Use COMgr to read Kernel Args Metadata (#1006)"
This reverts commit 10048a5631.
* Revert "improve program state commentary"
This reverts commit 5233d41c6c.
* Revert "load program state once per agent"
This reverts commit 9cee2c5311.
* start moving function_names() into the hip shared lib
* start moving code_object_blobs to a new "state" object
* Consolidate various program state related static objects into a
single program_state object
* minor clean up
* move more stuffs from functional_grid_launch into program_state
* debug make_kernarg
* moving lookup for kernargs size_align into program_state
* clean up old code for kernarg size and alignment
* update hip_module to use newer api in program_state
* Create public member functions for program_state
* move most program state functions into shared library
* Pass the data buffer size to load_executable
Otherwise, it can't figure what the data size is
just from the char* (since the data is not really a string)
* turning free functions in program state into members of program_state_impl
* change the free function globals() into a member of program_state_impl
* replace the static mutex used for populating globals
* moving associate_code_object_symbols_with_host_allocation into
program_state_impl
* move load_code_object_and_freeze_executable into program_state_impl
* moving executables and functions_names into program_state_impl
* moving kernels() into program_state_impl
* moving functions() into program_state_impl
* move get_kernargs into program_state_impl
* moving kernel_descriptor into program_state_impl
* moving kernargs_size_align calculation into program_state_impl
* Changing the handle to program_state_impl to a pointer
* moving program_state_impl into a separate inline source file
* fixing/cleaning up some header file includes
* moving member function for kernargs_size_align into program_state.cpp
* moving Kernel_descriptor into program_state.inl
* add a new class to manage agent globals
* moving all agent globals processing functions into agent_globals_impl
* load program state once per agent
re-merging PR991 against other program state changes
* fix per-agent program state member initialization
* cache executables based on elf name, isa, and agent.
This avoids program state reloading executables after a shared library is dlopened.
re-merging PR1057 against other program state changes
* protect executables cache by a global mutex
* return ref to executables cache
* adapt PR#981 Make hipModuleGetGlobal be in HIP runtime
[ROCm/clr commit: 05a1b696da]
* Return hipErrorInsufficientDriver status when CPU device not found - no exception thrown
* Return hipErrorInsufficientDriver status when CPU device not found
[ROCm/clr commit: ebf986dcee]
* In hipFree, if memory is associated with a device, synchronize that device's streams.
This changes the behavior from synchronizing the currently set TLS device.
* All devices sync in hipFree for _appId=-1 case.
* Revert "All devices sync in hipFree for _appId=-1 case."
This reverts commit 1efb34d6a8426661e45bc5f763422a1147aeac10.
* add HIP_SYNC_FREE env var
[ROCm/clr commit: cf8fb43e6b]
__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.
[ROCm/clr commit: 8f5c812a68]
* 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]
* 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]
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]
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.
[ROCm/clr commit: 972ca06c4c]
Siu Chi Chan
Evgeny Mankov
Alex Voicu
Wenkai Du
Maneesh Gupta
Rahul Garg
wkwchau
Rahul Garg
Sameer Sahasrabuddhe
Yaxun (Sam) Liu
Jeff Daily
Wen-Heng (Jack) Chung
Evgeny
Aaron Enye Shi
lmoriche
Laurent Morichetti