3951fccb4e
Fix memory leak introduced by previous change to Agent_global.
...
Make Agent_global manage the lifetime of the name string
[ROCm/clr commit: fa564a5345
2019-03-11 19:51:32 +00:00
3562ddf981
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
[ROCm/clr commit: fcde9fe3df
2019-03-11 19:51:25 +00:00
4326d67550
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*
[ROCm/clr commit: f8aca3159c
2019-03-11 19:51:18 +00:00
2a11a5529e
Merge pull request #949 from gargrahul/single_stream_concurrent_kernels
...
Add extension for kernel concurrency on same stream
[ROCm/clr commit: 3955f2c131
2019-03-06 17:34:54 +05:30
45f4ac5023
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.
[ROCm/clr commit: ed48847237
2019-03-06 17:31:44 +05:30
f364b32e29
Add extension for kernel concurrency on same stream
...
[ROCm/clr commit: 263e82a67a
2019-03-06 12:55:39 +05:30
2f64c76789
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.
[ROCm/clr commit: 6e68d44220
2019-01-08 17:33:40 +00:00
f2e2eb68e7
Let hip-clang support --genco
...
[ROCm/clr commit: 988dcd1e4a
2018-11-27 15:55:50 -05:00
938c731b27
Merge pull request #760 from eshcherb/roctracer-hip-frontend-181113
...
Roctracer hip frontend 181113
[ROCm/clr commit: 99bb89b756
2018-11-23 11:08:25 +05:30
4b77b5c1ba
Merge branch 'master' into getsymboladdress
...
[ROCm/clr commit: e9b88711e2
2018-11-20 12:03:22 +01:00
96b7f0404e
renaming HIP_INIT_CB_API to HIP_INIT_API
...
[ROCm/clr commit: e362688adf
2018-11-13 15:33:26 +00:00
6b45949728
adding lost i the merge change
...
[ROCm/clr commit: a13e035fd9
2018-11-13 15:33:26 +00:00
7519b3f75e
adding activity prof layer
...
[ROCm/clr commit: 084a68be63
2018-11-13 15:33:26 +00:00
ea9ea82fe1
Fixed symbol tracking device index
...
[ROCm/clr commit: 70bb04cc61
2018-11-13 07:01:17 +05:30
387e6de9e6
Merge pull request #721 from fronteer/my-fix
...
Make correct checking of the returned hipDeviceptr_t from read_global…
[ROCm/clr commit: 0bafc38135
2018-11-08 11:42:08 +05:30
a857354072
Introduce ihipModuleGetGlobal
...
[ROCm/clr commit: 4da2d92281
2018-11-06 09:54:34 +01:00
8856b89b21
Fixes global symbols tracking in hip_module
...
[ROCm/clr commit: 592efa017f
2018-10-31 03:22:38 +05:30
b12c16f72b
Make correct checking of the returned hipDeviceptr_t from read_global_description()
...
[ROCm/clr commit: 443698ce00
2018-10-23 21:13:11 +08:00
677e945726
Update hip_module.cpp
...
Typo.
[ROCm/clr commit: 43fca684c8
2018-05-18 17:50:45 +01:00
5a9d5a931b
Update hip_module.cpp
...
[ROCm/clr commit: 40a22d235e
2018-05-14 17:15:36 +01:00
2973f68465
Don't use magic constants, they're evil.
...
Also clarify that the register count cannot be queried at the moment.
[ROCm/clr commit: eded014abc
2018-05-11 11:31:46 +01:00
4e6996282d
Add support for the hipFuncGetAttributes interface.
...
[ROCm/clr commit: bf9529aaa8
2018-05-11 03:35:10 +01:00
181e0ee8ee
hip_module code cleanup
...
-Fixed missing ihipLogStatus in hipModuleLoad()
-Fixed some ihipXXX functions
[ROCm/clr commit: c23898f49b
2018-04-16 15:35:04 +05:30
46ddefedee
Apply .clangformat to all repo source files
...
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344
[ROCm/clr commit: 9e47fccc89
2018-03-12 11:29:03 +05:30
eab8624938
Re-sync with upstream.
...
[ROCm/clr commit: a704bd8b44
2018-02-12 20:20:24 +00:00
f15543d5cf
Fixed host allocated globals address lookup for host usage
...
Fixed texture driver APIs failure
[ROCm/clr commit: b8c23f979b
2018-01-30 18:06:31 +05:30
1ee7e7c2f9
Fixed build error
...
[ROCm/clr commit: d1dcc5025d
2017-12-28 16:15:45 +05:30
aa48cc7b55
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.
...
[ROCm/clr commit: 4e0739c68a
2017-12-08 04:22:57 +00:00
37e40fcd40
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.
...
[ROCm/clr commit: e186bd9533
2017-12-03 23:09:06 +00:00
cdac21c81a
Choose whether or not to use functional grid_launch based on the version of HCC used to compile.
...
[ROCm/clr commit: c6ab6f292b
2017-11-29 00:17:44 +00:00
914acbe88f
Re-sync with upstream and re-factor platform global management for texture references.
...
[ROCm/clr commit: 6a0efb7ed2
2017-11-28 19:15:29 +00:00
579a3187da
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
# Conflicts:
# src/hip_module.cpp
[ROCm/clr commit: d37a5a6008
2017-11-28 17:29:11 +00:00
212afdf56b
Fixed review comments
...
[ROCm/clr commit: 38029f2849
2017-11-21 21:19:06 +05:30
1366b2d5ae
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.
...
[ROCm/clr commit: f8c1c1b38e
2017-11-21 02:40:34 +00:00
9cee9705dc
-Moved coGlobals in hipModule class (takes care of multi module case)
...
-Used mutex scope for updating coGlobals
[ROCm/clr commit: eff09fcd86
2017-11-20 16:23:18 +05:30
7453c8b25c
Update hipModuleGetTexRef API
...
[ROCm/clr commit: 76f174b536
2017-11-19 22:10:46 +05:30
93df8a2a57
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
[ROCm/clr commit: 1518611976
2017-11-09 23:43:07 +00:00
bb8c31c198
Texture driver APIs support
...
[ROCm/clr commit: 1851c153f6
2017-11-09 22:10:55 +05:30
e35ce82eff
This introduces correct support for agent global variables, and implements hipModuleGetGlobal as an actual equivalent for cuModuleGetGlobal.
...
[ROCm/clr commit: d12cf0da7d
2017-11-03 01:44:48 +00:00
19935275ea
Merge pull request #245 from scchan/centos_fixes
...
various fixes for centos/rhel
[ROCm/clr commit: 5a80c4108e
2017-11-01 18:10:29 +01:00
b738633bcb
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.
...
[ROCm/clr commit: 28f87f7d2e
2017-11-01 15:09:59 +00:00
2094b92a3e
Centos/RHEL - remove usage of constexpr since libc++ doesn't enable ctor for constexpr pair in C++11
...
[ROCm/clr commit: c1a4ff7b20
2017-10-31 18:16:12 +00:00
4df6b371ab
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.
[ROCm/clr commit: dc7d993a02
2017-10-30 20:20:51 +00:00
941682c701
Make elfio headers private
...
Change-Id: I3ba174bb46e84a75380207d93a0da6fe3703689e
[ROCm/clr commit: 281c543bc3
2017-10-23 10:24:36 +05:30
1c8db33de9
This fixes incorrect usage of the reader object, which created arcane
...
mismatches when one reader accessed another's section.
[ROCm/clr commit: b222448efd
2017-10-09 15:46:38 +01:00
d9075301c3
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.
[ROCm/clr commit: 582779c6c6
2017-10-09 13:27:11 +01:00
51c48e4cef
Bump min hcc_workweek required for named kernel dispatch to 17312
...
Change-Id: I8c7b58306b279ed113d03260e4bc6086bb8b4e68
[ROCm/clr commit: dabd6d2bc6
2017-08-08 11:08:55 +05:30
004080bff8
Add workweek check to make sure we have a new enough compiler
...
[ROCm/clr commit: 2781293e89
2017-07-27 23:00:58 -05:00
c8b72cb926
Pass kernel name to HCC dispatch_hsa_kernel, for debug/profile
...
[ROCm/clr commit: a282351763
2017-07-27 22:00:15 -05:00
ccc085e4db
Add option to pass names to HCC dispatch API (for debug)
...
[ROCm/clr commit: c7382f7da6
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