17453a6278
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/hip commit: 76f86ef097
2018-07-05 20:38:46 +00:00
27e2af1773
Merge pull request #490 from ROCm-Developer-Tools/feature_decouple_atomics_from_hc
...
Switch the atomic implementation to use Clang builtins.
[ROCm/hip commit: 946c8da88a
2018-06-20 14:16:43 +05:30
3792417463
Revert "Switch over to using native vector types, for better codegen. Remove noise."
...
[ROCm/hip commit: 7a4aace13d
2018-06-08 16:48:22 -04:00
f7fd20ec17
Switch the atomic implementation to use Clang builtins.
...
[ROCm/hip commit: 089ab3b947
2018-06-02 12:27:17 +01:00
b9bf931765
Re-sync with upstream.
...
[ROCm/hip commit: 417869821d
2018-06-01 15:49:05 +01:00
121993cae8
Merge pull request #447 from ROCm-Developer-Tools/feature_native_vector_types
...
Switch over to using native vector types, for better codegen. Remove noise.
[ROCm/hip commit: 68c36f9125
2018-06-01 13:58:07 +05:30
63c8aa6fcb
Re-sync with upstream. Add integer abs.
...
[ROCm/hip commit: ab4b2a650b
2018-05-31 16:38:00 +01:00
90e7799bf4
Add missing interop with volatile. Fix unit tests.
...
[ROCm/hip commit: 0108819e2a
2018-05-31 15:27:12 +01:00
1d220f8867
Switch to using ROCDL directly, as opposed to via HC. Add missing bits.
...
[ROCm/hip commit: 59db16fd36
2018-05-31 03:17:26 +01:00
8b039f989c
Move converting constructor from _Float16 under macro guard. Refactor.
...
[ROCm/hip commit: 0eec115bcd
2018-05-25 19:46:41 +01:00
328e946c59
Update hipTestHalf to actually test behaviour. Add missing hipHostfree.
...
[ROCm/hip commit: 9948b5961e
2018-05-24 13:55:30 +01:00
4ceb9cbc09
Missing commit.
...
[ROCm/hip commit: 6f819f226b
2018-05-23 17:57:47 +01:00
bb8918faf7
Disable incomplete unit tests that don't work on nvcc path
...
Change-Id: If5823ec96a3b2497a08c46ab802c5a0158271053
[ROCm/hip commit: 305592d622
2018-05-21 11:35:03 +05:30
3da5fac84b
Merge pull request #444 from aaronenyeshi/vg20-initial
...
initial gfx906 support
[ROCm/hip commit: 67d45164fa
2018-05-18 13:18:07 +05:30
8712744ba0
Switch over to using native vector types, for better codegen. Remove noise.
...
[ROCm/hip commit: b8c98d2357
2018-05-12 00:14:17 +01:00
9a591d6c26
Fix hipMathFunction for gfx906
...
[ROCm/hip commit: 848a24b524
2018-05-11 10:53:07 -04:00
c05e2e2261
initial gfx906 support
...
[ROCm/hip commit: b898049412
2018-05-10 19:28:00 +00:00
923d831d5e
add intrinsics mbcnt_lo, mbcnt_hi, lane_id
...
[ROCm/hip commit: b285145966
2018-05-08 13:43:53 -04:00
43c2724907
Fixing a copy/paste error in my previous checkin
...
[ROCm/hip commit: 76a7d7e374
2018-04-27 12:59:51 +00:00
2b3332ac72
Checkin to add unit tests for fixes in my previous commit
...
This change adds unit tests for
1. __hisinf
2. __hisnan
2. abs(long long)
3. lgamma(double)
[ROCm/hip commit: 4be4cf644f
2018-04-25 17:53:21 +00:00
4f42ee762d
Apply .clangformat to all repo source files
...
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344
[ROCm/hip commit: 1ba06f63c4
2018-03-12 11:29:03 +05:30
75dcb768e0
[nvccWarnings] Fix -Wno-deprecated-declarations in hip_anyall and hip_ballot
...
[ROCm/hip commit: bd60ef9031
2017-12-20 12:05:21 +05:30
f8d610cf61
[nvccWarnings] Fix warnings seen with dtests on nvcc path
...
[ROCm/hip commit: eea7d495c7
2017-12-14 14:10:37 +05:30
8ef230eaaf
Temporarily disable a couple tests pending some HCC work
...
[ROCm/hip commit: 62cbe22d50
2017-12-01 21:46:28 +00:00
8d51eaafb6
Revert "Revert adoption of CUDA indexing in general - this can only work with later versions of the compiler, just like module based dispatch, and thus must be guarded against usage in earlier (e.g. 1.6) versions."
...
This reverts commit 1c5096832e11e7dc6
2017-11-29 21:49:10 +00:00
fcc42f035e
Revert "Revert adoption of CUDA indexing in general - this can only work with later versions of the compiler, just like module based dispatch, and thus must be guarded against usage in earlier (e.g. 1.6) versions."
...
This reverts commit d2fd1f5fbaf729f88
2017-11-29 21:36:29 +00:00
d8db534805
Modify the set component of the memcpy test (unclear why there is a memset component to begin with).
...
[ROCm/hip commit: 4131b47134
2017-11-21 17:52:01 +00:00
3c936777aa
Re-sync with upstream.
...
[ROCm/hip commit: c5f2b22d0d
2017-11-20 15:34:50 +00:00
d807f528e6
Merge pull request #266 from gargrahul/fix_half2_gfx900
...
Fixed half2 issue on gfx900
[ROCm/hip commit: db378fbc9e
2017-11-20 07:28:41 +05:30
962bf7bfda
This implements the trivial change needed to move back from the hip{Something}_{x, y, z} macros to the natural CUDA syntax of Something.{x, y, z}. This is contained in lines 384-404 in hip_runtime.h. All of the other changes have to do with changing unit tests to use this syntax. The macros are retained for backwards compatibility.
...
[ROCm/hip commit: cffd0e14eb
2017-11-19 01:54:12 +00:00
5de28551c1
Fixed test case for GFX900
...
[ROCm/hip commit: 9af0f9cbc1
2017-11-16 09:34:52 +05:30
d0ee6e00fe
Fixed half2 issue on gfx900
...
[ROCm/hip commit: fef496d4f1
2017-11-15 18:52:59 +05:30
b8a81d2b75
This corrects some interesting choices that were present in the HIP
...
unit tests such as e.g. de-allocating memory allocated with new[] using
free. All of these were identified via cppcheck.
[ROCm/hip commit: a619ac6f92
2017-11-01 22:51:22 +00:00
70a41e7dac
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.
...
[ROCm/hip commit: c2482d1255
2017-11-01 15:09:59 +00:00
52a97377ac
[dtests] Fix hip_threadfence_system test's BUILD command
...
Change-Id: I037c523369e4bddd7a386ffe117fe97a178900fd
[ROCm/hip commit: 4bcb3372b5
2017-08-11 14:49:17 +05:30
588ca89c92
add C++11 compilation flags and minor bug fixes
...
[ROCm/hip commit: 9a04c5cdd4
2017-08-10 17:58:36 +00:00
7390368693
[dtests] Disable hip_threadfence_system.tst on nvcc
...
The test needs some more modifications before it is ready for building
on nvcc path. Disabling it for now so that the pull request can be
merged.
Change-Id: I19a894fcda7b2159f86a4e4e95a409c5655d3760
[ROCm/hip commit: 69e431f035
2017-08-10 12:05:58 +05:30
21c1d69e77
implement __threadfence_system
...
[ROCm/hip commit: 96e79f97e2
2017-08-02 08:50:18 +00:00
9ce8abb1ed
Renable frexp(f) device math function
...
Change-Id: I53c022b8ddf38cd17ddb42eba457b9020db66395
[ROCm/hip commit: 7a446dcbc5
2017-07-20 14:41:30 +05:30
298fa00e8c
Enable memcpytosymbol family functions test
...
[ROCm/hip commit: 039f62213b
2017-07-18 12:37:37 +05:30
f28a1365f8
Remove explicit address spaces declarations
...
[ROCm/hip commit: 30ee797dd4
2017-07-07 12:46:29 -05:00
ee92e19e7a
hipDeviceMemcpy test: make it functional on nvcc path
...
Change-Id: Id10c79b48747ed701adbd0a233c53cd60cfa743b
[ROCm/hip commit: e6cafbf342
2017-06-07 15:24:44 +05:30
be2c61b72a
Disable rcbrtf, scalblnf, scalbnf in single precision device test
...
Change-Id: I8a250a64a0cb05132d022a11d9766ced9cdf11a7
[ROCm/hip commit: 5cdd1b2bf5
2017-05-31 10:16:19 +05:30
ed41fcd51a
Disable rcbrt, scalbln and scalbn double precision device test
...
Change-Id: I46bd895701c46d3592b553090cafba99e41a2e2d
[ROCm/hip commit: 4919863d3e
2017-05-31 10:15:41 +05:30
26f6ce992d
added __host__ for complex functions and corrected memset and memcpy test
...
Change-Id: I9ffefb7a0025aa111a54d20d2766982df15532e7
[ROCm/hip commit: 6dff2714e9
2017-04-06 09:29:44 -05:00
39550b0f4c
fix hipVectorTypesDevice direct test with GGL enabled
...
Change-Id: I7a63b87348f08f094cd709e87397d9e0fc24e4c2
[ROCm/hip commit: c865151e50
2017-03-30 17:14:55 -05:00
83ce6f48a9
Update hipTestDeviceSymbol sample
...
Change-Id: If5ba99c60cd30c4491ca3a4856764224163d3ddf
[ROCm/hip commit: ec0d334354
2017-03-24 10:39:11 +05:30
a9ec705ee4
Add USE_PROMOTE_FREE_HCC for smooth transition to new HCC caps.
...
ADDRESS_SPACE_1 defines
[ROCm/hip commit: e85c167171
2017-03-17 11:04:39 -05:00
25e4ace28b
Add simple device-side assert macro
...
Currently swallows asserts but will compile.
[ROCm/hip commit: ecd8179a71
2017-03-17 10:25:56 -05:00
86e2a5d11d
Disable broken tests on hcc path
...
Change-Id: Id6234da576566faa32d5fdf42dca6d6267596823
[ROCm/hip commit: 65bb22eefc
2017-03-15 12:03:44 +05:30
Aaron Enye Shi
Maneesh Gupta
Siu Chi Chan
Alex Voicu
Deven Desai
Phaneendr-kumar Lanka
Phaneendr-kumar Lanka
Ben Sander
Rahul Garg
Wen-Heng (Jack) Chung
Aditya Atluri
Sun, Peng