SWDEV-423317 - Enable GPU wait for hip sync calls

hipStreamSynchronize and hipDeviceSynchronize won't longer wait
for CPU commands in DD mode

Change-Id: I079c8bbfc34ddc6d3e2d74c92a34665877e512a5

hipamd/src/hip_device_runtime.cpp

@@ -509,7 +509,7 @@ hipError_t hipDeviceSetLimit ( hipLimit_t limit, size_t value ) {
 hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) {
   HIP_INIT_API(hipDeviceSetSharedMemConfig, config);
   if (config != hipSharedMemBankSizeDefault &&
-      config != hipSharedMemBankSizeFourByte && 
+      config != hipSharedMemBankSizeFourByte &&
       config != hipSharedMemBankSizeEightByte) {
     HIP_RETURN(hipErrorInvalidValue);
   }
@@ -520,7 +520,8 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) {
 
 hipError_t hipDeviceSynchronize() {
   HIP_INIT_API(hipDeviceSynchronize);
-  hip::Stream::SyncAllStreams(hip::getCurrentDevice()->deviceId());
+  constexpr bool kDontWaitForCpu = false;
+  hip::Stream::SyncAllStreams(hip::getCurrentDevice()->deviceId(), kDontWaitForCpu);
   HIP_RETURN(hipSuccess);
 }
 

hipamd/src/hip_internal.hpp

@@ -293,7 +293,7 @@ namespace hip {
     const std::vector<uint32_t> GetCUMask() const { return cuMask_; }
 
     /// Sync all streams
-    static void SyncAllStreams(int deviceId);
+    static void SyncAllStreams(int deviceId, bool cpu_wait = true);
 
     /// Check whether any blocking stream running
     static bool StreamCaptureBlocking();

hipamd/src/hip_stream.cpp

@@ -122,7 +122,7 @@ int Stream::DeviceId(const hipStream_t hStream) {
 }
 
 // ================================================================================================
-void Stream::SyncAllStreams(int deviceId) {
+void Stream::SyncAllStreams(int deviceId, bool cpu_wait) {
   // Make a local copy to avoid stalls for GPU finish with multiple threads
   std::vector<hip::Stream*> streams;
   streams.reserve(streamSet.size());
@@ -136,7 +136,7 @@ void Stream::SyncAllStreams(int deviceId) {
     }
   }
   for (auto it : streams) {
-    it->finish();
+    it->finish(cpu_wait);
     it->release();
   }
 }
@@ -442,8 +442,9 @@ hipError_t hipStreamSynchronize_common(hipStream_t stream) {
     }
   }
   bool wait = (stream == nullptr) ? true : false;
+  constexpr bool kDontWaitForCpu = false;
   // Wait for the current host queue
-  hip::getStream(stream, wait)->finish();
+  hip::getStream(stream, wait)->finish(kDontWaitForCpu);
   return hipSuccess;
 }