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
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
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
81354999e8
Set event->_stream on hipHccModuleLaunchKernel path if start/stop used
...
Ensure _stream is always non-null in recorded events.
Fixes isDefaultStream fault.
[ROCm/clr commit: 6cc5dc0326
2017-05-30 21:55:46 -05:00
052f630bd3
[HIP] [HIPIFY] [FIX] cuModuleLoadDataEx -> hipModuleLoadDataEx
...
https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/81
1. Do not use JIT options on HCC path, call hipModuleLoadData instead.
2. NVCC path is unchanged, to call cuModuleLoadDataEx with all options.
3. Get rid of manual hipification, based on #ifdef #else for NVCC/HIP.
4. Update documentation accordingly.
[ROCm/clr commit: ae9f14ef9c
2017-05-19 17:22:14 +03:00
4ac6ac9d1d
Add initial HIP_SYNC_NULL_STREAM=0 mode.
...
This eliminates host-synchronization for null stream. Instead, the
null-stream uses GPU-side events to wait for other streams.
Default is OFF pending additional testing.
Add enhanced null-stream test.
Also refine HIP_TRACE_API.
[ROCm/clr commit: 8bc6ee5932
2017-05-16 19:04:25 -05:00
dc9e957a36
Finish adding start/stop event support to hipHccModuleLaunchKernel.
...
Change interface to use hipEvent_t rather than hipEvent_t*
Change-Id: I259062dc087a13d51dc27f84e1e8861f332a104d
[ROCm/clr commit: e9eaadd135
2017-04-06 21:02:50 -05:00
1897ca4a20
Refactor events and add initial event option for hipHccModuleLaunchKernel
...
- Change hipEvent_t to a class.
- Move event logic inside the class.
- Add _type to support Independent, StartCommand, StopCommand events.
StartCommand returns start timestamp from events.
Change-Id: I4ddd694f2645a3ff7170c9111dc1d3e39931ca21
[ROCm/clr commit: f442e975c6
2017-04-06 21:02:50 -05:00
28c8d15412
added debug support for HIP sample
...
Change-Id: Ia7265234082039b68114f7421f4dbcb7149d4d2b
[ROCm/clr commit: 81d3725a16
2017-03-31 14:13:46 -05:00
3651cb672a
added new api hipHccModuleLaunchKernel
...
1. hipHccModuleLaunchKernel is same as hipModuleLaunchKernel with OpenCL workitem model
2. Added copy right
3. Fixed header naming
Change-Id: I6a7c35a3566e2f8d3f5056613e34193775d4b236
[ROCm/clr commit: 1cead6a4cd
2017-03-31 12:11:34 -05:00
bdaf6a3aaa
disabled metadata apis
...
Change-Id: Ifb8839c581644cccc2afcd18c38a866f649a4144
[ROCm/clr commit: 4e5861df7a
2017-03-28 10:46:31 -05:00
c3f1a37c28
Re-enabled metadata parsing in HIP
...
Change-Id: If8caa844571cb8581450df9ffdb76e2445c75f13
[ROCm/clr commit: abf369049b
2017-03-27 11:00:39 -05:00
478eb4e350
removed llvm dependency and metadata functionality
...
Change-Id: Ib9783b75d326559ed29c5aa2218aed40d20ad0fb
[ROCm/clr commit: 5a0685c137
2017-03-23 10:16:37 -05:00
b4b87b8786
Added default module launch api functionality
...
1. As in hipModuleLaunchKernel(..., kernelParams, nullptr); works with this commit
2. Added headers AMDGPUPTNote.h, AMDGPURuntimeMetadata.h to do code object meta data parsing
3. Changed CMake to look at llvm link libraries
4. HIP developer should set env variable LLVM_HOME to remove link errors
5. HIP depends on installed LLVM (not source, not build)
6. Added sample to test out the feature
7. Right now HCC does not support embedding metadata in code object. Use clang opencl
8. Changed HIPCC to read LLVM_HOME env var
9. New argument to CMake should be given -DLLVM_HOME=<where llvm 5.0 is installed>
Change-Id: Iba38194aa872d97cc2c90a8e5ff746c48055c868
[ROCm/clr commit: 9645b0e0dc
2017-03-17 13:11:34 -05:00
0710dbde55
Fix bug in hipModuleGetFunction.
...
Modules with > 1 function didn't return the function correctly.
Also fix coding convention issues
[ROCm/clr commit: 09df0977c0
2017-03-08 13:52:38 -06:00
Evgeny
Maneesh Gupta
Rahul Garg
Qianfeng Zhang
Alex Voicu
Ben Sander
Siu Chi Chan
Evgeny Mankov
sunway513
Aditya Atluri