SWDEV-486602 - Optimize HSA callback performance

- Don't generate callbacks for HIP events
- Don't process profiling info in the callback for HIP events
- Wait for CPU status update of the submitted commands
every 50 calls. That will allow to drain the commands and
destroy HSA signals.

Change-Id: Ib601a350e7e7c2b6c6209a172385389baccf73a9

hipamd/src/hip_stream.cpp

@@ -23,6 +23,7 @@
 #include "hip_event.hpp"
 #include "thread/monitor.hpp"
 #include "hip_prof_api.h"
+#include <atomic>
 
 namespace hip {
 
@@ -358,11 +359,15 @@ hipError_t hipStreamSynchronize_common(hipStream_t stream) {
     }
   }
   bool wait = (stream == nullptr || stream == hipStreamLegacy) ? true : false;
-  constexpr bool kDontWaitForCpu = false;
-
   auto hip_stream = hip::getStream(stream, wait);
+  bool wait_for_cpu = false;
+  // Force blocking wait if requested. That allows to avoid a build up of unreleased CPU commands
+  if (DEBUG_HIP_BLOCK_SYNC != 0) {
+    static std::atomic<uint64_t> flush = 0;
+    wait_for_cpu = ((++flush % DEBUG_HIP_BLOCK_SYNC) == 0) ? true : false;
+  }
   // Wait for the current host queue
-  hip_stream->finish(kDontWaitForCpu);
+  hip_stream->finish(wait_for_cpu);
   if (stream == nullptr) {
     // null stream will sync with other streams.
     ReleaseGraphExec(hip_stream->DeviceId());