ce9ca744dc
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*.
2018-07-04 20:52:55 +00:00
f8109c5b06
Merge branch 'master' into hipclang-add-amdgcn-funcs
2018-07-04 17:36:08 +00:00
e54ffe52a0
Merge pull request #530 from ROCm-Developer-Tools/hipclang-compute-mode
...
Add HIP Compute Mode
2018-07-04 10:55:48 +05:30
1340adce1a
Merge pull request #503 from ROCm-Developer-Tools/fix-complex
...
Fix hip_complex.h
2018-06-28 12:15:22 +05:30
b6f5c61a2c
Merge pull request #495 from ROCm-Developer-Tools/hip-clang-upstream
...
add more device functions for hip-clang
2018-06-28 11:42:01 +05:30
f426491f17
Fix typo
2018-06-28 11:19:22 +05:30
06c1c93945
Merge pull request #536 from ROCm-Developer-Tools/hotfix_remove_unimplemented_math
...
Removes use of unimplemented OCML functionality.
2018-06-28 11:15:57 +05:30
e44eaa1a1e
Implement __shfl_* funcs into HIP headers
2018-06-26 18:32:11 +00:00
c74682ef1b
Include host_defines.h in hip_fp16.h since it uses __host__ __device__ attributes
2018-06-25 15:34:34 -04:00
9452298fbe
Removes use of unimplemented OCML functionality.
2018-06-25 19:16:27 +01:00
c73cd64f5b
Add HIP Compute Mode
2018-06-22 14:37:19 +00:00
5d4986d470
Replace __hip_hc_ir_ inline asm with __ockl_* functions
2018-06-20 20:40:14 +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
e8a18a5473
Move hipclang funcs into corresponding headers
2018-06-20 17:39:15 +00:00
366207f8af
Split __llvm and device lib funcs into new headers
2018-06-20 17:39:06 +00:00
0694d0300a
Add get_dynamicgroupbaseptr def and remove hc_
2018-06-20 17:38:59 +00:00
7cade79d98
Add prefix __ to memory scope and order
2018-06-20 17:38:52 +00:00
4298ed308e
Add __llvm_fence funcs and __ prefixes
2018-06-20 17:38:45 +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
272c16d34e
Merge pull request #507 from ROCm-Developer-Tools/fix-forward
...
Add __device__ to device functions in hip_fp16_math_fwd.h
2018-06-20 14:21:46 +05:30
29aa297899
Merge pull request #504 from ROCm-Developer-Tools/fix-vector3
...
Fix channel_descriptor.h about vector 3 for gcc
2018-06-20 14:20:29 +05:30
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
2b5242b2c3
Add conj, operator-,==,!= for hipFloatComplex/hipDoubleComplex
2018-06-19 10:49:59 -04:00
0de3ef7b32
Add abs/real/imag functions for hipFloatComplex/hipDoubleComplex
2018-06-18 11:57:57 -04:00
40d2d8a0bd
Add missing __device__ __host__ to complex constructor
...
Also add missing typedef value_type
2018-06-17 20:20:32 -04:00
5d0d2ea455
Includes <cmath> or <math.h> by __cplusplus in hip_complex.h
2018-06-17 20:20:32 -04:00
cc44a3d234
Include cmath instead of math.h in hip_complex.h
2018-06-17 20:20:32 -04:00
8ddadf835d
Add missing macro MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT to hip_complex.h
2018-06-17 20:20:32 -04:00
6815c7bdd7
added missing hipCmul() to nvcc_detail/hip_complex.h
2018-06-14 21:49:54 +02:00
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
181cde1899
Merge pull request #506 from ROCm-Developer-Tools/fix-extern-shared
...
Add support of extern __shared__ for hip-clang
2018-06-11 11:59:58 +05:30
d137271083
Revert "Switch over to using native vector types, for better codegen. Remove noise."
2018-06-08 16:48:22 -04:00
9566616dd2
Add __device__ to device functions in hip_fp16_math_fwd.h
2018-06-08 11:23:52 -04:00
93fbbfc8ad
Add more function declarations for hip-clang
2018-06-08 11:20:42 -04:00
6215587197
Fix channel_descriptor.h about vector 3 for gcc
2018-06-08 11:18:41 -04:00
9398c9c927
Add support of extern __shared__ for hip-clang
2018-06-08 11:17:25 -04:00
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
a1ad2b9c65
Merge pull request #496 from gargrahul/add_gettexresdesc_nvcc
...
Add getTextureResourceDescriptor on NVCC
2018-06-06 15:12:11 +05:30
a432847b1b
Merge pull request #487 from gargrahul/fix_hiparray_alloc_flag_nvcc
...
Map hipArray alloc flags on NVCC
2018-06-06 15:11:40 +05:30
6f8bcf53b0
Add getTextureResourceDescriptor on NVCC
2018-06-05 18:46:25 +05:30
92c106bcf5
Fix hideous typos.
2018-06-03 03:03:55 +01:00
13a1917741
Add missing __device__ for forward declares.
2018-06-02 17:46:37 +01:00
513ba4eb8a
Switch the atomic implementation to use Clang builtins.
2018-06-02 12:27:17 +01:00
94f086e9cd
Add integrated device property
2018-06-02 13:11:16 +05:30
3a17e2ad06
Rename for minimal confusion.
2018-06-01 22:55:33 +01:00
b72d82f982
Missing __device__.
2018-06-01 19:48:36 +01:00
91d9ec75d7
Fix typos / address review comments.
2018-06-01 16:20:21 +01:00
Aaron Enye Shi
Maneesh Gupta
Yaxun (Sam) Liu
Alex Voicu
Nico Trost
Wen-Heng (Jack) Chung
Siu Chi Chan
Rahul Garg