d9db9dabeb
Check for null event in hipEventElapsedTime
...
[ROCm/hip commit: f278f67d2d
2017-11-06 23:49:31 +00:00
465b24123b
hipStreamWaitEvent returns success if event created but not recorded
...
[ROCm/hip commit: 16708dd2e0
2017-11-06 23:49:31 +00:00
f80896d58b
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.
[ROCm/hip commit: 4a2e6f8955
2017-11-06 23:49:25 +00:00
33836c9521
Merge pull request #251 from ROCm-Developer-Tools/fix_event_state
...
Set event state AFTER it is recorded.
[ROCm/hip commit: 1131c9e41a
2017-11-06 07:28:11 +05:30
8845e38efd
Merge pull request #249 from bensander/warn_event
...
Add HIP_DB=warn + message if sync on dangerous event.
[ROCm/hip commit: a62d5aa875
2017-11-06 07:25:40 +05:30
6a4c50cf6c
Set event state AFTER it is recorded.
...
[ROCm/hip commit: 4c3b65a5cd
2017-11-05 10:33:18 -06:00
1f911cd23a
Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe
...
# Conflicts:
# src/hip_module.cpp
[ROCm/hip commit: bb1176001f
2017-11-03 10:53:39 +00:00
28eb8e2c3e
This introduces correct support for agent global variables, and implements hipModuleGetGlobal as an actual equivalent for cuModuleGetGlobal.
...
[ROCm/hip commit: 328c18b886
2017-11-03 01:44:48 +00:00
dab971370e
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).
...
[ROCm/hip commit: 2cacda91bb
2017-11-01 22:33:13 +00:00
2979a99888
Merge pull request #237 from bensander/use_ctxptr_for_p2p
...
Use ctxptr for p2p
[ROCm/hip commit: 09d866a639
2017-11-01 18:55:25 +01:00
6c60773e82
Add HIP_DB=warn + message if sync on dangerous event.
...
[ROCm/hip commit: 70c25bdf8e
2017-11-01 10:44:34 -07:00
6f3b28d27c
Merge pull request #245 from scchan/centos_fixes
...
various fixes for centos/rhel
[ROCm/hip commit: 86f62accfd
2017-11-01 18:10:29 +01: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
d938703f9b
Centos/RHEL - remove usage of constexpr since libc++ doesn't enable ctor for constexpr pair in C++11
...
[ROCm/hip commit: 99d32a195f
2017-10-31 18:16:12 +00:00
e88ef63bc8
Add ns-level timer for HIP API routines
...
Refactor some miuses of ihipLogStatus, these should only be in top-level
HIP APIs and should be paired with HIP_API_INIT calls.
[ROCm/hip commit: 7e908bdec8
2017-10-30 20:20:51 +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
ccb8b441a7
Check for null copyEngine before looking at peers.
...
[ROCm/hip commit: d610f16c47
2017-10-30 16:58:03 +00:00
6cc7f10e84
Merge remote-tracking branch 'origin/master' into HEAD
...
[ROCm/hip commit: a9789ddcda
2017-10-27 01:18:28 -04:00
ad9a636b90
Merge pull request #198 from AlexVlx/feature_support_globals_for_module_api
...
Feature support globals for module api
[ROCm/hip commit: f288f24e95
2017-10-27 01:53:34 +02:00
d7153f792d
Fix bug with peer-to-peer combined with context API
...
- Store context inside the tracker rather than using int deviceID that
was always mapped to primary context
- IsPeerWatcher now based on device IDs rather than specific peers.
[ROCm/hip commit: 7d30f32332
2017-10-26 19:44:22 +00:00
3ddd48ab86
Enhance debug for copy pointers
...
- show more pointer tracking fields
- show pointer info before and after "tailoring'
[ROCm/hip commit: 5d646d0fe3
2017-10-26 19:44:22 +00:00
c595b8f660
add HC_FEATURE_PRINTF around the printf buffer definition
...
[ROCm/hip commit: d91a4f5bd6
2017-10-25 12:00:02 -04:00
638f62b0c4
printf support for module API
...
[ROCm/hip commit: 1ddee10c2f
2017-10-24 00:55:41 -04:00
2bf3a98d83
replace __hcc_workweek__ with HC_FEATURE_PRINTF flag
...
[ROCm/hip commit: 5b9ce032d6
2017-10-23 18:30:08 -04:00
f52f6e53bb
Make elfio headers private
...
Change-Id: I3ba174bb46e84a75380207d93a0da6fe3703689e
[ROCm/hip commit: b792f9f507
2017-10-23 10:24:36 +05:30
3a0a22746d
Remove printf
...
[ROCm/hip commit: dd24983571
2017-10-20 13:24:04 -07: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
30872c6020
Modify device properties to use pool API.
...
- Also better error code checking
[ROCm/hip commit: c9f906c2ce
2017-10-20 14:49:29 +00:00
4dc24498af
hipDeviceReset(): make sure to reinitialize the printf buffer in hcc RT
...
[ROCm/hip commit: ccef1cbd6e
2017-10-18 16:26:13 -04:00
322cc5f7aa
This fixes incorrect usage of the reader object, which created arcane
...
mismatches when one reader accessed another's section.
[ROCm/hip commit: 70786a5563
2017-10-09 15:46:38 +01:00
fbcd3b0aa6
This adds cursory support for globals to the HIP module loading API. The
...
style is purposefully alien so as to signal that HIP experts should turn
it into HIP worthy code as soon as possible.
[ROCm/hip commit: ffffe052b6
2017-10-09 13:27:11 +01:00
8c67dd3b2e
Fixed hipTexture fetch issue
...
[ROCm/hip commit: 9929e8e717
2017-09-25 06:54:29 +05:30
9c0dfd3538
Bump device major version from 2 to 3
...
This would significantly improve performance for certain apps in kernel
selection logic.
[ROCm/hip commit: c74d3fe2cb
2017-09-15 15:47:39 +00:00
ab4128a53f
Merge pull request #179 from gargrahul/fix_hipmallocarray
...
Fixed hipMallocArray for 1D cases
[ROCm/hip commit: 3ed5f42dad
2017-09-14 12:59:33 -05:00
b446d8ebdf
Merge branch 'master' into hip_init_alloc
...
[ROCm/hip commit: fbd22c3e49
2017-09-14 11:53:33 -05:00
5418e3de45
Add HIP_INIT_ALLOC to init allocated memory.
...
[ROCm/hip commit: cea80cd8b3
2017-09-13 23:31:48 +00:00
3178dc04cd
Fixed hipMallocArray for 1D cases
...
[ROCm/hip commit: 2a915b2790
2017-09-12 21:52:11 +05:30
cbbb0826f2
hipStreamQuery uses av::is_empty. Add HIP_FORCE_NULL_STREAM.
...
[ROCm/hip commit: 4ac6d643c1
2017-08-31 03:00:14 +00:00
2a23d6dbe7
Refactor hipStreamWaitEvent
...
- Null streams use same flow as non-null.
- Add HIP_SYNC_STREAM_WAIT
- Resolve null stream.
[ROCm/hip commit: 882dab4536
2017-08-31 03:00:14 +00:00
f66a03fa9c
Merge pull request #167 from bensander/event_safety
...
Event safety
[ROCm/hip commit: cc08b36647
2017-08-29 13:55:46 -05:00
bd434502db
Lock streams when waiting on event completion or querying event safety.
...
[ROCm/hip commit: 6ff74d0e97
2017-08-28 18:40:16 -05:00
870228006e
Null check on input pointer arguments
...
[ROCm/hip commit: 38c9febe21
2017-08-25 08:46:34 +05:30
e673304653
Device count check on input device id
...
[ROCm/hip commit: 38ecfb9384
2017-08-22 14:05:44 +05:30
542bd863a2
Add some new HIP_TRACE_API options.
...
[ROCm/hip commit: e3161bb40e
2017-08-16 04:00:33 +00:00
c26e1d8c3d
Merge pull request #148 from weixingzhang/vmaddr1
...
[HIP Texture] The GPU virtual address for texture memory needs to be
[ROCm/hip commit: 8b7dde6519
2017-08-15 14:19:45 -05:00
c036a1b25c
[texture] guard new HCC APIs under workweek
...
Change-Id: I4f60a64fb0b0496ca1eb01ffe6ddda121c25d976
[ROCm/hip commit: e40047f2bf
2017-08-15 15:51:38 +05:30
0a25b3e972
Merge pull request #138 from scchan/threadfence_system
...
implement __threadfence_system
[ROCm/hip commit: f14e31d9d7
2017-08-10 12:51:05 +05:30
d4d9836a96
Merge pull request #145 from gargrahul/context_mgmt_changes
...
Updated context management logic:
[ROCm/hip commit: 660061c3b4
2017-08-09 17:47:45 -05:00
5bf177340a
[HIP Texture] The GPU virtual address for texture memory needs to be
...
aligned.
In hcc_am, a bigger buffer will be allocated for alignment purpose
and _unalignedDevicePointer is added in struct AmPointerInfo for
original allocated address.
[ROCm/hip commit: 4264a4ce56
2017-08-08 11:18:00 -04:00
c050396510
Bump min hcc_workweek required for named kernel dispatch to 17312
...
Change-Id: I8c7b58306b279ed113d03260e4bc6086bb8b4e68
[ROCm/hip commit: 8d6acec135
2017-08-08 11:08:55 +05:30
Ben Sander
Maneesh Gupta
Alex Voicu
Siu Chi Chan
Aditya Atluri
Rahul Garg
Wen-Heng (Jack) Chung
Weixing Zhang