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

@@ -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();
   }