Handle GraphExec instance is destroyed before async launch completes
GraphExec instance is destroyed after async launch completes
GraphExec instance is destroyed without a launch
Change-Id: I45a7c82295fea916c7559bd8f796df710513aea1
- Application is passing null for parameter stream in API hipStreamWaitEvent
- When event stream isn't capturing and event is not recorded, causes segfault because we are accessing deviceId() from waitStream
Change-Id: I8b87ffd6f234677f68b66dcb7ef44b2ff04a7c91
=> hipDeviceSynchronize is not allowed during capture.
=> hipEventSynchronize during capture should return hipErrorCapturedEvent error
=> hipEventQuery during capture should return hipErrorCapturedEvent error
hipStreamSynchronize, hipEventSynchronize, hipStreamWaitEvent, hipStreamQuery
For Side Stream(Stream that is not currently under capture):
=> If current thread is capturing in relaxed mode, calls are allowed
=> If any stream in current/concurrent thread is capturing in global mode, calls are not allowed
=> If any stream in current thread is capturing in ThreadLocal mode, calls are not allowed
For Stream that is currently under capture
=> calls are not allowed
=> Any call that is not allowed during capture invalidates the capture sequence
=> It is invalid to call synchronous APIs during capture. Synchronous APIs,
such as hipMemcpy(), enqueue work to the legacy stream and synchronize it before returning.
Change-Id: I201c6e63e1a5d93fd416a3b520264c0fdbe31237
Mempool has capability to track dependency between streams for
faster memory reuse. Enable that capability.
Change-Id: I28266a7e38d0fc4c5d027b9542d3719653840821
Dont track the status of fence_dirty_ flag on the host, instead clear it
when we submit a barrier on the respective stream.
Change-Id: I4d98dbf20c81379c9c5da9f5b67629a8f9f6dfcd
- Use isFenceDirty() check in hipStreamQuery.
In case if there is a stopEvent for hipExtLaunch* API we avoid system
scope if there are consecutive hipExtLaunch* and replace them with agent
scope, but then if one uses hipStreamQuery it should detect we need a
barrier with system scope.
Change-Id: I539810f659d4e1c49104e3a0da6ff2803c6b3dee
Runtime has to release extra memory, held by the pools,
in synchronization points for event, stream or device.
Change-Id: Id533a5e1d137812aa72bdfe101b4b333c6a43d66
- Allow capture to be less rectrictive if we set the global thread
interaction mode for the thread.
Change-Id: I84f65d9418ac26ada0477c85a45a3831c2351ce4
- correct error for hipStreamWaitEvent when event recorded before
capture
- correct hipEventSync when event is synced during capture
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I7ecbed5621eaf323846d4ccb20ec112aaa8a5757
With global flag, In addition to global vector, we need to check
thread local vector too. If both of them are empty then only
return false.
Change-Id: I3a24d16a35c702b34b9a8867dc410ebd80108f25
Avoid syncing blocking streams with the default stream,
since that introduces extra command dependencies and
doesn't allow to destroy memory after last submission
Change-Id: I618e9bd2091c4cf9157125612d8c4759030c5a80
- invalid to synchronize or query the execution status of a capturing stream but non-capturing streams can.
- If any local thread has an ongoing or concurrent capture sequence initiated with hipStreamCaptureModeGlobal, it is prohibited from unsafe calls
Change-Id: Ifa641e807216d3b7c3e8c2fb1be2f7a50bd641df
- cuda shows a different behaviour and different error need more
investigation
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I68771102ba4dff6157bca34a4135cb245f023d08
This reverts commit f39655c0c8.
Reason for revert: patch does not fix all stream hangs. So another patch was merged which fixes all issues.
Change-Id: I332d1ea29c23747b46b7667fe3e34e0ceefd2b23
- if the stream is null ptr and there is capture ongoing return an error
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: Iff33e4aebc253fb7bb56daf5c455722d726c6705
- stream capture should be done before any sync APIs.
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I3d65f67ee68777be71f97f48d460ccaefdd4e1af
- null stream should return error on host function
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I17659ad7af0814091dbf6cb4a2cb2644514551f0
Skip destruction of the null queue for MT. Windows can destroy worker
thread on exit during the stream destruction, causing a race condition.
Change-Id: I3e4bc3ab2b477d38c0aec1fc59e6f1af0f8f6b0f
Parallel streams can be destroyed during capture. Make sure runtime
removes the destoryed streams from the list of parallel captured streams
Change-Id: I46b6cbb951e8711f0bf02f2826c3b890b2541ffe
Remove extra barrier, since ROCR backend in DD mode blocks HW queue now when a callback is injected
Add a notificaiton for MT mode about possible waiter for a callback
Change-Id: Ifd70ce5597e1ba868e4197ad1850ace11a4f90ae
- Aggregate all TLS(Thread Local Storage) variables into a single class
- This is to improve cache accesses per thread
Change-Id: Ic8361eaeae290fff00254684e309471958365eb9
Remove the api_callbacks_table_t that was holding the API activities and
user callbacks. Instead use a single roctracer callback (TracerCallback)
used to report both API activities and callbacks.
Remove the hipInitActivityCallback that was setting the ROCtracer
callback and memory pool for asynchronous activities as it did not
allow disctinct pools to be used for each activity. Instead, use
hipRegisterTracerCallback to set the single roctracer callback.
Change-Id: I4c10f04f29a6e4cce8caf15db3016c3f72c86b04
sdashmiz
Anusha GodavarthySurya
Rakesh Roy
Ioannis Assiouras
German Andryeyev
Saleel Kudchadker
Jatin Chaudhary
Sourabh Betigeri
Jaydeep Patel
Christophe Paquot
Ajay GunaShekar
Laurent Morichetti
Jason Tang
Sarbojit Sarkar