fa564a5345
Fix memory leak introduced by previous change to Agent_global.
...
Make Agent_global manage the lifetime of the name string
2019-03-11 19:51:32 +00:00
fcde9fe3df
Fix Agent_global variables failing hipTestDeviceSymbol
...
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
2019-03-11 19:51:25 +00:00
f8aca3159c
Fix hash_for undefined reference in hipTestConstant test
...
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*
2019-03-11 19:51:18 +00:00
3955f2c131
Merge pull request #949 from gargrahul/single_stream_concurrent_kernels
...
Add extension for kernel concurrency on same stream
2019-03-06 17:34:54 +05:30
ed48847237
dlopen() fixes ( #929 )
...
* 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.
2019-03-06 17:31:44 +05:30
263e82a67a
Add extension for kernel concurrency on same stream
2019-03-06 12:55:39 +05:30
6e68d44220
Introduce hash key to HIP module implementation
...
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.
2019-01-08 17:33:40 +00:00
988dcd1e4a
Let hip-clang support --genco
2018-11-27 15:55:50 -05:00
99bb89b756
Merge pull request #760 from eshcherb/roctracer-hip-frontend-181113
...
Roctracer hip frontend 181113
2018-11-23 11:08:25 +05:30
e9b88711e2
Merge branch 'master' into getsymboladdress
2018-11-20 12:03:22 +01:00
e362688adf
renaming HIP_INIT_CB_API to HIP_INIT_API
2018-11-13 15:33:26 +00:00
a13e035fd9
adding lost i the merge change
2018-11-13 15:33:26 +00:00
084a68be63
adding activity prof layer
2018-11-13 15:33:26 +00:00
70bb04cc61
Fixed symbol tracking device index
2018-11-13 07:01:17 +05:30
0bafc38135
Merge pull request #721 from fronteer/my-fix
...
Make correct checking of the returned hipDeviceptr_t from read_global…
2018-11-08 11:42:08 +05:30
4da2d92281
Introduce ihipModuleGetGlobal
2018-11-06 09:54:34 +01:00
592efa017f
Fixes global symbols tracking in hip_module
2018-10-31 03:22:38 +05:30
443698ce00
Make correct checking of the returned hipDeviceptr_t from read_global_description()
2018-10-23 21:13:11 +08:00
43fca684c8
Update hip_module.cpp
...
Typo.
2018-05-18 17:50:45 +01:00
40a22d235e
Update hip_module.cpp
2018-05-14 17:15:36 +01:00
eded014abc
Don't use magic constants, they're evil.
...
Also clarify that the register count cannot be queried at the moment.
2018-05-11 11:31:46 +01:00
bf9529aaa8
Add support for the hipFuncGetAttributes interface.
2018-05-11 03:35:10 +01:00
c23898f49b
hip_module code cleanup
...
-Fixed missing ihipLogStatus in hipModuleLoad()
-Fixed some ihipXXX functions
2018-04-16 15:35:04 +05:30
9e47fccc89
Apply .clangformat to all repo source files
...
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344
2018-03-12 11:29:03 +05:30
a704bd8b44
Re-sync with upstream.
2018-02-12 20:20:24 +00:00
b8c23f979b
Fixed host allocated globals address lookup for host usage
...
Fixed texture driver APIs failure
2018-01-30 18:06:31 +05:30
d1dcc5025d
Fixed build error
2017-12-28 16:15:45 +05:30
4e0739c68a
This introduces LipoProteinLipase (lpl), a simple tool for creating fat binaries. It represents a direct replacement of the creaky hccgenco.sh script, which had various issues. The format it uses is that of a code object bundle, generated by the Clang Offload Bundler. The output is always suffixed with the ".adipose" extension. It is shared with HCC. The hipcc script and associated tests are modified to use lpl. Help can be obtained by invoking lpl --help. A more computer-sciency / corporate friendly name is likely to be beneficial, which is a reason for choosing easily searchable/replaceable names such as lpl or adipose.
2017-12-08 04:22:57 +00:00
e186bd9533
This is primarily intended as an additional cleanup of the module functionality, in the aftermath of adopting module based dispatch. The main effort was associated with refactoring the questionable ihipModuleGetSymbol. It was quaintly written and misleading, in that it had little to do with getting symbols, and was exactly retrieving a kernel object. Error handling is modified so as to reduce branching depth. Functions which serve as interfaces to the HSA RT are moved in a separate helper header. Code object readers are properly deleted. Some leftover dead functionality pertaining to associating namespace scope variables with their allocated memory is removed. Executable loading is changed to use a string which holds the ELF image of the code object being loaded, thus avoiding some corner cases where using a istream would fail.
2017-12-03 23:09:06 +00:00
c6ab6f292b
Choose whether or not to use functional grid_launch based on the version of HCC used to compile.
2017-11-29 00:17:44 +00:00
6a0efb7ed2
Re-sync with upstream and re-factor platform global management for texture references.
2017-11-28 19:15:29 +00:00
d37a5a6008
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
# Conflicts:
# src/hip_module.cpp
2017-11-28 17:29:11 +00:00
38029f2849
Fixed review comments
2017-11-21 21:19:06 +05:30
f8c1c1b38e
Refactor the __device__ versions of memset and memcpy to be less awkward i.e. not return nullptr as opposed to the destination pointer (it can only be assumed it was done for maximum confusion) and actually unroll as they claim to. Change all of the {to, from}Symbol functions to use hipModuleGetGlobal, as opposed to hc::accelerator::get_symbol_address which is no longer valid with module based dispatch.
2017-11-21 02:40:34 +00:00
eff09fcd86
-Moved coGlobals in hipModule class (takes care of multi module case)
...
-Used mutex scope for updating coGlobals
2017-11-20 16:23:18 +05:30
76f174b536
Update hipModuleGetTexRef API
2017-11-19 22:10:46 +05:30
1518611976
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
2017-11-09 23:43:07 +00:00
1851c153f6
Texture driver APIs support
2017-11-09 22:10:55 +05:30
d12cf0da7d
This introduces correct support for agent global variables, and implements hipModuleGetGlobal as an actual equivalent for cuModuleGetGlobal.
2017-11-03 01:44:48 +00:00
5a80c4108e
Merge pull request #245 from scchan/centos_fixes
...
various fixes for centos/rhel
2017-11-01 18:10:29 +01:00
28f87f7d2e
This switches HIP from its currently convoluted macro + pfe based dispatch mechanism to a more natural one partially based on the existing module API. The basic idea is that HCC will always correctly emit __global__ functions: as empty-bodied stubs, on host, and as kernels, on device. It then becomes trivial to obtain the mangled name on host, at dispatch, from the function's address, and then to use the mangled name to retrieve the kernel. This should address all problems stemming from serialisation, dubious mismatches due to the manufactured functor, macro-isms et al. It also immediately enables support for generalised globals as a consequence of that being available in the module API. Finally, it will make debug much easier, since the actual names of the __global__ functions will automatically be used in traces etc. One detail is that due to how dispatch works now (hipLaunchKernel and hipLaunchKernelGGL are themselves variadic function templates which deduce the function type of the callee), in certain cases it may be necesssary to insert explicit casts to ensure that the variadic argument list selects a viable overload - this can be observed in some unit tests. Eventually we may be able to remove this limitation, but for now it does not appear terribly onerous. The code is not extremely HIPpie, nor is it fully optimised, but rather is intended as a starting point for the HIP team to make its own.
2017-11-01 15:09:59 +00:00
c1a4ff7b20
Centos/RHEL - remove usage of constexpr since libc++ doesn't enable ctor for constexpr pair in C++11
2017-10-31 18:16:12 +00:00
dc7d993a02
Add ns-level timer for HIP API routines
...
Refactor some miuses of ihipLogStatus, these should only be in top-level
HIP APIs and should be paired with HIP_API_INIT calls.
2017-10-30 20:20:51 +00:00
281c543bc3
Make elfio headers private
...
Change-Id: I3ba174bb46e84a75380207d93a0da6fe3703689e
2017-10-23 10:24:36 +05:30
b222448efd
This fixes incorrect usage of the reader object, which created arcane
...
mismatches when one reader accessed another's section.
2017-10-09 15:46:38 +01:00
582779c6c6
This adds cursory support for globals to the HIP module loading API. The
...
style is purposefully alien so as to signal that HIP experts should turn
it into HIP worthy code as soon as possible.
2017-10-09 13:27:11 +01:00
dabd6d2bc6
Bump min hcc_workweek required for named kernel dispatch to 17312
...
Change-Id: I8c7b58306b279ed113d03260e4bc6086bb8b4e68
2017-08-08 11:08:55 +05:30
2781293e89
Add workweek check to make sure we have a new enough compiler
2017-07-27 23:00:58 -05:00
a282351763
Pass kernel name to HCC dispatch_hsa_kernel, for debug/profile
2017-07-27 22:00:15 -05:00
c7382f7da6
Add option to pass names to HCC dispatch API (for debug)
2017-06-23 18:05:30 -05:00
Siu Chi Chan
Aaron Enye Shi
Maneesh Gupta
Alex Voicu
Rahul Garg
Wen-Heng (Jack) Chung
Yaxun Sam Liu
Michael Kuron
Evgeny
Michael Kuron
Qianfeng Zhang
Ben Sander