introduce LockedAccessor option so destructor does not unlock.
Allows locks to exist across function boundaries, required
for hipLaunchKernel macro which has several unusual requirements.
(including C comppatibility, must use variadic macro, more).
- Remove "call-to-call" for hipStreamCreate and hipEventCreate.
These now call an internal functions rather than calling through
hipStreamCreateWithFalgs and hipEventCreateWithFlags.
- Add HIP_INIT_API for more functions so they trace correctly.
- Use stream#DEVICE.STREAMID in debug messages via new specialization in
tace_helper.
Move critical data into separate class and protect with LockAccessor
wrapper class.
For device, the streams list is the critical data since it is modified when
streams are created or destroyed. The streams list is accessed in
several places including when synchronizing across all streams on the
device (ie from the default stream).
Other device data is set once by the device cosntructor and is not critical
so
All functions which acquire the LockAccessor now named with "locked_" prefix.
- Some comments in hip_hcc.cpp
- document hipHostRegister* flags.
- expand docs on hipHostRegister/Unregister. Use "register" rather than
"pin" to describe action these take.
- change required CUDA version to 6 (require unified memory)
- remake doxygen.
- Validate compile-time disables.
- Add README.md section explain how to install/use CodeXL tracing
- Add code docs on trace_helper.h
- fix color on hipLaunchKernel to green.
Note hipHostMalloc (not hipHostAlloc or hipMallocHost).
- the hipHost* is used for all HIP APIs dealing with Host memory.
(including hipHostMalloc, hipHostFree, hipHostRegister,
hipHostUnregister, hipHostGetFlags, hipHostGetDevicePointer).
- hipMallocHost is consistent with "hipMalloc" for allocating device
memory. Enumerations hipHostMalloc* also used as optional
flags parm to hipHostMalloc.
These will print compiler warnings if used, so we can weed them out
before removing.
Also add a default flags args for hipHostAlloc, in the C++ functioin
headers. So you can replace hipMallocHost(&ptr, size( with hipHostAlloc(&ptr, size)
-Move staging buffer locks inside the staging buffer code.
-Remove dedicated per-device completion_signal + per-device lock -
instead allocated signal from the per-stream pool. This elimintes
the lock and allows more concurrency.
-remove switch HIP_DISABLE_BIDIR_MEMCPY
- refactor staging buffer to operate on hsa* data structures not
hc::accelerator.
- use hsa_memory_allocate to allocate staging buffers rather than
am_alloc.
- Refactor device reset with single member function. Don't reallocate
staging buffers on reset.
- Properly track dependencies based on command type. Add new deps for
H2D and D2D rather than overloading H2D.
Ben Sander
Aditya Atluri
scchan
pensun
Maneesh Gupta
Aditya Atluri