From 2d7912dc0135550589b3091f68e58828365708ef Mon Sep 17 00:00:00 2001
From: Saleel Kudchadker <Saleel.Kudchadker@amd.com>
Date: Thu, 2 May 2024 20:27:33 +0000
Subject: [PATCH] SWDEV-460080 - Do not wait for host side command complete

- awaitCompletion would wait for host side command compelete(aka
cpuWait). The correct way is to check the completion signal and if not
dispatch a marker that has a signal.

Change-Id: I0f4f23c7ea68c329bf1d5f05e9735f631e5e3808
---
 hipamd/src/hip_memory.cpp | 10 ++++------
 1 file changed, 4 insertions(+), 6 deletions(-)

diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index bfe00f6eca..edb7178e76 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -548,7 +548,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   }
   command->enqueue();
   if (!isHostAsync) {
-    command->awaitCompletion();
+    stream.finish();
   } else if (!isGPUAsync) {
     hip::Stream* pStream = hip::getNullStream(dstMemory->GetDeviceById()->context());
     amd::Command::EventWaitList waitList;
@@ -2349,9 +2349,7 @@ inline hipError_t ihipMemcpyCmdEnqueue(amd::Command* command, bool isAsync = fal
   }
   command->enqueue();
   if (!isAsync) {
-    if (!command->awaitCompletion()) {
-      status = hipErrorUnknown;
-    }
+    command->queue()->finish();
   }
   command->release();
   return status;
@@ -3217,7 +3215,7 @@ hipError_t ihipMemset(void* dst, int64_t value, size_t valueSize, size_t sizeByt
     for (auto command : commands) {
       command->enqueue();
       if (!isAsync) {
-        command->awaitCompletion();
+        hip_stream->finish();
       }
       command->release();
     }
@@ -3388,7 +3386,7 @@ hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent
   for (auto& command : commands) {
     command->enqueue();
     if (!isAsync) {
-      command->awaitCompletion();
+      hip_stream->finish();
     }
     command->release();
   }