Add isDefaultStream() accessor.

Fix code that checked for stream==nullptr after stream had been
resolved to a "true stream".

hipamd/src/hip_event.cpp

@@ -53,14 +53,19 @@ void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf, ihip
 
 void ihipEvent_t::setTimestamp()
 {
+    bool isReady0 = _marker.is_ready();
+    bool isReady1;
+    int val = 0;
     if (_state == hipEventStatusRecorded) {
         // already recorded, done:
         return;
     } else {
         // TODO - use completion-future functions to obtain ticks and timestamps:
         hsa_signal_t *sig  = static_cast<hsa_signal_t*> (_marker.get_native_handle());
+        isReady1 = _marker.is_ready();
         if (sig) {
-            if (hsa_signal_load_acquire(*sig) == 0) {
+            val = hsa_signal_load_acquire(*sig);
+            if (val == 0) {
 
                 if ((_type == hipEventTypeIndependent) || (_type == hipEventTypeStopCommand)) {
                     _timestamp =  _marker.get_end_tick();
@@ -75,6 +80,10 @@ void ihipEvent_t::setTimestamp()
             }
         }
     }
+
+    if (_state != hipEventStatusRecorded) {
+        printf (" not ready isReady0=%d val=%d isReady1=%d\n", isReady0, val, isReady1);
+    }
 }
 
 
@@ -118,11 +127,11 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
 
         event->_stream = stream;
 
-        if (HIP_SYNC_NULL_STREAM && stream == NULL) {
+        if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
 
             // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
             
-            // If stream == NULL, wait on all queues.
+            // If default stream , then wait on all queues.
             ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
             ctx->locked_syncDefaultStream(true, true);
 
@@ -167,7 +176,7 @@ hipError_t hipEventSynchronize(hipEvent_t event)
         } else if (event->_state == hipEventStatusCreated ) {
             // Created but not actually recorded on any device:
             return ihipLogStatus(hipSuccess);
-        } else if (HIP_SYNC_NULL_STREAM && (event->_stream == NULL)) {
+        } else if (HIP_SYNC_NULL_STREAM && (event->_stream->isDefaultStream() )) {
             auto *ctx = ihipGetTlsDefaultCtx();
             // TODO-HIP_SYNC_NULL_STREAM - can remove this code
             ctx->locked_syncDefaultStream(true, true);
@@ -175,6 +184,8 @@ hipError_t hipEventSynchronize(hipEvent_t event)
         } else {
             event->_marker.wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
 
+            assert (event->_marker.is_ready());
+
             return ihipLogStatus(hipSuccess);
         }
     } else {

hipamd/src/hip_hcc_internal.h

@@ -538,10 +538,12 @@ public:
     const ihipDevice_t *     getDevice() const;
     ihipCtx_t *              getCtx() const;
 
+    bool isDefaultStream() const { return _id == 0; };
+
 public:
     //---
     //Public member vars - these are set at initialization and never change:
-    SeqNum_t                    _id;   // monotonic sequence ID
+    SeqNum_t                    _id;   // monotonic sequence ID.  0 is the default stream.
     unsigned                    _flags;
 
 
@@ -560,6 +562,7 @@ private:
 
     void addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t sizeBytes);
 
+
 public: // TODO - move private
     // Critical Data - MUST be accessed through LockedAccessor_StreamCrit_t
     ihipStreamCritical_t        _criticalData;

hipamd/src/hip_stream.cpp

@@ -146,7 +146,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream)
 
     hipError_t e = hipSuccess;
 
-    if (stream == NULL) {
+    if (stream == hipStreamNull) {
         ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
         ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/);
     } else {
@@ -198,7 +198,7 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags)
 
     if (flags == NULL) {
         return ihipLogStatus(hipErrorInvalidValue);
-    } else if (stream == NULL) {
+    } else if (stream == hipStreamNull) {
         return ihipLogStatus(hipErrorInvalidResourceHandle);
     } else {
         *flags = stream->_flags;