Commit Graph

110 Commits

Author SHA1 Message Date
Alex Voicu ccfb764a59 Add HIPRTC, glorious ersatz for NVRTC (#1097)
* Add ersatz for NVRTC.

* Fix extraneous paren and use correct namespace.

* Use lowerCamelCase (yuck, yuck) consistently.

* Link against FS when building hiprtc lib.

* Correctly mark Manipulators. Fix dual compile.

* Add unit tests. Extend HIT to accept linker options.

* Make sure the HIPRTC library is installed.

* Better logging. Try to auto-detect the target.

* Stop specifying the target explicitly.

* Add missing flavour of `hipModuleLaunchKernel`.

* Program was already destroyed.

* Don't use `--genco`. Fix mangled name trimming.

* Fix HIPRTC breakage due to upstream noise.

* [dtests] Replace RUN -> TEST in hiprtc tests

Change-Id: Ie499e92dfe4e5c94634b1c2b76cf52d241bcfea3

* [hit] Set HIP_PATH to HIP_ROOT_DIR for all tests

Change-Id: Ib0ad1f99bc71c03e363e055dd508a7a4a210680a
2019-05-16 18:28:54 +05:30
Siu Chi Chan f5eb91d53d migrate program_state logic from header into shared library (phase I) (#1077)
* Revert "Revert "Use COMgr to read Kernel Args Metadata (#1006)""

This reverts commit a3d118eaa8.

* Revert "Use COMgr to read Kernel Args Metadata (#1006)"

This reverts commit 8a548bf40b.

* Revert "improve program state commentary"

This reverts commit 7aada87cbd.

* Revert "load program state once per agent"

This reverts commit c9117de8eb.

* start moving function_names() into the hip shared lib

* start moving code_object_blobs to a new "state" object

* Consolidate various program state related static objects into a
single program_state object

* minor clean up

* move more stuffs from functional_grid_launch into program_state

* debug make_kernarg

* moving lookup for kernargs size_align into program_state

* clean up old code for kernarg size and alignment

* update hip_module to use newer api in program_state

* Create public member functions for program_state

* move most program state functions into shared library

* Pass the data buffer size to load_executable
Otherwise, it can't figure what the data size is
just from the char* (since the data is not really a string)

* turning free functions in program state into members of program_state_impl

* change the free function globals() into a member of program_state_impl

* replace the static mutex used for populating globals

* moving associate_code_object_symbols_with_host_allocation into
program_state_impl

* move load_code_object_and_freeze_executable into program_state_impl

* moving executables and functions_names into program_state_impl

* moving kernels() into program_state_impl

* moving functions() into program_state_impl

* move get_kernargs into program_state_impl

* moving kernel_descriptor into program_state_impl

* moving kernargs_size_align calculation into program_state_impl

* Changing the handle to program_state_impl to a pointer

* moving program_state_impl into a separate inline source file

* fixing/cleaning up some header file includes

* moving member function for kernargs_size_align into program_state.cpp

* moving Kernel_descriptor into program_state.inl

* add a new class to manage agent globals

* moving all agent globals processing functions into agent_globals_impl

* load program state once per agent

re-merging PR991 against other program state changes

* fix per-agent program state member initialization

* cache executables based on elf name, isa, and agent.

This avoids program state reloading executables after a shared library is dlopened.

re-merging PR1057 against other program state changes

* protect executables cache by a global mutex

* return ref to executables cache

* adapt PR#981 Make hipModuleGetGlobal be in HIP runtime
2019-05-12 19:24:03 +05:30
Yaxun (Sam) Liu bb5c620b13 Fix missing arg in HIP_INIT_API 2019-04-18 16:18:31 -04:00
Yaxun (Sam) Liu 271fdc4e4d Fix regression on multi-gpu due to PR#997 2019-04-05 22:54:41 -04:00
Yaxun Sam Liu 98b9e92908 hip-clang: fix kernel not found on multi-gpu
__hipRegisterFunction is called during by .init functions during program initialization.
It calls hipModuleGetFunction to locate kernel symbol in code objects. hipModuleGetFunction
assumes current device when locating kernel symbols. This works for HCC but not for hip-clang,
since hip-clang needs to locate kernel symbols for different devices without switching
between devices.

This patch introduces a new hsa agent parameter to ihipModuleGetFunction, which allows
__hipRegisterFunction to choose the correct hsa agent when locating kernel symbols. By
default it uses this_agent(), therefore this patch has no impact on HCC.
2019-03-31 10:08:20 -04:00
Wen-Heng (Jack) Chung 4b7177ac42 Make hipModuleGetGlobal be in HIP runtime so it can be discovered at runtime (#981)
* 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
2019-03-29 03:45:04 +00:00
Jeff Daily c9117de8eb load program state once per agent 2019-03-27 18:19:10 +00:00
Evgeny 31475c5ac8 tracing callback layer update 2019-03-14 22:43:52 -05:00
Siu Chi Chan bf1d48bf78 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
Aaron Enye Shi 00d24d254d 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
Aaron Enye Shi 23e9968752 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
Maneesh Gupta 352b17346c 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
Alex Voicu ea0fcf3e61 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
Rahul Garg 59081c69fc Add extension for kernel concurrency on same stream 2019-03-06 12:55:39 +05:30
Wen-Heng (Jack) Chung b4d658a48f 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
Yaxun Sam Liu 450f093231 Let hip-clang support --genco 2018-11-27 15:55:50 -05:00
Maneesh Gupta 160c509e23 Merge pull request #760 from eshcherb/roctracer-hip-frontend-181113
Roctracer hip frontend 181113
2018-11-23 11:08:25 +05:30
Michael Kuron 8610128c3e Merge branch 'master' into getsymboladdress 2018-11-20 12:03:22 +01:00
Evgeny e5ba097afd renaming HIP_INIT_CB_API to HIP_INIT_API 2018-11-13 15:33:26 +00:00
Evgeny cba2d42bbf adding lost i the merge change 2018-11-13 15:33:26 +00:00
Evgeny b8b1637ef7 adding activity prof layer 2018-11-13 15:33:26 +00:00
Rahul Garg 6b3cbc65ad Fixed symbol tracking device index 2018-11-13 07:01:17 +05:30
Maneesh Gupta e672dc8a55 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
Michael Kuron 31acf1c268 Introduce ihipModuleGetGlobal 2018-11-06 09:54:34 +01:00
Rahul Garg b270313129 Fixes global symbols tracking in hip_module 2018-10-31 03:22:38 +05:30
Qianfeng Zhang de5f47a984 Make correct checking of the returned hipDeviceptr_t from read_global_description() 2018-10-23 21:13:11 +08:00
Alex Voicu cd6c979c27 Update hip_module.cpp
Typo.
2018-05-18 17:50:45 +01:00
Alex Voicu 5325b6535e Update hip_module.cpp 2018-05-14 17:15:36 +01:00
Alex Voicu 1ba8a35dba 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
Alex Voicu 13274ce559 Add support for the hipFuncGetAttributes interface. 2018-05-11 03:35:10 +01:00
Rahul Garg 1446f78799 hip_module code cleanup
-Fixed missing ihipLogStatus in hipModuleLoad()
-Fixed some ihipXXX functions
2018-04-16 15:35:04 +05:30
Maneesh Gupta 1ba06f63c4 Apply .clangformat to all repo source files
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344
2018-03-12 11:29:03 +05:30
Alex Voicu baf50a5311 Re-sync with upstream. 2018-02-12 20:20:24 +00:00
Rahul Garg 24ab820a11 Fixed host allocated globals address lookup for host usage
Fixed texture driver APIs failure
2018-01-30 18:06:31 +05:30
Rahul Garg 8a060d194e Fixed build error 2017-12-28 16:15:45 +05:30
Alex Voicu b842394957 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
Alex Voicu 5127ce67e8 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
Alex Voicu 89e9399427 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
Alex Voicu 02c2bfc7ef Re-sync with upstream and re-factor platform global management for texture references. 2017-11-28 19:15:29 +00:00
Alex Voicu dc67ca3feb 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
Rahul Garg 56862b1c35 Fixed review comments 2017-11-21 21:19:06 +05:30
Alex Voicu 9d088d2283 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
Rahul Garg f97c5f9a64 -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
Rahul Garg c7d60a7a75 Update hipModuleGetTexRef API 2017-11-19 22:10:46 +05:30
Alex Voicu f7726cd416 Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 2017-11-09 23:43:07 +00:00
Rahul Garg ef09c4918d Texture driver APIs support 2017-11-09 22:10:55 +05:30
Alex Voicu 328c18b886 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
Ben Sander 86f62accfd Merge pull request #245 from scchan/centos_fixes
various fixes for centos/rhel
2017-11-01 18:10:29 +01:00
Alex Voicu c2482d1255 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
Siu Chi Chan 99d32a195f 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