Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe

src/hip_event.cpp

@@ -45,10 +45,10 @@ ihipEvent_t::ihipEvent_t(unsigned flags)
 void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf, 
                                            hipStream_t stream, ihipEventType_t eventType)
 {
-    _state  = hipEventStatusRecording;
     _marker = *cf;
     _type   = eventType;
     _stream = stream;
+    _state  = hipEventStatusRecording;
 }
 
 
@@ -157,13 +157,12 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
             event->_state = hipEventStatusComplete;
             return ihipLogStatus(hipSuccess);
         } else {
-            event->_state  = hipEventStatusRecording;
             // Clear timestamps
             event->_timestamp = 0;
 
             // Record the event in the stream:
             stream->locked_recordEvent(event);
-
+            event->_state  = hipEventStatusRecording;
             return ihipLogStatus(hipSuccess);
         }
     } else {
@@ -191,6 +190,10 @@ hipError_t hipEventSynchronize(hipEvent_t event)
 {
     HIP_INIT_SPECIAL_API(TRACE_SYNC, event);
 
+    if (!(event->_flags & hipEventReleaseToSystem)) {
+        tprintf(DB_WARN, "hipEventSynchronize on event without system-scope fence ; consider creating with hipEventReleaseToSystem\n");
+    }
+
     if (event) {
         if (event->_state == hipEventStatusUnitialized) {
             return ihipLogStatus(hipErrorInvalidResourceHandle);
@@ -268,6 +271,10 @@ hipError_t hipEventQuery(hipEvent_t event)
 {
     HIP_INIT_SPECIAL_API(TRACE_QUERY, event);
 
+    if (!(event->_flags & hipEventReleaseToSystem)) {
+        tprintf(DB_WARN, "hipEventQuery on event without system-scope fence ; consider creating with hipEventReleaseToSystem\n");
+    }
+
     if ((event->_state == hipEventStatusRecording) && !event->locked_isReady()) {
         return ihipLogStatus(hipErrorNotReady);
     } else {

src/hip_hcc_internal.h

@@ -218,7 +218,8 @@ extern const char *API_COLOR_END;
 #define DB_SYNC   1 /* 0x02 - trace synchronization pieces */
 #define DB_MEM    2 /* 0x04 - trace memory allocation / deallocation */
 #define DB_COPY   3 /* 0x08 - trace memory copy and peer commands. . */
-#define DB_MAX_FLAG 4
+#define DB_WARN   4 /* 0x10 - warn about sub-optimal or shady behavior */
+#define DB_MAX_FLAG 5
 // When adding a new debug flag, also add to the char name table below.
 //
 //
@@ -235,6 +236,7 @@ static const DbName dbName [] =
     {KYEL, "sync"},
     {KCYN, "mem"},
     {KMAG, "copy"},
+    {KRED, "warn"},
 };
 
 

src/hip_stream.cpp

@@ -103,7 +103,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
             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.
+            // This will use create_blocking_marker to wait on the specified queue.
             stream->locked_streamWaitEvent(event);
         }