Commit graph

923 Commits

Autor SHA1 Nachricht Datum
Maneesh Gupta 36d37296d5 Merge pull request #1081 from mangupta/swdev-181624
Implement hipExtGetLinkTypeAndHopCount for ROCm devices

[ROCm/hip commit: d5abe65668]
2019-05-07 16:15:41 +05:30
wkwchau 236834c679 Return hipErrorInsufficientDriver status when CPU device not found (#1064)
* Return hipErrorInsufficientDriver status when CPU device not found - no exception thrown

* Return hipErrorInsufficientDriver status when CPU device not found


[ROCm/hip commit: 29b3b46b42]
2019-05-07 15:58:25 +05:30
Maneesh Gupta f7eaff474b Implement hipExtGetLinkTypeAndHopCount for ROCm devices
Change-Id: Ie5bb4f640ac6d189c7fceeab22627a7494fd10bd


[ROCm/hip commit: 035ef04e1f]
2019-05-06 15:54:31 +05:30
Sameer Sahasrabuddhe c3db7b929a minor cleanup: eliminate repetition
[ROCm/hip commit: abb9375707]
2019-04-25 20:41:16 +05:30
Rahul Garg c01236f679 Add hipMallocManaged default functional support (#1036)
* Add hipMallocManaged default functional support

* Fix build error

* Add dtest


[ROCm/hip commit: 2bc2c46d4d]
2019-04-24 16:50:03 +05:30
Yaxun (Sam) Liu d4bce6c019 Fix missing arg in HIP_INIT_API
[ROCm/hip commit: bb5c620b13]
2019-04-18 16:18:31 -04:00
Maneesh Gupta 8b7cd90f0c Merge pull request #1019 from scchan/lazy_binding
minor workaround for lazy binding

[ROCm/hip commit: d789aef46e]
2019-04-16 08:36:10 +05:30
Jeff Daily a0172ca884 In hipFree, synchronize owner of memory (#1018)
* In hipFree, if memory is associated with a device, synchronize that device's streams.

This changes the behavior from synchronizing the currently set TLS device.

* All devices sync in hipFree for _appId=-1 case.

* Revert "All devices sync in hipFree for _appId=-1 case."

This reverts commit 1efb34d6a8426661e45bc5f763422a1147aeac10.

* add HIP_SYNC_FREE env var


[ROCm/hip commit: 2b3037a6ea]
2019-04-16 08:35:55 +05:30
Yaxun (Sam) Liu cf4bdb8b55 Fix regression on multi-gpu due to PR#997
[ROCm/hip commit: 271fdc4e4d]
2019-04-05 22:54:41 -04:00
Siu Chi Chan 23df3b5879 minor workaround for lazy binding
[ROCm/hip commit: e126f7254c]
2019-04-02 17:28:06 -04:00
Yaxun Sam Liu 5072c98f32 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.


[ROCm/hip commit: 98b9e92908]
2019-03-31 10:08:20 -04:00
Wen-Heng (Jack) Chung 1cc94f9369 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


[ROCm/hip commit: 4b7177ac42]
2019-03-29 03:45:04 +00:00
Maneesh Gupta f63c4be10e Merge pull request #992 from gargrahul/handle_d2d_memcpy2d
Handle D2D in memcpy2D

[ROCm/hip commit: d0e5fbeb72]
2019-03-28 04:41:36 +00:00
Rahul Garg 0eaa29ad06 Handle D2D in memcpy2D
[ROCm/hip commit: 0c55db8552]
2019-03-28 02:21:45 +05:30
Jeff Daily 21f5e142f5 load program state once per agent
[ROCm/hip commit: c9117de8eb]
2019-03-27 18:19:10 +00:00
Rahul Garg d98d5ca12a Let hipHostMalloc always share/map pinned host ptr
[ROCm/hip commit: f0af073793]
2019-03-26 10:19:13 +05:30
Rahul Garg c6ef785464 Avoid double mapping of devices to hostMalloc buffer
[ROCm/hip commit: 5e917d70f3]
2019-03-25 23:07:05 +05:30
Maneesh Gupta 82fd86e63f Merge pull request #970 from mangupta/swdev-172995
hipExtMallocWithFlags implementation

[ROCm/hip commit: 30b5c02ec4]
2019-03-25 07:46:53 +00:00
Maneesh Gupta 8c6b5bf266 Merge pull request #962 from gargrahul/add_2d_copy_fallback
Add 2D fallback to use copy kernel

[ROCm/hip commit: eb03d50de9]
2019-03-25 07:46:43 +00:00
Rahul Garg 63e926cec7 2D Fallback needs hcc workweek 19101 or higher
[ROCm/hip commit: 9bbfbceb64]
2019-03-25 12:07:28 +05:30
Maneesh Gupta 67819c0395 hipExtMallocWithFlags needs hcc workweek 19115 or higher
[ROCm/hip commit: cab119c8b2]
2019-03-25 11:41:20 +05:30
Siu Chi Chan 99c4f43a6c reimplement HIP_INIT as hip_impl::hip_init(), add hip_init() to some of the inlined API (#966)
* reimplement HIP_INIT as a function, expose it as hip_impl::hip_init()
so that it could be called from hipLaunchKernelGGL and other inlined
HIP functions

* Don't call hip_init from ihipPreLaunchKernel


[ROCm/hip commit: 24d08beef8]
2019-03-20 05:11:15 +00:00
Maneesh Gupta 9ac6005d35 hipExtMallocWithFlags implementation
Change-Id: Iee9e119796472200b2933d5e23be60813f33bc75


[ROCm/hip commit: 73ec5d54b5]
2019-03-19 11:59:22 +05:30
Evgeny 36b5313d65 tracing callback layer update
[ROCm/hip commit: 31475c5ac8]
2019-03-14 22:43:52 -05:00
Rahul Garg a3fb908a0a Add 2D fallback to use copy kernel
[ROCm/hip commit: 918d7e3a40]
2019-03-14 13:03:06 +05:30
Siu Chi Chan fc9af438a4 move triple_to_hsa_isa into the header
[ROCm/hip commit: 824ee1aa72]
2019-03-11 19:51:44 +00:00
Siu Chi Chan 15061ddfcc Fix memory leak introduced by previous change to Agent_global.
Make Agent_global manage the lifetime of the name string


[ROCm/hip commit: bf1d48bf78]
2019-03-11 19:51:32 +00:00
Aaron Enye Shi 1e07be3ab3 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/hip commit: 00d24d254d]
2019-03-11 19:51:25 +00:00
Aaron Enye Shi 4b87bd25e8 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/hip commit: 23e9968752]
2019-03-11 19:51:18 +00:00
Maneesh Gupta 3f5e937afc Merge pull request #949 from gargrahul/single_stream_concurrent_kernels
Add extension for kernel concurrency on same stream

[ROCm/hip commit: 352b17346c]
2019-03-06 17:34:54 +05:30
Alex Voicu 0c16497abd 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/hip commit: ea0fcf3e61]
2019-03-06 17:31:44 +05:30
Rahul Garg 5d6fd17fbe Add extension for kernel concurrency on same stream
[ROCm/hip commit: 59081c69fc]
2019-03-06 12:55:39 +05:30
Wen-Heng (Jack) Chung da589e38ed Address code review comments to use hipDeviceptr_t
[ROCm/hip commit: 5cbd28f29b]
2019-03-05 05:51:05 +00:00
Wen-Heng (Jack) Chung 0b7f38d100 Add hipMemsetD32 and hipMemsetD32Async
Add 2 extra memset functions which fills memory with integer-typed data

Also change the parameters of ihipMemset to better explain the semantic


[ROCm/hip commit: 7ebbbd3525]
2019-03-04 17:00:33 +00:00
Wilkin Chau 4a0d68ba3f Fix hipMemset3D test
Calculate the allocated size based on the width, height and depth.


[ROCm/hip commit: 8d92d1ebd7]
2019-02-28 22:42:46 +00:00
Yaxun Sam Liu 7ef8b1138f Add __gnu_h2f_ieee and __gnu_f2h_ieee
The implementation is copied from HCC runtime.

For hcc it has no effect since apps can find them in either hcc runtime or HIP
runtime.

hip-clang needs it in HIP/HCC runtime so that HIP/HCC and HIP/VDI runtime are
swappable.


[ROCm/hip commit: c5e813f64c]
2019-02-21 12:48:28 -05:00
lmoriche a9d82ed57b Merge pull request #899 from lmoriche/hip_clang
Add code-object-v3 support for the HIP-Clang path.

[ROCm/hip commit: 005e09f5b9]
2019-02-06 19:50:51 -08:00
Laurent Morichetti 1bbb9d455b Add code-object-v3 support for the HIP-Clang path.
[ROCm/hip commit: 6ab8711f02]
2019-02-05 14:43:26 -08:00
Rahul Garg 089077a8b7 Revert "Fixed issue of GPU device losing access to host pinned memory"
[ROCm/hip commit: 3120db15e4]
2019-02-04 22:45:13 -08:00
Rahul Garg 7aafa0da9f Fixed issue of GPU device losing access to host pinned memory
[ROCm/hip commit: bab3a94b33]
2019-01-22 06:25:43 +05:30
Evgeny 3f7ff3450e fixing HSA_INIT_API cid args
[ROCm/hip commit: 0164464bcc]
2019-01-16 23:45:44 -06:00
Alex Voicu e2911291dd Ensure that static structures are populated
[ROCm/hip commit: 0b57f50855]
2019-01-09 17:21:53 +00:00
Wen-Heng (Jack) Chung 2f9021afd3 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/hip commit: b4d658a48f]
2019-01-08 17:33:40 +00:00
Alex Voicu 587745b841 Hook into the creaky lazy-reinit machinery. Try to minimise race-risk.
[ROCm/hip commit: ec14daa7ce]
2019-01-01 11:01:16 +05:30
Alex Voicu e127990e23 More blobs, more problems.
[ROCm/hip commit: 340674ceb6]
2019-01-01 11:01:07 +05:30
Alex Voicu 7bcb83a05f Start re-working 731 for 2.0.
[ROCm/hip commit: 25c7e5d609]
2019-01-01 11:00:57 +05:30
Maneesh Gupta 3cf96f31d0 Merge pull request #797 from gargrahul/fixhipPointerGetAttributes
Fixed hipPointerGetAttributes for hostmalloced ptr

[ROCm/hip commit: 56ce3e37d5]
2018-12-12 10:16:07 +05:30
Maneesh Gupta 07dcdff9e5 Merge pull request #608 from gargrahul/add_pinned_2d_sdma_copy
Added support for pinned 2D SDMA copy

[ROCm/hip commit: 0dd26b4f63]
2018-12-12 07:44:16 +05:30
Rahul Garg b304ff5210 Fixed hipPointerGetAttributes for hostmalloced ptr
[ROCm/hip commit: 5f12067708]
2018-12-08 01:42:08 +05:30
Siu Chi Chan bca8022aa0 Revert "Rely on code object metadat for kernarg arguments alignof and sizeof."
This reverts commit 6d64f5e112.


[ROCm/hip commit: 1fbf639962]
2018-11-29 11:38:37 -05:00