- It's a common mistake by assuming 1 << shamt would be promoted to
64-bit, if shamt is a 64-bit integer. That's not the case. Replace
that left shift to a 64-bit one to ensure it won't fall into undefined
behavior.
- Fix the host-side implementation as well for device function testing.
- Separate the definition of `__HCC_OR_HIP_CLANG__`, `__HCC_ONLY__`, and
`__HIP_CLANG_ONLY__` into hip_common.h so that it could be included in
hip_fp16.h, which may be included separately in app.
* Update hip_runtime_api.h
when i try to use mpicc or gcc to compile a c language code which call some hip runtime api , error occured as
> /path/to/hcc_detail/hip_runtime_api.h:2268:33: error: unknown type name ‘hipFuncAttributes’;
> hipFuncGetAttributes(hipFuncAttributes* attr, const void* func);
add ' struct ' for the first parameter of hipFuncGetAttributes will get ride of this problem.
* Add CMAKE dep to amd_comgr
* Use COMGR for read_kernarg_metadata in COV2
* Do not assume kernargs exist
* Add proper metadata destroy cleanup
* Use a process function for easier destroy
* Remove old read_kernarg_metadata
* Clean up HCC, prints, names
* Use COMGR in CMAKE by default
* Move metadata lookup for keyword values into helper
* Remove C string usage for lookup_keyword_value
* Guard COMGR for non-NVCC path
* Add hip_hcc dependency on comgr package
* Add lifetime to metadata nodes
* Find COMGR config file for amd_comgr target
* Move set_active data earlier
* 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
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.
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
* adding prof primitives generator
* minor change, renaming
* minor cosmetic changes, comments correcting and dead code removing
* minor changes and renaming
* minor chane, fixing comments
* 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
Maneesh Gupta
Rahul Garg
Michael LIAO
Aaron Enye Shi
Rahul Garg
Evgeny
Konstantin Pyzhov
David Salinas
Mr-LiuSw
Yaxun (Sam) Liu
Siu Chi Chan
Wen-Heng (Jack) Chung
Jeff Daily
Nico Trost
eshcherb