AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
2,161 Commits 3 Branches 0 Tags
e186bd9533bce6c11dc7d4a447dbde683d4bd1fc
Commit Graph

33 Commits

Author SHA1 Message Date
Alex Voicu e186bd9533 This is primarily intended as an additional cleanup of the module functionality, in the aftermath of adopting module based dispatch. The main effort was associated with refactoring the questionable ihipModuleGetSymbol. It was quaintly written and misleading, in that it had little to do with getting symbols, and was exactly retrieving a kernel object. Error handling is modified so as to reduce branching depth. Functions which serve as interfaces to the HSA RT are moved in a separate helper header. Code object readers are properly deleted. Some leftover dead functionality pertaining to associating namespace scope variables with their allocated memory is removed. Executable loading is changed to use a string which holds the ELF image of the code object being loaded, thus avoiding some corner cases where using a istream would fail. 2017-12-03 23:09:06 +00:00
Alex Voicu d37a5a6008 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
Rahul Garg eff09fcd86 -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
Alex Voicu f7919a8321 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
Alex Voicu 8854dec897 Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 2017-11-07 00:01:22 +00:00
Ben Sander 955cfbfdc7 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
Maneesh Gupta 958ca62c63 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
Alex Voicu 17753cbd92 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
Ben Sander fe32685fbc Merge pull request #237 from bensander/use_ctxptr_for_p2p 
Use ctxptr for p2p
 2017-11-01 18:55:25 +01:00
Ben Sander 172b9f96e6 Add HIP_DB=warn + message if sync on dangerous event. 2017-11-01 10:44:34 -07:00
Alex Voicu 28f87f7d2e 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. 2017-11-01 15:09:59 +00:00
Ben Sander dc7d993a02 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.
 2017-10-30 20:20:51 +00:00
Ben Sander a417241507 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.
 2017-10-26 19:44:22 +00:00
Ben Sander 1ea468e279 Merge branch 'master' into hip_init_alloc 2017-09-14 11:53:33 -05:00
Ben Sander fff42fd591 Add HIP_INIT_ALLOC to init allocated memory. 2017-09-13 23:31:48 +00:00
Ben Sander fff74eee21 hipStreamQuery uses av::is_empty. Add HIP_FORCE_NULL_STREAM. 2017-08-31 03:00:14 +00:00
Ben Sander ed8c3ba7e7 Refactor hipStreamWaitEvent 
- Null streams use same flow as non-null.
- Add HIP_SYNC_STREAM_WAIT
- Resolve null stream.
 2017-08-31 03:00:14 +00:00
Ben Sander bc9ba7cd81 Lock streams when waiting on event completion or querying event safety. 2017-08-28 18:40:16 -05:00
Ben Sander d6aee9ad8f Add some new HIP_TRACE_API options. 2017-08-16 04:00:33 +00:00
Rahul Garg c4e9323877 Updated context management logic: 
1) hipSetDevice sets a flag so that next call to hipCtxGetCurrent returns primary context on current device
2) hipCtxGetCurrent returns primary context on current device if TLS context stack is empty
3) hipCtxPopCurrent falls back to primary context on current device as default
4) hipCtxPushCurrent, hipCtxSetCurrent and hipCtxCreate reset the flag set in hipSetDevice
 2017-08-08 07:02:22 +05:30
Ben Sander 6576201ec2 Make host memory allocations coherent by default. 
Associated change is to optimize event recording so it uses
agent-scope releaes (since it was only using system-scope release
to support non-coherent host mem).

Flags and environment variables exist to obtain previous behavior
if desired.  Options are documented in new performance guide.
 2017-07-26 19:20:34 -05:00
Sun, Peng e5ce585307 Add support of HIP_HIDDEN_FREE_MEM, to deduct the returned available 
memory from hipMemGetInfo API, measured in MB.

Change-Id: I7a8260c12e032e04e26611db4c38c893a29f2653
 2017-06-26 15:29:38 -05:00
Ben Sander 823281dcba 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
 2017-06-05 08:50:22 -05:00
Ben Sander 6cc5dc0326 Set event->_stream on hipHccModuleLaunchKernel path if start/stop used 
Ensure _stream is always non-null in recorded events.
Fixes isDefaultStream fault.
 2017-05-30 21:55:46 -05:00
Ben Sander b2b620c12b Add isDefaultStream() accessor. 
Fix code that checked for stream==nullptr after stream had been
resolved to a "true stream".
 2017-05-26 13:46:48 -05:00
Ben Sander d43d57d39c Remove HIP_MAX_QUEUES (replaced with HCC_MAX_QUEUES) 2017-05-23 23:48:01 -05:00
Ben Sander 8bc6ee5932 Add initial HIP_SYNC_NULL_STREAM=0 mode. 
This eliminates host-synchronization for null stream.  Instead, the
null-stream uses GPU-side events to wait for other streams.
Default is OFF pending additional testing.

Add enhanced null-stream test.

Also refine HIP_TRACE_API.
 2017-05-16 19:04:25 -05:00
Ben Sander 7e7ba5027f Add HIP_TRACE_API=4. Only display memory allocation/free apis. 2017-05-16 19:04:25 -05:00
Ben Sander ff9bed6535 hipHostMalloc allocation are mapped to all devices by default. 
Support hipHostMallocPortable flag.
Default flags are hipHostMallocPortable | hipHostMallocMapped.

Also:
-refactor tests to move addCount and addCountReverse into HipTest
namespace.
-test multi-GPU host memory.
 2017-05-10 17:34:36 -05:00
Ben Sander e9eaadd135 Finish adding start/stop event support to hipHccModuleLaunchKernel. 
Change interface to use hipEvent_t rather than hipEvent_t*

Change-Id: I259062dc087a13d51dc27f84e1e8861f332a104d
 2017-04-06 21:02:50 -05:00
sunway513 f442e975c6 Refactor events and add initial event option for hipHccModuleLaunchKernel 
- Change hipEvent_t to a class.
- Move event logic inside the class.
- Add _type to support Independent, StartCommand, StopCommand events.
  StartCommand returns start timestamp from events.

Change-Id: I4ddd694f2645a3ff7170c9111dc1d3e39931ca21
 2017-04-06 21:02:50 -05:00
Aditya Atluri 3859f59189 Fixed copyright and header names 
Change-Id: Id595c65ea3b7289e87be4c42db5d8a31905a4fdd
 2017-03-31 12:40:29 -05:00
Aditya Atluri 1cead6a4cd added new api hipHccModuleLaunchKernel 
1. hipHccModuleLaunchKernel is same as hipModuleLaunchKernel with OpenCL workitem model
2. Added copy right
3. Fixed header naming

Change-Id: I6a7c35a3566e2f8d3f5056613e34193775d4b236
 2017-03-31 12:11:34 -05:00