From 2a23d6dbe710bf2e1901a16677be47e76fc80c12 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 30 Aug 2017 19:56:47 +0000 Subject: [PATCH] Refactor hipStreamWaitEvent - Null streams use same flow as non-null. - Add HIP_SYNC_STREAM_WAIT - Resolve null stream. [ROCm/hip commit: 882dab4536736945e574344674bc83f890116ae4] --- projects/hip/src/hip_hcc.cpp | 3 ++ projects/hip/src/hip_hcc_internal.h | 1 + projects/hip/src/hip_stream.cpp | 15 +++---- .../runtimeApi/stream/hipStreamWaitEvent.cpp | 40 +++++++++++++++++-- 4 files changed, 46 insertions(+), 13 deletions(-) diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index dc72714e3e..7c9f0f966f 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -88,6 +88,8 @@ int HIP_HOST_COHERENT = 1; int HIP_SYNC_HOST_ALLOC = 1; +int HIP_SYNC_STREAM_WAIT = 0; + #if (__hcc_workweek__ >= 17300) // Make sure we have required bug fix in HCC @@ -1249,6 +1251,7 @@ void HipReadEnv() READ_ENV_I(release, HIP_SYNC_HOST_ALLOC, 0, "Sync before and after all host memory allocations. May help stability"); READ_ENV_I(release, HIP_SYNC_NULL_STREAM, 0, "Synchronize on host for null stream submissions"); + READ_ENV_I(release, HIP_SYNC_STREAM_WAIT, 0, "hipStreamWaitEvent will synchronize to host"); READ_ENV_I(release, HIP_HOST_COHERENT, 0, "If set, all host memory will be allocated as fine-grained system memory. This allows threadfence_system to work but prevents host memory from being cached on GPU which may have performance impact."); diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 56ca37b3e2..af855861c6 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -65,6 +65,7 @@ extern int HIP_HIDDEN_FREE_MEM; //--- // Chicken bits for disabling functionality to work around potential issues: extern int HIP_SYNC_HOST_ALLOC; +extern int HIP_SYNC_STREAM_WAIT; extern int HIP_SYNC_NULL_STREAM; diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index 51aeb01412..0b645ec658 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -93,18 +93,15 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int } else if (event->_state != hipEventStatusUnitialized) { - if (stream != hipStreamNull) { - + if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) { + // conservative wait on host for the specified event to complete: + event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); + } else { + stream = ihipSyncAndResolveStream(stream); // This will user create_blocking_marker to wait on the specified queue. stream->locked_streamWaitEvent(event); - - } else { - // TODO-hcc Convert to use create_blocking_marker(...) functionality. - // Currently we have a super-conservative version of this - block on host, and drain the queue. - // This should create a barrier packet in the target queue. - // TODO-HIP_SYNC_NULL_STREAM - stream->locked_wait(); } + } // else event not recorded, return immediately and don't create marker. return ihipLogStatus(e); diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp index 9bbd43828c..f5b1b79550 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp @@ -446,9 +446,41 @@ int main(int argc, char *argv[]) if (p_tests & 0x1000) { - printf ("==> Test 0x1000 try null stream\n"); - hipStreamQuery(0/* try null stream*/); + printf ("==> Test 0x1000 simple null stream tests\n"); + // try some null stream: + hipStreamQuery(0); + + + hipStream_t s1; + hipEvent_t e1; + + { + // stream null waits on event in s1 stream: + HIPCHECK(hipStreamCreate(&s1)); + HIPCHECK(hipEventCreate(&e1)); + + HIPCHECK(hipEventRecord(e1, s1)) + + HIPCHECK(hipStreamWaitEvent(hipStream_t(0), e1, 0/*flags*/)); + + HIPCHECK(hipStreamDestroy(s1)); + HIPCHECK(hipEventDestroy(e1)); + } + + { + // stream s1 waits on event in null stream: + HIPCHECK(hipStreamCreate(&s1)); + HIPCHECK(hipEventCreate(&e1)); + + HIPCHECK(hipEventRecord(e1, hipStream_t(0))) + + HIPCHECK(hipStreamWaitEvent(s1, e1, 0/*flags*/)); + + HIPCHECK(hipStreamDestroy(s1)); + HIPCHECK(hipEventDestroy(e1)); + } + } @@ -471,8 +503,8 @@ int main(int argc, char *argv[]) } - { - printf ("test: alternating memcpy/count-reverse followed by event\n"); + if (p_tests & 0x4000 ) { + printf ("test: %x alternating memcpy/count-reverse followed by event\n", p_tests); RUN_SYNC_TEST(0x4000, streamersDev0, sync_queryAllUntilComplete(streamersDev0), true); RUN_SYNC_TEST(0x8000, streamersDev0, sync_streamWaitEvent(streamersDev0.back()->event(), 0, sideStreams[0], false), true); }