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
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
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
c96e9341f7
fixed default args for symbol memcpy apis
...
Change-Id: Ie0b63f8b9c5535eb3946bd6af3f30fe71a015244
[ROCm/hip commit: 522e059a79
2017-06-23 19:05:44 -05:00
ed1f242c60
Clean up old USE_* and RELEASE.md notes.
...
[ROCm/hip commit: 7912e61502
2017-06-23 18:05:30 -05:00
1678be3d74
Added device side abort function in HIP/NVCC
...
Change-Id: I6ae35a72a8b9c34852619f02da1a046c8d3b2ed3
[ROCm/hip commit: e6e4fe613c
2017-06-20 09:38:56 +05:30
3cb1061f65
Abort device function in HIP/HCC, need new HCC
...
Change-Id: I4195ab75e9b7b48c8b8128d6925ddc0fa5e9e009
[ROCm/hip commit: 3c73229916
2017-06-18 12:31:31 +05:30
f5dff37b8b
removed bad copy constructor
...
Change-Id: I661991d9d43941a61848b0b8e9879c0bfa811b40
[ROCm/hip commit: 64bb8d154a
2017-06-16 09:02:26 -05:00
7e9cfda45e
Additional GGL make_kernel_functor_* macros, contributed by Alex
...
Change-Id: I01aabb7d2b5418fcefb1bbf78eb5d1888dbc5c96
[ROCm/hip commit: fd36303c24
2017-06-14 11:10:52 -05:00
47921e39d2
Fix error related to undefined reference of __get_dynamicgroupbaseptr().
...
Change-Id: I14951e1725e35dd5f5e53805f81cdb58661f59f2
[ROCm/hip commit: 43df5ba660
2017-06-08 19:24:32 -05:00
8b14507650
Add clang version guard so the hip_fp16.h header won't be picked up by gcc
...
Change-Id: Ia21335a455bc93210901b44bc8c76a7f4a385b55
[ROCm/hip commit: 06816fb68b
2017-06-08 19:24:32 -05:00
3f9b16c397
Improve HIP kernel names, attributes and codegen, contributed by Alex Voicu
...
Change-Id: I2cafbdc5a98e26c7f4fad84739c915e7dc09993c
[ROCm/hip commit: a9808961bd
2017-06-05 11:39:00 -05:00
445042f916
Fix HIP_SYNC_NULL_STREAM=0 mode.
...
- Fix null-stream sync
- hipStreamDestroy of null stream returns hipErrorInvalidResourceHandle
- Update documentation.
- Add tests for null stream sync, hipEventElapsedTime.
- Rename internal enum hipEventStatusRecorded to hipEventStatusComplete
- refactor hipStreamWaitEvent to streamline control-flow
[ROCm/hip commit: 39c18e5e5f
2017-06-05 08:50:22 -05:00
47ec040140
added half data type and vector destructors
...
1. Added half data types to hip_fp16.h
2. Added destructor to vector data types
Change-Id: Id5ae76a663bb90a4bde2839ec79c58fbaee5072f
[ROCm/hip commit: cb7c4c423c
2017-06-02 11:19:33 -05:00
7a3befc555
[HIP] [HIPIFY] CUDA Driver API 8.0.44 JIT options support.
...
[ROCm/hip commit: 6e99e388ea
2017-05-31 18:55:29 +03:00
b3e01ed74f
Add event controls for release fences.
...
Env var : HIP_EVENT_SYS_RELEASE
Event allocation flags : hipEventReleaseToDevice, hipEventReleaseToSystem
(remove hipEventDisableSystemRelease)
Update test for new functionality.
[ROCm/hip commit: ec7102f461
2017-05-27 16:02:34 -05:00
Alex Voicu
Ben Sander
Siu Chi Chan
Chris Kitching
Maneesh Gupta
Rahul Garg
sriharikarnam
wsttiger
Weixing Zhang
Wen-Heng (Jack) Chung
Sun, Peng
Aditya Atluri
Evgeny Mankov