e85f49a57e
Remove leftover agent allocated globals.
...
[ROCm/clr commit: 45a9f4f7b1
2017-11-28 19:56:04 +00:00
abf92dd6e7
Change memset kernel to use memcpy instead of placement new. Simplify indexers.
...
[ROCm/clr commit: d5c8de3f41
2017-11-28 19:45:47 +00:00
914acbe88f
Re-sync with upstream and re-factor platform global management for texture references.
...
[ROCm/clr commit: 6a0efb7ed2
2017-11-28 19:15:29 +00:00
579a3187da
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
# Conflicts:
# src/hip_module.cpp
[ROCm/clr commit: d37a5a6008
2017-11-28 17:29:11 +00:00
621cc5e5ca
Fix float2int rounding functions
...
Change-Id: I67943859a6344c5eec0eaa23418c9b802ef72468
[ROCm/clr commit: cc418da654
2017-11-28 17:23:43 +00:00
a59daf78d6
This actually (tries) to do the right thing all the way, by using memcpy for bitcasting, and not rely on undefined behaviour of a different flavour as a substitute for the original undefined behaviour. Note that the compiler will (should) optimise down to the same emitted code, since this is a pattern it understands.
...
[ROCm/clr commit: bb35299560
2017-11-28 17:23:06 +00:00
0de6e0be2b
This fixes some outright quaint choices made when implementing HIP's bitwise conversion functions, by using simple reinterpret_casts, as is idiomatic. These functions are supposed to be re-entrant, correct and efficient. Sadly, they were neither: they hid a massive race condition against a value stored in global memory, which means that they were also unreasonably slow if they ever managed to be correct, and relied on union based type punning which is in a grey area of the standard. It is difficult to ascertain what may have been the reason for coming up with this quirky solution.
...
[ROCm/clr commit: 4b7cb0624e
2017-11-28 17:23:06 +00:00
8042277672
Add hipify mappings for all CUDA headers that have HIP equivalents
...
I'm particularly running into issues with `device_types.h` in real
CUDA code...
[ROCm/clr commit: 7761f63bbc
2017-11-28 17:23:06 +00:00
0de96b6e3e
Merge pull request #256 from gargrahul/texture_driver_api_support
...
Texture driver APIs support
[ROCm/clr commit: e93a24bdbe
2017-11-27 13:52:39 -06:00
37ceb9a1a8
Added note in Release notes for texture driver APIs
...
[ROCm/clr commit: 6c18df987d
2017-11-27 22:06:50 +05:30
2df33086a8
Merge pull request #262 from ChrisKitching/frontendaction
...
[HIPIFY] Mostly fix preprocessor-or-template induced issues
[ROCm/clr commit: b25d199111
2017-11-27 17:30:11 +03:00
896c146fef
Porting guides update for texture APIs usage
...
[ROCm/clr commit: 3e9a4cfdd1
2017-11-24 12:00:55 +05:30
84802e5d68
Merge pull request #273 from mangupta/swdev-129574
...
Fix float2int rounding functions
[ROCm/clr commit: 77de4b96cf
2017-11-23 12:04:36 -06:00
05963822e7
Fix float2int rounding functions
...
Change-Id: I67943859a6344c5eec0eaa23418c9b802ef72468
[ROCm/clr commit: 822743845c
2017-11-23 09:57:24 +05:30
36951c298d
Remove leftover comment.
...
[ROCm/clr commit: dfa532db98
2017-11-22 19:37:03 +00:00
d7e1419d8b
Merge pull request #263 from ChrisKitching/headers
...
[HIPIFY] Add hipify mappings for all CUDA headers that have HIP equivalents
[ROCm/clr commit: 6bddac1cfa
2017-11-22 21:24:21 +03:00
f9b177c0b3
Modify the set component of the memcpy test (unclear why there is a memset component to begin with).
...
[ROCm/clr commit: 0755f1fc26
2017-11-21 17:52:01 +00:00
212afdf56b
Fixed review comments
...
[ROCm/clr commit: 38029f2849
2017-11-21 21:19:06 +05:30
31deac0768
This corrects how addresses are formed for symbols which reside in shared objects. For this case, the .value component of an ELF symbol holds the offset from the base VA where the shared object was loaded. Thus, to correctly obtain the VA of the object refered by the symbol, we must add the offset to the VA where the shared object is loaded. We were already doing this correctly for symbols denoting functions, but we were incorrect for those denoting objects.
...
[ROCm/clr commit: fb1021cc0a
2017-11-21 13:15:13 +00:00
dc2c56c43d
Changed function hipMemcpy_2D to hipMemcpyParam2D
...
[ROCm/clr commit: 24307fe5c4
2017-11-21 12:36:24 +05:30
1366b2d5ae
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/clr commit: f8c1c1b38e
2017-11-21 02:40:34 +00:00
98ea26e613
Clean-up some remaining noise in program_state.cpp.
...
[ROCm/clr commit: 24f8a93ff7
2017-11-20 22:41:46 +00:00
f27cd5ad33
Correct ill-formed merge in earlier commit and adjust for differences with the new CUDA natural indexing mechanism.
...
[ROCm/clr commit: 27bc487635
2017-11-20 16:33:52 +00:00
9cf73ef515
Re-sync with upstream.
...
[ROCm/clr commit: 30d90dab38
2017-11-20 15:34:50 +00:00
fdb4ea06f7
Merge pull request #264 from pzins/missing_end_marker
...
Fix missing MARKER_END
[ROCm/clr commit: a1e0c96689
2017-11-20 06:08:01 -06:00
9cee9705dc
-Moved coGlobals in hipModule class (takes care of multi module case)
...
-Used mutex scope for updating coGlobals
[ROCm/clr commit: eff09fcd86
2017-11-20 16:23:18 +05:30
02cbb93ec1
Merge pull request #266 from gargrahul/fix_half2_gfx900
...
Fixed half2 issue on gfx900
[ROCm/clr commit: 4477d3d314
2017-11-20 07:28:41 +05:30
9498ff3cae
Merge pull request #265 from phani544/nvccTests
...
[nvccTests]Enabled inline_asm_vadd on nvcc
[ROCm/clr commit: 29c0ab8401
2017-11-20 07:28:29 +05:30
809f575305
Fix test on cuda
...
[ROCm/clr commit: aeadc3f18f
2017-11-19 15:31:02 -06:00
80021d3c68
Merge branch 'feature_natural_indexing' of https://github.com/AlexVlx/HIP
...
[ROCm/clr commit: a43262e699
2017-11-19 15:25:17 -06:00
6b1fb439f6
Temporarily disable P2P on nvidia (fails on dual GPU)
...
[ROCm/clr commit: fc34fd6f03
2017-11-19 15:21:37 -06:00
7453c8b25c
Update hipModuleGetTexRef API
...
[ROCm/clr commit: 76f174b536
2017-11-19 22:10:46 +05:30
af0ae5421c
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/clr commit: 0cc921f103
2017-11-19 01:54:12 +00:00
3b413e7ee7
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
[ROCm/clr commit: 45ac5b70df
2017-11-16 23:20:15 +00:00
b0461a2176
Fixed test case for GFX900
...
[ROCm/clr commit: cfc9f87465
2017-11-16 09:34:52 +05:30
c3ba309562
Fixed half2 issue on gfx900
...
[ROCm/clr commit: 20947f80fb
2017-11-15 18:52:59 +05:30
0f8296b0c2
Removed redundant desc variable
...
[ROCm/clr commit: e2e499461a
2017-11-15 18:28:27 +05:30
a0f838a75e
-Fixed texture driver API sample
...
-Added hipTexRefSetAddress and hipTexRefSetAddress2D APIs
[ROCm/clr commit: edbfeeb141
2017-11-15 18:23:28 +05:30
3bfe30e22c
[nvccTests]Enabled inline_asm_vadd on nvcc
...
[ROCm/clr commit: 9fa3ed4f6f
2017-11-14 16:37:59 +05:30
c154921e10
Texture code reorganized
...
[ROCm/clr commit: e536eb038d
2017-11-14 11:09:35 +05:30
75a1e24519
Fix missing MARKER_END
...
Logging status of hipCtxSynchronize was missing
Test if hip profiling is active for MARKER_END in ihipPostLaunchKernel
Add MARKER_END after the completion of a kernel launched through
the "grid launch"
[ROCm/clr commit: e60a95d7dd
2017-11-13 16:13:19 -05:00
c3dc74be1b
Add an explicit check for proper rewriting of CUDA includes
...
[ROCm/clr commit: 642af94492
2017-11-13 21:02:42 +00:00
98034c4e3f
Add a preprocessor conditional to one of the tests
...
Hurrah, we can cope with ifdefs now (except for kernel launches)
[ROCm/clr commit: e8c1e7bb11
2017-11-13 20:58:55 +00:00
83e450e2c7
Use proper clang diagnostics for printing warnings
...
Much pretty. Very wow
This gives users all the usual power when it comes to manipulating
clang diagnostics. People can pass -Werror can have hipify fail if
it doesn't completely translate a file, for example. Much nicer
than reinventing the wheel.
[ROCm/clr commit: 506deca07f
2017-11-13 20:58:55 +00:00
842c8c28d7
Use a custom FrontendAction to simplify identifier translation
...
Most of what hipify does is really just replacing CUDA idenitifers
with HIP ones. CUDA function calls, preprocessor macro calls,
enum references, types, etc.
This is problematic: calls/types/enum-refs require name resolution
for the AST matcher to work. This fails in the presence of code
deleted by the preprocessor, and in two-pass template compilation.
Instead, we can simply hook the lexer and have it rewrite the
identifiers for us.
This approach means identifier transformations will work correctly
regardless of where they appear (and we get to delete lots of code)
- Fixes #260
- Helps a bit with #207 - it will still fail to translate kernel
calls in preprocessor-ignored code, but everything except kerel
launches should translate correctly now, even in
preprocessor-deleted code.
[ROCm/clr commit: 7a441aef10
2017-11-13 20:58:54 +00:00
e00c92d828
Add hipify mappings for all CUDA headers that have HIP equivalents
...
I'm particularly running into issues with `device_types.h` in real
CUDA code...
[ROCm/clr commit: 9014c89f04
2017-11-13 17:20:07 +00:00
41d7a05974
Add a test that exposes #260
...
[ROCm/clr commit: 0afc6b2e95
2017-11-13 16:18:15 +00:00
5ced4544eb
Add a couple of missing CHECK directives to concurrentKernels.cu
...
[ROCm/clr commit: 4353fc06d7
2017-11-13 16:17:19 +00:00
a857b650ad
Merge pull request #261 from gargrahul/fix_module_api_sample
...
Fix module_api sample
[ROCm/clr commit: 7f3653af7a
2017-11-13 11:55:54 +05:30
4ff4a71867
Fix module_api sample
...
[ROCm/clr commit: e434a23be5
2017-11-13 08:56:39 +05:30
Alex Voicu
Maneesh Gupta
Chris Kitching
Ben Sander
Rahul Garg
Evgeny Mankov
Phaneendr-kumar Lanka
Pierre