AI/rocm-systems
2
0 Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add initial HIP_SYNC_NULL_STREAM=0 mode.

Browse Source 
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.


[ROCm/clr commit: 8bc6ee5932]
This commit is contained in:
Ben Sander committed 2017-05-16 19:04:25 -05:00
1 parent db102ab82f
commit 4ac6ac9d1d
11 files changed
+320 -96
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

projects/clr/hipamd/src/hip_stream.cpp
+2 -2
View File
@@ -150,7 +150,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream)
if (stream == NULL) {
ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
ctx->locked_syncDefaultStream(true/*waitOnSelf*/);
ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/);
} else {
stream->locked_wait();
e = hipSuccess;
@@ -174,7 +174,7 @@ hipError_t hipStreamDestroy(hipStream_t stream)
//--- Drain the stream:
if (stream == NULL) {
ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
ctx->locked_syncDefaultStream(true/*waitOnSelf*/);
ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true /*syncToHost*/);
} else {
stream->locked_wait();
e = hipSuccess;