633ce36ea7
It is unclear what I was thinking when authoring the original code...
2018-07-17 14:04:57 +01:00
1fc3ef3cd4
Merge branch 'master' into move-memcpy
2018-07-17 10:51:42 +05:30
7d8c1ae761
Merge pull request #515 from ROCm-Developer-Tools/hipclang-add-amdgcn-funcs
...
Add hipclang amdgcn functions
2018-07-17 09:25:09 +05:30
4702a6c86e
Merge pull request #562 from ROCm-Developer-Tools/fix-build-failure
...
Fix build failure in code_object_bundle.cpp
2018-07-12 07:49:41 +05:30
d63db71b61
Merge pull request #546 from gargrahul/fix_bindtex_offset_null_check
...
Fixed offset null check in bind texture functions
2018-07-11 12:52:31 +05:30
404bd43a0e
Merge pull request #545 from ROCm-Developer-Tools/revert-521-temp_fixmemcpy2dasync_trsmissue
...
Revert "Use memcpy kernel for all pinned memory cases in hipMemcpy2DAsync"
2018-07-11 12:52:09 +05:30
903dfa94eb
Move __hip_hc_memcpy and __hip_hc_memset from device_utils.cpp to device_functions.h as inline functions
2018-07-10 18:12:41 -04:00
e6d155bcd5
Fix build failure in code_object_bundle.cpp
2018-07-10 16:49:59 -04:00
d39508b4fd
Implement hip_ldg Functions into HIP header
...
Move all the function definitions for hip_ldg.cpp into hip_ldg.h header and enable for HIP clang path.
2018-07-05 20:38:46 +00:00
b362c53475
Implement min/max functions in HIP header
...
Remove using hc::precise_math min and max. Instead we can use ocml directly for device and std:: for host.
2018-07-05 20:15:41 +00:00
96ab7c7b25
Implement Memory Fence Functions in header
...
Enabled __llvm_fence_* functions for seq_cst.
2018-07-04 23:35:24 +00:00
f8109c5b06
Merge branch 'master' into hipclang-add-amdgcn-funcs
2018-07-04 17:36:08 +00:00
d372d0eaad
Update hip_hcc_internal.h
...
Adding missing include for hip_hcc_internal in order to build with HCC
2018-07-04 09:33:51 +05:30
77cbd4476d
Fixed offset null check in bind texture functions
2018-07-03 08:54:17 +05:30
f554e48db3
Revert "Use memcpy kernel for all pinned memory cases in hipMemcpy2DAsync"
2018-07-02 14:32:11 +05:30
e44eaa1a1e
Implement __shfl_* funcs into HIP headers
2018-06-26 18:32:11 +00:00
9a295a732d
Implement __ballot, __any, __all into HIP headers
2018-06-20 17:39:39 +00:00
175c87f2bf
Implement hip_hc.ll into HIP headers
...
Move all __hip_hc_ir_* functions from hip_hc.ll into HIP header as inline asm. Remove hip_hc.ll and build dependencies from HIP.
2018-06-20 17:39:31 +00:00
ad5f1dcbf2
Implement device_functions.cpp into HIP headers
...
Move all Integer Intrinsics, device_functions.cpp definitions and HIP specific device functions into HIP headers. Implement the device functions using llvm_intrinsics and device-libs functions instead of calling hc::__* functions. Remove device_functions.cpp since everything is now defined in header.
2018-06-20 17:39:23 +00:00
f4399d57a3
Add hipclang amdgcn functions
...
These are moving from hipclang in device library to hip headers. These are required for the functionality of HIPclang project.
2018-06-20 17:38:37 +00:00
d27991e1c8
Merge pull request #490 from ROCm-Developer-Tools/feature_decouple_atomics_from_hc
...
Switch the atomic implementation to use Clang builtins.
2018-06-20 14:16:43 +05:30
8ff33a83b1
Merge pull request #457 from whchung/hip-reinit
...
HIP program state re-initialization logic
2018-06-20 09:37:27 +05:30
8f521edff1
Keep the map which tracks GPU kernel symbols to grow monotonically
2018-06-18 16:54:18 -05:00
b883ea759d
Improve performance of re-initialization logic
...
Keep track of shared libaries already discovered. Do not build HSA executables
for them.
2018-06-15 18:07:33 -05:00
007e2a4b5f
TEMP- fix memcpy2dAsync for trsm issue
2018-06-15 16:08:29 +05:30
04640992dc
HIP program state re-initialization logic
...
This commit is to support kernels dynamically loaded thru means such as
dlopen() after HIP runtime initializes.
2018-06-14 15:46:49 +00:00
2ae3be9773
Fix stream resolution in memcpy2dasync
2018-06-14 11:58:56 +05:30
68554e155b
Fix retrieved locked ptr offset
2018-06-13 23:10:05 +05:30
cb642f14ab
Merge pull request #482 from ROCm-Developer-Tools/feature_clean_up_hip_math
...
Switch to using ROCDL directly, as opposed to via HC. Add missing bits.
2018-06-06 16:07:22 +05:30
53037472ff
Merge pull request #497 from gargrahul/fix_memcpy3d_fastpath
...
Fix hipMemcpy3D for fast path
2018-06-06 14:44:02 +05:30
28c5b15d88
Merge pull request #492 from gargrahul/fix_depth_3d_alloc
...
Fix depth value for 3D allocations
2018-06-06 14:41:23 +05:30
f7c49dde38
Merge pull request #491 from scchan/fix_wait
...
callback handling: don't need to wait for the thread to become ready
2018-06-06 14:38:25 +05:30
163d4a5b03
Fix hipMemcpy3D for fast path
2018-06-05 18:54:33 +05:30
0d719c514f
remove the _ready flag in ihipStreamCallback_t and the mutex that protects it.
2018-06-04 17:29:04 -04:00
fa6ce7a724
Fix depth value for 3D allocations
2018-06-04 18:00:22 +05:30
e21e6ed3a0
callback handler: don't need to wait for the thread to become ready
2018-06-02 17:55:37 -04:00
513ba4eb8a
Switch the atomic implementation to use Clang builtins.
2018-06-02 12:27:17 +01:00
68a0dd826d
Remove vestigial implementations.
2018-06-02 11:37:08 +01:00
94f086e9cd
Add integrated device property
2018-06-02 13:11:16 +05:30
f2d7f112ab
Re-sync with upstream.
2018-06-01 15:49:05 +01:00
095d4dd91e
Merge pull request #484 from gargrahul/fix_malloc_hiphostreg
...
Fix memcpy2D for malloc+ hostRegister
2018-06-01 16:53:25 +05:30
8ecb3eeb55
Merge pull request #466 from ROCm-Developer-Tools/feature_use_Float16
...
Feature use _Float16 and match CUDA __half behaviour.
2018-06-01 13:50:12 +05:30
e03ca1a72e
Re-sync with upstream. Add integer abs.
2018-05-31 16:38:00 +01:00
a3609eaf61
Fix memcpy2D for malloc+ hostRegister
2018-05-31 13:14:27 +05:30
14e6a04387
Switch to using ROCDL directly, as opposed to via HC. Add missing bits.
2018-05-31 03:17:26 +01:00
3327a9a2de
Merge pull request #472 from Jorghi12/patch-3
...
Adding double/long int signatures for abs
2018-05-30 08:32:14 +05:30
61ff40a1cf
Update math_functions.cpp
...
CUDA also has a function named labs.
2018-05-26 16:21:14 -04:00
13f37d550f
Adding double/long int signatures for abs
...
Adding overloads for abs that are found in cuda's math_functions.
2018-05-26 00:40:14 -04:00
27e0566c3a
Use 64x4 grid dims
2018-05-24 23:51:52 +05:30
aed2653857
Clean up and fix remaining bytes copy
2018-05-24 23:30:27 +05:30
Alex Voicu
Maneesh Gupta
Yaxun (Sam) Liu
Aaron Enye Shi
Rahul Garg
Rahul Garg
Wen-Heng (Jack) Chung
Siu Chi Chan
Jorghi12