02fc25de4c
Remove stray leftover comment.
2017-12-02 00:03:10 +00:00
6027d3f332
Replace archaic use of homebrew functionality with calls to the HC maths library. This fixes a hang observed when building hipTestDeviceDouble.
2017-12-02 00:01:47 +00:00
7c0b9a005b
Fix legacy mode detection of the address of an agent allocated variable. In this mode, there exist two executables per each code object, one created by HCC and one created by HIP. Since we dispatch through HCC in legacy mode, we should obtain the address for an agent allocated variable from the latter's executable. Also add two omitted validity checks, whose absence could lead to segfaults when the current process had no .kernel section and / or when an invalid or empty blob was extracted from the latter.
2017-11-30 03:29:04 +00:00
32e11e7dc6
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 d2fd1f5
2017-11-29 21:49:10 +00:00
fbaf729f88
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 d2fd1f5
2017-11-29 21:36:29 +00:00
b881cf713c
Fix compiler version check.
2017-11-29 03:05:53 +00:00
3ed8897a5a
Add missing file.
2017-11-29 02:16:44 +00:00
faa546d194
Fix oversight in selection mechanism which led to erroneous code to be compiled for the grid_launch_GGL component.
2017-11-29 01:37:52 +00:00
89e9399427
Choose whether or not to use functional grid_launch based on the version of HCC used to compile.
2017-11-29 00:17:44 +00:00
5aeb5dcd6f
Remove leftover agent allocated globals.
2017-11-28 19:56:04 +00:00
6e4ca3fbb4
Change memset kernel to use memcpy instead of placement new. Simplify indexers.
2017-11-28 19:45:47 +00:00
02c2bfc7ef
Re-sync with upstream and re-factor platform global management for texture references.
2017-11-28 19:15:29 +00:00
dc67ca3feb
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
# Conflicts:
# src/hip_module.cpp
2017-11-28 17:29:11 +00:00
265c3b224e
Fix float2int rounding functions
...
Change-Id: I67943859a6344c5eec0eaa23418c9b802ef72468
2017-11-28 17:23:43 +00:00
a6ccaf3d57
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.
2017-11-28 17:23:06 +00:00
a401ce6e5d
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.
2017-11-28 17:23:06 +00:00
0da0426f94
Merge pull request #256 from gargrahul/texture_driver_api_support
...
Texture driver APIs support
2017-11-27 13:52:39 -06:00
4c96882366
Fix float2int rounding functions
...
Change-Id: I67943859a6344c5eec0eaa23418c9b802ef72468
2017-11-23 09:57:24 +05:30
56862b1c35
Fixed review comments
2017-11-21 21:19:06 +05:30
5e16ee0d1f
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.
2017-11-21 13:15:13 +00:00
9866fa250d
Changed function hipMemcpy_2D to hipMemcpyParam2D
2017-11-21 12:36:24 +05:30
9d088d2283
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.
2017-11-21 02:40:34 +00:00
1824fb7698
Clean-up some remaining noise in program_state.cpp.
2017-11-20 22:41:46 +00:00
7d5a45ac1a
Correct ill-formed merge in earlier commit and adjust for differences with the new CUDA natural indexing mechanism.
2017-11-20 16:33:52 +00:00
c5f2b22d0d
Re-sync with upstream.
2017-11-20 15:34:50 +00:00
e8ede28ec4
Merge pull request #264 from pzins/missing_end_marker
...
Fix missing MARKER_END
2017-11-20 06:08:01 -06:00
f97c5f9a64
-Moved coGlobals in hipModule class (takes care of multi module case)
...
-Used mutex scope for updating coGlobals
2017-11-20 16:23:18 +05:30
c7d60a7a75
Update hipModuleGetTexRef API
2017-11-19 22:10:46 +05:30
cffd0e14eb
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.
2017-11-19 01:54:12 +00:00
ae1eb7a03a
Removed redundant desc variable
2017-11-15 18:28:27 +05:30
4b19c2aa0c
-Fixed texture driver API sample
...
-Added hipTexRefSetAddress and hipTexRefSetAddress2D APIs
2017-11-15 18:23:28 +05:30
63680edd30
Texture code reorganized
2017-11-14 11:09:35 +05:30
6baaed8e48
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"
2017-11-13 16:13:19 -05:00
f7726cd416
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
2017-11-09 23:43:07 +00:00
ef09c4918d
Texture driver APIs support
2017-11-09 22:10:55 +05:30
31bcb59f62
Merge pull request #250 from AlexVlx/feature_add_agent_global_support
...
Support for agent globals
2017-11-09 07:52:09 +05:30
3d248927e4
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
# Conflicts:
# tests/src/runtimeApi/stream/hipStreamSync2.cpp
2017-11-08 10:26:30 +00:00
d8e323d4b5
Clean up trailing whitespace so as to reduce noise in #246 .
2017-11-08 00:08:55 +00:00
adaf6b8dff
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
2017-11-07 00:01:22 +00:00
f278f67d2d
Check for null event in hipEventElapsedTime
2017-11-06 23:49:31 +00:00
16708dd2e0
hipStreamWaitEvent returns success if event created but not recorded
2017-11-06 23:49:31 +00:00
4a2e6f8955
Make hipEvent_t thread safe.
...
Support re-recording of same event by different threads.
- Add criticalData structure to hipEvent_t, similar to mechanism used
for streams, contexts, device. Events are always locked
after streams to avoid deadlock.
- ihipEvent_t::locked_copyCrit can be used to copy critical state
including marker. The critical state in the event can then
be re-recorded.
- refactor hipEventElapsedTime. Remmove stale debug code, native signal
refs.
2017-11-06 23:49:25 +00:00
1131c9e41a
Merge pull request #251 from ROCm-Developer-Tools/fix_event_state
...
Set event state AFTER it is recorded.
2017-11-06 07:28:11 +05:30
a62d5aa875
Merge pull request #249 from bensander/warn_event
...
Add HIP_DB=warn + message if sync on dangerous event.
2017-11-06 07:25:40 +05:30
4c3b65a5cd
Set event state AFTER it is recorded.
2017-11-05 10:33:18 -06:00
bb1176001f
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
# Conflicts:
# src/hip_module.cpp
2017-11-03 10:53:39 +00:00
328c18b886
This introduces correct support for agent global variables, and implements hipModuleGetGlobal as an actual equivalent for cuModuleGetGlobal.
2017-11-03 01:44:48 +00:00
2cacda91bb
Correctly deal with functions from shared objects, wherein the program visible VA == so_base_va + st_value(function_symbol). Remove quaint usage of pfe for hipMemset (which is actually fill_n).
2017-11-01 22:33:13 +00:00
09d866a639
Merge pull request #237 from bensander/use_ctxptr_for_p2p
...
Use ctxptr for p2p
2017-11-01 18:55:25 +01:00
70c25bdf8e
Add HIP_DB=warn + message if sync on dangerous event.
2017-11-01 10:44:34 -07:00
Alex Voicu
Maneesh Gupta
Ben Sander
Rahul Garg
Pierre