Detect when an explicit printf buffer flush is required (#1766)

* Detect when an explicit printf buffer flush is required
in a device/stream synchronization function.

* hip_module.cpp: add missing hc_am.hpp header


[ROCm/hip commit: f4555c835a]

projects/hip/src/hip_hcc.cpp

@@ -304,18 +304,28 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t& crit) {
 
 //---
 // Wait for all kernel and data copy commands in this stream to complete.
-void ihipStream_t::locked_wait() {
+inline void ihipStream_t::locked_wait(bool& waited) {
     // create a marker while holding stream lock,
     // but release lock prior to waiting on the marker
     hc::completion_future marker;
     {
         LockedAccessor_StreamCrit_t crit(_criticalData);
         // skipping marker since stream is empty
-        if (crit->_av.get_is_empty()) return;
+        if (crit->_av.get_is_empty()) {
+            waited = false;
+            return;
+        }
         marker = crit->_av.create_marker(hc::no_scope);
     }
 
     marker.wait(waitMode());
+    waited = true;
+    return;
+};
+
+void ihipStream_t::locked_wait() {
+    bool waited;
+    locked_wait(waited);
 };
 
 // Causes current stream to wait for specified event to complete:
@@ -1040,6 +1050,7 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost) {
     // Vector of ops sent to each stream that will complete before ops sent to null stream:
     std::vector<hc::completion_future> depOps;
 
+    bool last_stream_waited = false;
     for (auto streamI = crit->const_streams().begin(); streamI != crit->const_streams().end();
          streamI++) {
         ihipStream_t* stream = *streamI;
@@ -1051,7 +1062,8 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost) {
 
         if (HIP_SYNC_NULL_STREAM) {
             if (waitThisStream) {
-                stream->locked_wait();
+                last_stream_waited = false;
+                stream->locked_wait(last_stream_waited);
             }
         } else {
             if (waitThisStream) {
@@ -1082,6 +1094,14 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost) {
             defaultCf.wait();  // TODO - account for active or blocking here.
         }
     }
+    else if ( (HIP_SYNC_NULL_STREAM && !last_stream_waited) ||
+              (!HIP_SYNC_NULL_STREAM && depOps.empty()) ) {
+        // This catches all the conditions where the printf buffer
+        // need to be explicitly flushed
+        if (syncHost) {
+            Kalmar::getContext()->flushPrintfBuffer();
+        }
+    }
 
     tprintf(DB_SYNC, "  syncDefaultStream depOps=%zu\n", depOps.size());
 }
@@ -1101,9 +1121,18 @@ void ihipCtx_t::locked_waitAllStreams() {
     LockedAccessor_CtxCrit_t crit(_criticalData);
 
     tprintf(DB_SYNC, "waitAllStream\n");
+    bool need_printf_flush = false;
     for (auto streamI = crit->const_streams().begin(); streamI != crit->const_streams().end();
          streamI++) {
-        (*streamI)->locked_wait();
+        bool waited;
+        (*streamI)->locked_wait(waited);
+        need_printf_flush = !waited;
+    }
+
+    // When synchronizing with the last stream, if we didn't do an explicit wait,
+    // then we need to an extra flush to the printf buffer
+    if (need_printf_flush) {
+        Kalmar::getContext()->flushPrintfBuffer();
     }
 }
 
@@ -1492,7 +1521,11 @@ hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream) {
         ctx->locked_syncDefaultStream(true /*waitOnSelf*/, true /*syncToHost*/);
     } else {
         // note this does not synchornize with the NULL stream:
-        stream->locked_wait();
+        bool waited;
+        stream->locked_wait(waited);
+        if (!waited) {
+            Kalmar::getContext()->flushPrintfBuffer();
+        }
         e = hipSuccess;
     }
 

projects/hip/src/hip_hcc_internal.h

@@ -570,7 +570,7 @@ class ihipStream_t {
     LockedAccessor_StreamCrit_t lockopen_preKernelCommand();
     void lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av, bool unlockNotNeeded = 0);
 
-
+    void locked_wait(bool& waited);
     void locked_wait();
 
     hc::accelerator_view* locked_getAv() {

projects/hip/src/hip_module.cpp

@@ -28,6 +28,7 @@ THE SOFTWARE.
 #include "hip/hip_ext.h"
 #include "program_state.inl"
 #include "trace_helper.h"
+#include "hc_am.hpp"
 
 #include <hsa/amd_hsa_kernel_code.h>
 #include <hsa/hsa.h>