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
6cc5dc0326
Set event->_stream on hipHccModuleLaunchKernel path if start/stop used
...
Ensure _stream is always non-null in recorded events.
Fixes isDefaultStream fault.
2017-05-30 21:55:46 -05:00
ae9f14ef9c
[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.
2017-05-19 17:22:14 +03:00
8bc6ee5932
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.
2017-05-16 19:04:25 -05:00
e9eaadd135
Finish adding start/stop event support to hipHccModuleLaunchKernel.
...
Change interface to use hipEvent_t rather than hipEvent_t*
Change-Id: I259062dc087a13d51dc27f84e1e8861f332a104d
2017-04-06 21:02:50 -05:00
f442e975c6
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
2017-04-06 21:02:50 -05:00
81d3725a16
added debug support for HIP sample
...
Change-Id: Ia7265234082039b68114f7421f4dbcb7149d4d2b
2017-03-31 14:13:46 -05:00
1cead6a4cd
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
2017-03-31 12:11:34 -05:00
4e5861df7a
disabled metadata apis
...
Change-Id: Ifb8839c581644cccc2afcd18c38a866f649a4144
2017-03-28 10:46:31 -05:00
abf369049b
Re-enabled metadata parsing in HIP
...
Change-Id: If8caa844571cb8581450df9ffdb76e2445c75f13
2017-03-27 11:00:39 -05:00
5a0685c137
removed llvm dependency and metadata functionality
...
Change-Id: Ib9783b75d326559ed29c5aa2218aed40d20ad0fb
2017-03-23 10:16:37 -05:00
9645b0e0dc
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
2017-03-17 13:11:34 -05:00
09df0977c0
Fix bug in hipModuleGetFunction.
...
Modules with > 1 function didn't return the function correctly.
Also fix coding convention issues
2017-03-08 13:52:38 -06:00
1cffe49775
v3: added free for ihipModuleSymbol_t structures inside tracker
...
Change-Id: Ib8041a05312c08cbdf2d4fee5e7cbae17df6efff
2017-02-10 13:42:10 -06:00
554fe191e1
v2: Fixed hipModule memory management
...
1. Changed test to assert for same hipFunction values
2. Added better memory management for hipModule
Change-Id: I10d7aef13c215a2211e262f3c79017f26a17d9a7
2017-02-10 13:32:13 -06:00
6110fcdfb6
fixed hipFunction memory management
...
Change-Id: I7ebb323419bcd220ebd6466a8eb38e7bfdb1520a
2017-02-09 17:22:55 -06:00
2f7a8ec39c
Read HCC_OPT_FLUSH and optimize dispatch accordingly.
...
If HCC is in this mode, we can use less aggressive flushes in some
cases.
2017-01-25 21:50:52 -06:00
27447f10c2
Show dynamic shared mem usage not static.
2017-01-23 22:34:41 -06:00
2ffc9f4e22
Fix debug display for Module launch kernels
2017-01-19 12:40:45 -06:00
e9ff23e5f9
changed copyright year from 2016 to 2017 in src directory
...
Change-Id: Idb97db509b2b4b1656b2df7a14a62ade38c9d574
2017-01-11 18:05:41 -06:00
5d815937de
Add name for function
2016-12-17 08:54:09 -06:00
2bd70ff345
Remove HSA dependency from hipFunction_t
...
Place _groupSegmentSize and _privateSegmentSize inside Function,
remove hsa_executable_symbol_t.
2016-12-17 07:22:56 -06:00
06d382bc6d
Remove USE_DISPATCH_HSA_KERNEL=0 path.
2016-12-17 07:22:56 -06:00
3f9404d0e1
Refactor Module and Function APIs.
...
- hipFunction_t is now returned by value. This eliminates dynamic
allocation / memory management complexity in the module. Removed
the kernel
name so the structure is just 16 bytes now.
- Moved the hsa_executable_load_module and hsa_executable_freeze
calls to the hipModuleLoad and hipModuleLoadData calls.
- Apply sharedMemBytes in hipModuleLaunchKernel to group segment
size (not private).
2016-12-17 07:22:33 -06:00
3f810a203b
Enabled USE_DISPATCH_HSA_KERNEL, with serialization in hipModuleUnload.
2016-12-06 10:14:01 -06:00
4cb1579d4a
HIP resource leaks fix from Jack
...
Change-Id: I93f3ad7cb94ff1cba1577bd8acc90e826693d12e
2016-12-05 20:21:33 -06:00
ac93376c26
Revert "Enable USE_DISPATCH_HSA_KERNEL."
...
This reverts commit f8bcbe8680
2016-12-05 16:55:26 +05:30
Rahul Garg
Alex Voicu
Ben Sander
Siu Chi Chan
Maneesh Gupta
Evgeny Mankov
sunway513
Aditya Atluri