dd8a589893
Choose whether or not to use functional grid_launch based on the version of HCC used to compile.
...
[ROCm/hip commit: 89e9399427
2017-11-29 00:17:44 +00:00
9668003fe3
Change memset kernel to use memcpy instead of placement new. Simplify indexers.
...
[ROCm/hip commit: 6e4ca3fbb4
2017-11-28 19:45:47 +00:00
aef26d3477
Re-sync with upstream and re-factor platform global management for texture references.
...
[ROCm/hip commit: 02c2bfc7ef
2017-11-28 19:15:29 +00:00
11f7d895f4
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
# Conflicts:
# src/hip_module.cpp
[ROCm/hip commit: dc67ca3feb
2017-11-28 17:29:11 +00:00
d31ef2b462
Merge pull request #256 from gargrahul/texture_driver_api_support
...
Texture driver APIs support
[ROCm/hip commit: 0da0426f94
2017-11-27 13:52:39 -06:00
acc05cdff7
Remove leftover comment.
...
[ROCm/hip commit: 08f252e4bf
2017-11-22 19:37:03 +00:00
13879387e5
Fixed review comments
...
[ROCm/hip commit: 56862b1c35
2017-11-21 21:19:06 +05:30
03552f2e94
Changed function hipMemcpy_2D to hipMemcpyParam2D
...
[ROCm/hip commit: 9866fa250d
2017-11-21 12:36:24 +05:30
45ff3c31c4
Refactor the __device__ versions of memset and memcpy to be less awkward i.e. not return nullptr as opposed to the destination pointer (it can only be assumed it was done for maximum confusion) and actually unroll as they claim to. Change all of the {to, from}Symbol functions to use hipModuleGetGlobal, as opposed to hc::accelerator::get_symbol_address which is no longer valid with module based dispatch.
...
[ROCm/hip commit: 9d088d2283
2017-11-21 02:40:34 +00:00
69f8043d12
Correct ill-formed merge in earlier commit and adjust for differences with the new CUDA natural indexing mechanism.
...
[ROCm/hip commit: 7d5a45ac1a
2017-11-20 16:33:52 +00:00
3c936777aa
Re-sync with upstream.
...
[ROCm/hip commit: c5f2b22d0d
2017-11-20 15:34:50 +00:00
62f20c6dc8
Update hipModuleGetTexRef API
...
[ROCm/hip commit: c7d60a7a75
2017-11-19 22:10:46 +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
58b37dce5f
-Fixed texture driver API sample
...
-Added hipTexRefSetAddress and hipTexRefSetAddress2D APIs
[ROCm/hip commit: 4b19c2aa0c
2017-11-15 18:23:28 +05:30
0263292087
Texture code reorganized
...
[ROCm/hip commit: 63680edd30
2017-11-14 11:09:35 +05:30
49a2c6468e
Texture driver APIs support
...
[ROCm/hip commit: ef09c4918d
2017-11-09 22:10:55 +05:30
35d8af2001
Clean up trailing whitespace so as to reduce noise in #246 .
...
[ROCm/hip commit: d8e323d4b5
2017-11-08 00:08:55 +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
51ee7807db
Merge pull request #222 from bensander/fix_device_prop
...
Fix device prop
[ROCm/hip commit: 2e8ec71e40
2017-10-30 17:58:48 +01:00
6cc7f10e84
Merge remote-tracking branch 'origin/master' into HEAD
...
[ROCm/hip commit: a9789ddcda
2017-10-27 01:18:28 -04:00
ccc193cf98
add a printf device function
...
[ROCm/hip commit: f8bdb9f688
2017-10-23 18:30:08 -04:00
e249e3c30e
Update device properties.
...
- clear properties to defined initial state.
- enable some property flags which are now supported.
[ROCm/hip commit: acf89b43d4
2017-10-20 15:52:13 +00:00
49feb372f5
Add [[nodiscard]] attribute to hipError_t in C++17 mode
...
This will yield a compiler warning any time you discard an error
return value from an API call.
[ROCm/hip commit: 1d205ab130
2017-10-19 13:34:02 +01:00
710d095874
Merge pull request #194 from bensander/fix_genco_math_path
...
Fix math ordering for --genco mode.
[ROCm/hip commit: c1803c1951
2017-10-03 16:17:34 +05:30
d1be93bbd9
Fix math ordering for --genco mode.
...
[ROCm/hip commit: 68c2146dc3
2017-10-02 21:52:16 +00:00
be4b7099bb
Fixed review comment
...
[ROCm/hip commit: b8aae66424
2017-10-02 17:29:08 +05:30
28b92618b7
Fixed review comments
...
[ROCm/hip commit: 19d574d939
2017-10-02 10:57:25 +05:30
71552c63ac
Fixed build issue
...
[ROCm/hip commit: f62254a8d8
2017-10-02 10:40:01 +05:30
69508f52cc
Fix texture support on HIP/NVCC path
...
[ROCm/hip commit: 470224cf25
2017-10-02 10:40:01 +05:30
6884635ae0
ensure the condition is evaluated before applying the negation operator
...
[ROCm/hip commit: 83d9e889a9
2017-09-20 01:38:28 -04:00
ed0f94c051
Merge pull request #155 from wsttiger/master
...
Modify hip headers such that they are more compatible with strict C compiler
[ROCm/hip commit: 8c92d28ff3
2017-09-15 10:49:55 +05:30
569ec04d0c
Fix typo in nvcc_detail/hip_runtime_api.h
...
[ROCm/hip commit: c3c94dcd8d
2017-09-14 16:19:57 +05:30
0ba15f6faf
Merge pull request #175 from gargrahul/assert_support
...
Device side assert support
[ROCm/hip commit: c7e5fd24ac
2017-09-14 16:08:00 +05:30
ba8fd5cbee
Add hipDevice[Get/Set]SharedMemConfig API on nvcc details
...
[ROCm/hip commit: 3e69b2ed5a
2017-09-14 16:03:19 +05:30
d1425f3cf5
Device side assert support
...
[ROCm/hip commit: 40ff42e555
2017-09-07 20:45:27 +05:30
a497059d78
Update docs.
...
- Review unsupported comments and refresh.
- Review experimental tag in API.
[ROCm/hip commit: 1a9e588932
2017-08-28 15:44:38 -05:00
09fb0fce42
Modify header such that it's compatible with strict C compiler
...
[ROCm/hip commit: b6a3d8c946
2017-08-15 13:23:09 -05:00
db3ebbcbae
Add missing CUDA tags
...
[ROCm/hip commit: 512214e930
2017-08-15 13:22:20 -05:00
21c1d69e77
implement __threadfence_system
...
[ROCm/hip commit: 96e79f97e2
2017-08-02 08:50:18 +00:00
9ec4b0b66f
Merge pull request #100 from weixingzhang/texture
...
HIP Texture Support
[ROCm/hip commit: 4f93b99f23
2017-07-26 19:28:15 -05:00
9ce8abb1ed
Renable frexp(f) device math function
...
Change-Id: I53c022b8ddf38cd17ddb42eba457b9020db66395
[ROCm/hip commit: 7a446dcbc5
2017-07-20 14:41:30 +05:30
65bec11393
HIP Texture Support
...
[ROCm/hip commit: e671cec367
2017-07-17 15:16:12 -04:00
f28a1365f8
Remove explicit address spaces declarations
...
[ROCm/hip commit: 30ee797dd4
2017-07-07 12:46:29 -05:00
042ed2c38f
Add hipdbPrintMem - wrapper for hcc memory tracker.
...
[ROCm/hip commit: 0459f3db72
2017-07-03 15:05:30 -05:00
9498754cf6
Fixed build error in hipDeviceGetByPCIBusId
...
Change-Id: I33be362910c4c5dca7435fb9c41f427b54f0e9a7
[ROCm/hip commit: 34fcaa23bb
2017-07-01 07:56:30 +05:30
15527645f1
Refactor hip_hcc.
...
Don't require setting __HIP_PLATFORM_HCC__ ; Instead check if compiler is HCC.
The functions here use HCC-specific functionality so only work with HCC.
ANd they work regardless of the __HIP_PLATFORM_HCC__ setting so don't
require that.
Also remove the "hcc_detail" extra hierarchy level for hip_hcc.h as
it is not necessary. Remove hcc_detail/hip_hcc.h.
[ROCm/hip commit: 9c229f7530
2017-06-30 08:26:15 -05:00
83f3ed16de
Fixed hipDeviceGetByPCIBusId
...
Change-Id: Ia36bb9425671ef7659541c9aeedae4098456a31b
[ROCm/hip commit: 65cb024b61
2017-06-30 10:11:41 +05:30
9e55cdfa90
Fixed hipDeviceGetPCIBusId for HIP/NVCC
...
Change-Id: I662efa148257a710f09002850b41d57bef00dff3
[ROCm/hip commit: d6db525f51
2017-06-29 00:13:02 +05:30
ff48608355
Fixed hipDeviceGetPCIBusId for HIP/HCC
...
Change-Id: I3688fa2476e1baada2d3c5fc3735cec3f15a1e21
[ROCm/hip commit: f4ea3b76c8
2017-06-28 23:48:27 +05:30
5a6d4a30c4
Remove some warning debug info and add weak attribute back to GGL __global__ define
...
Change-Id: I2021b107dda697b1262d44fa1506465e94a3916b
[ROCm/hip commit: 1e1654c225
2017-06-27 14:15:16 -05:00
Alex Voicu
Ben Sander
Rahul Garg
Siu Chi Chan
Chris Kitching
Maneesh Gupta
sriharikarnam
wsttiger
Weixing Zhang
Wen-Heng (Jack) Chung
Sun, Peng