2d5e831c7f
Add declare of __get_dynamicgroupbaseptr for host compilation
...
[ROCm/clr commit: 1f34993ac1
2018-07-09 17:56:25 +00:00
0e5ef747e6
Enable surface functions on HIP clang path
...
Fix surface test on HIP clang path.
[ROCm/clr commit: 5862ae4b6a
2018-07-05 20:49:47 +00:00
eab9cacef2
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.
[ROCm/clr commit: d39508b4fd
2018-07-05 20:38:46 +00:00
94cc27d4d6
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.
[ROCm/clr commit: b362c53475
2018-07-05 20:15:41 +00:00
f78a8d7fd6
Implement Memory Fence Functions in header
...
Enabled __llvm_fence_* functions for seq_cst.
[ROCm/clr commit: 96ab7c7b25
2018-07-04 23:35:24 +00:00
303a8be7d3
Workaround cast warning of smaller integer type for __to_local
...
For now, guard the __to_local function for device compile only since a local pointer should be same size as unsigned int on GPU compile. Also change to void* instead of char*.
[ROCm/clr commit: ce9ca744dc
2018-07-04 20:52:55 +00:00
4b941a0db8
Merge branch 'master' into hipclang-add-amdgcn-funcs
...
[ROCm/clr commit: f8109c5b06
2018-07-04 17:36:08 +00:00
b1b4302bc4
Merge pull request #530 from ROCm-Developer-Tools/hipclang-compute-mode
...
Add HIP Compute Mode
[ROCm/clr commit: e54ffe52a0
2018-07-04 10:55:48 +05:30
aaee068ac0
Merge pull request #503 from ROCm-Developer-Tools/fix-complex
...
Fix hip_complex.h
[ROCm/clr commit: 1340adce1a
2018-06-28 12:15:22 +05:30
3ed562a12b
Merge pull request #495 from ROCm-Developer-Tools/hip-clang-upstream
...
add more device functions for hip-clang
[ROCm/clr commit: b6f5c61a2c
2018-06-28 11:42:01 +05:30
90a6b96c5f
Fix typo
...
[ROCm/clr commit: f426491f17
2018-06-28 11:19:22 +05:30
e33d1d0dda
Merge pull request #536 from ROCm-Developer-Tools/hotfix_remove_unimplemented_math
...
Removes use of unimplemented OCML functionality.
[ROCm/clr commit: 06c1c93945
2018-06-28 11:15:57 +05:30
c60682797a
Implement __shfl_* funcs into HIP headers
...
[ROCm/clr commit: e44eaa1a1e
2018-06-26 18:32:11 +00:00
9d35209f8f
Include host_defines.h in hip_fp16.h since it uses __host__ __device__ attributes
...
[ROCm/clr commit: c74682ef1b
2018-06-25 15:34:34 -04:00
4fcfb64a88
Removes use of unimplemented OCML functionality.
...
[ROCm/clr commit: 9452298fbe
2018-06-25 19:16:27 +01:00
4ad4f6a489
Add HIP Compute Mode
...
[ROCm/clr commit: c73cd64f5b
2018-06-22 14:37:19 +00:00
3f19f0979b
Replace __hip_hc_ir_ inline asm with __ockl_* functions
...
[ROCm/clr commit: 5d4986d470
2018-06-20 20:40:14 +00:00
7cb3cec513
Implement __ballot, __any, __all into HIP headers
...
[ROCm/clr commit: 9a295a732d
2018-06-20 17:39:39 +00:00
d09a2f9237
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.
[ROCm/clr commit: 175c87f2bf
2018-06-20 17:39:31 +00:00
7f2693010b
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.
[ROCm/clr commit: ad5f1dcbf2
2018-06-20 17:39:23 +00:00
8f558387fc
Move hipclang funcs into corresponding headers
...
[ROCm/clr commit: e8a18a5473
2018-06-20 17:39:15 +00:00
133a800e94
Split __llvm and device lib funcs into new headers
...
[ROCm/clr commit: 366207f8af
2018-06-20 17:39:06 +00:00
a769701c15
Add get_dynamicgroupbaseptr def and remove hc_
...
[ROCm/clr commit: 0694d0300a
2018-06-20 17:38:59 +00:00
69bdb11cf1
Add prefix __ to memory scope and order
...
[ROCm/clr commit: 7cade79d98
2018-06-20 17:38:52 +00:00
9020ef1b53
Add __llvm_fence funcs and __ prefixes
...
[ROCm/clr commit: 4298ed308e
2018-06-20 17:38:45 +00:00
1e6f00bf67
Add hipclang amdgcn functions
...
These are moving from hipclang in device library to hip headers. These are required for the functionality of HIPclang project.
[ROCm/clr commit: f4399d57a3
2018-06-20 17:38:37 +00:00
8e2a1e0acd
Merge pull request #507 from ROCm-Developer-Tools/fix-forward
...
Add __device__ to device functions in hip_fp16_math_fwd.h
[ROCm/clr commit: 272c16d34e
2018-06-20 14:21:46 +05:30
f1571c9750
Merge pull request #504 from ROCm-Developer-Tools/fix-vector3
...
Fix channel_descriptor.h about vector 3 for gcc
[ROCm/clr commit: 29aa297899
2018-06-20 14:20:29 +05:30
bdc67ac1b4
Merge pull request #490 from ROCm-Developer-Tools/feature_decouple_atomics_from_hc
...
Switch the atomic implementation to use Clang builtins.
[ROCm/clr commit: d27991e1c8
2018-06-20 14:16:43 +05:30
721a3f60be
Merge pull request #457 from whchung/hip-reinit
...
HIP program state re-initialization logic
[ROCm/clr commit: 8ff33a83b1
2018-06-20 09:37:27 +05:30
b93308449e
Add conj, operator-,==,!= for hipFloatComplex/hipDoubleComplex
...
[ROCm/clr commit: 2b5242b2c3
2018-06-19 10:49:59 -04:00
1709777f1b
Add abs/real/imag functions for hipFloatComplex/hipDoubleComplex
...
[ROCm/clr commit: 0de3ef7b32
2018-06-18 11:57:57 -04:00
37fe9e0346
Add missing __device__ __host__ to complex constructor
...
Also add missing typedef value_type
[ROCm/clr commit: 40d2d8a0bd
2018-06-17 20:20:32 -04:00
2dfbd89f5c
Includes <cmath> or <math.h> by __cplusplus in hip_complex.h
...
[ROCm/clr commit: 5d0d2ea455
2018-06-17 20:20:32 -04:00
92959fa686
Include cmath instead of math.h in hip_complex.h
...
[ROCm/clr commit: cc44a3d234
2018-06-17 20:20:32 -04:00
0163a94106
Add missing macro MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT to hip_complex.h
...
[ROCm/clr commit: 8ddadf835d
2018-06-17 20:20:32 -04:00
3eb353ec44
added missing hipCmul() to nvcc_detail/hip_complex.h
...
[ROCm/clr commit: 6815c7bdd7
2018-06-14 21:49:54 +02:00
865fe01b37
HIP program state re-initialization logic
...
This commit is to support kernels dynamically loaded thru means such as
dlopen() after HIP runtime initializes.
[ROCm/clr commit: 04640992dc
2018-06-14 15:46:49 +00:00
ad3c3314bc
Merge pull request #506 from ROCm-Developer-Tools/fix-extern-shared
...
Add support of extern __shared__ for hip-clang
[ROCm/clr commit: 181cde1899
2018-06-11 11:59:58 +05:30
c1cd1513aa
Revert "Switch over to using native vector types, for better codegen. Remove noise."
...
[ROCm/clr commit: d137271083
2018-06-08 16:48:22 -04:00
7f5b121990
Add __device__ to device functions in hip_fp16_math_fwd.h
...
[ROCm/clr commit: 9566616dd2
2018-06-08 11:23:52 -04:00
3536fbeb2b
Add more function declarations for hip-clang
...
[ROCm/clr commit: 93fbbfc8ad
2018-06-08 11:20:42 -04:00
e25084dd6d
Fix channel_descriptor.h about vector 3 for gcc
...
[ROCm/clr commit: 6215587197
2018-06-08 11:18:41 -04:00
ef8c4b51ae
Add support of extern __shared__ for hip-clang
...
[ROCm/clr commit: 9398c9c927
2018-06-08 11:17:25 -04:00
236413158d
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.
[ROCm/clr commit: cb642f14ab
2018-06-06 16:07:22 +05:30
95907c967d
Merge pull request #496 from gargrahul/add_gettexresdesc_nvcc
...
Add getTextureResourceDescriptor on NVCC
[ROCm/clr commit: a1ad2b9c65
2018-06-06 15:12:11 +05:30
c4b344dab4
Merge pull request #487 from gargrahul/fix_hiparray_alloc_flag_nvcc
...
Map hipArray alloc flags on NVCC
[ROCm/clr commit: a432847b1b
2018-06-06 15:11:40 +05:30
8a88d182f5
Add getTextureResourceDescriptor on NVCC
...
[ROCm/clr commit: 6f8bcf53b0
2018-06-05 18:46:25 +05:30
d45814f84d
Fix hideous typos.
...
[ROCm/clr commit: 92c106bcf5
2018-06-03 03:03:55 +01:00
6acd8d4cf2
Add missing __device__ for forward declares.
...
[ROCm/clr commit: 13a1917741
2018-06-02 17:46:37 +01:00
Yaxun (Sam) Liu
Aaron Enye Shi
Maneesh Gupta
Alex Voicu
Nico Trost
Wen-Heng (Jack) Chung
Siu Chi Chan
Rahul Garg