Refactor hipStreamWaitEvent

- Null streams use same flow as non-null.
- Add HIP_SYNC_STREAM_WAIT
- Resolve null stream.


[ROCm/hip commit: 882dab4536]

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.");

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;
 

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);

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);
     }