From 4ac6d643c12cf0da8a07726c52d2d0eb72ecd781 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 31 Aug 2017 02:50:30 +0000
Subject: [PATCH] hipStreamQuery uses av::is_empty.  Add HIP_FORCE_NULL_STREAM.

---
 src/hip_hcc.cpp        |  4 ++++
 src/hip_hcc_internal.h |  1 +
 src/hip_stream.cpp     | 38 ++++++++++++++++++++++----------------
 3 files changed, 27 insertions(+), 16 deletions(-)

diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 7c9f0f966f..860c6f87a8 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -90,6 +90,8 @@ int HIP_SYNC_HOST_ALLOC = 1;
 
 int HIP_SYNC_STREAM_WAIT = 0;
 
+int HIP_FORCE_NULL_STREAM=0;
+
 
 #if (__hcc_workweek__ >= 17300)
 // Make sure we have required bug fix in HCC
@@ -1251,6 +1253,8 @@ 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_FORCE_NULL_STREAM,  0, "Force all stream allocations to secretly return the null stream");
+
     READ_ENV_I(release, HIP_SYNC_STREAM_WAIT, 0, "hipStreamWaitEvent will synchronize to host");
 
 
diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h
index af855861c6..88717e48bf 100644
--- a/src/hip_hcc_internal.h
+++ b/src/hip_hcc_internal.h
@@ -68,6 +68,7 @@ extern int HIP_SYNC_HOST_ALLOC;
 extern int HIP_SYNC_STREAM_WAIT;
 
 extern int HIP_SYNC_NULL_STREAM;
+extern int HIP_FORCE_NULL_STREAM;
 
 // TODO - remove when this is standard behavior.
 extern int HCC_OPT_FLUSH;
diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp
index 0b645ec658..7dd6efd39c 100644
--- a/src/hip_stream.cpp
+++ b/src/hip_stream.cpp
@@ -38,21 +38,26 @@ hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags)
     hipError_t e = hipSuccess;
 
     if (ctx) {
-        hc::accelerator acc = ctx->getWriteableDevice()->_acc;
 
-        // TODO - se try-catch loop to detect memory exception?
-        //
-        //Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete:
-        //This matches CUDA stream behavior:
+        if (HIP_FORCE_NULL_STREAM) {
+            *stream = 0; 
+        } else {
+            hc::accelerator acc = ctx->getWriteableDevice()->_acc;
 
-        {
-            // Obtain mutex access to the device critical data, release by destructor
-            LockedAccessor_CtxCrit_t  ctxCrit(ctx->criticalData());
+            // TODO - se try-catch loop to detect memory exception?
+            //
+            //Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete:
+            //This matches CUDA stream behavior:
 
-            auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
+            {
+                // Obtain mutex access to the device critical data, release by destructor
+                LockedAccessor_CtxCrit_t  ctxCrit(ctx->criticalData());
 
-            ctxCrit->addStream(istream);
-            *stream = istream;
+                auto istream = new ihipStream_t(ctx, acc.create_view(), flags);
+
+                ctxCrit->addStream(istream);
+                *stream = istream;
+            }
         }
 
         tprintf(DB_SYNC, "hipStreamCreate, %s\n", ToString(*stream).c_str());
@@ -119,15 +124,14 @@ hipError_t hipStreamQuery(hipStream_t stream)
         stream =  device->_defaultStream;
     }
 
-    int pendingOps = 0;
+    bool isEmpty = 0;
 
     {
         LockedAccessor_StreamCrit_t crit(stream->_criticalData);
-        pendingOps = crit->_av.get_pending_async_ops();
+        isEmpty = crit->_av.get_is_empty();
     }
 
-
-    hipError_t e = (pendingOps > 0) ? hipErrorNotReady : hipSuccess;
+    hipError_t e = isEmpty ? hipSuccess : hipErrorNotReady ;
 
     return ihipLogStatus(e);
 }
@@ -167,7 +171,9 @@ hipError_t hipStreamDestroy(hipStream_t stream)
 
     //--- Drain the stream:
     if (stream == NULL) {
-        e = hipErrorInvalidResourceHandle; // TODO - review - what happens if try to destroy null stream
+        if (!HIP_FORCE_NULL_STREAM) {
+            e = hipErrorInvalidResourceHandle; 
+        } 
     } else {
         stream->locked_wait();