From 993f52cd8b3894fb2589f3b5da3abd38039d4b34 Mon Sep 17 00:00:00 2001
From: Sarbojit Sarkar <Sarbojit.Sarkar@amd.com>
Date: Wed, 31 Mar 2021 00:26:23 -0700
Subject: [PATCH] SWDEV-254441 - Fixed review comments

Change-Id: Ica13f23638be04216d152821127b2b7316f96af0


[ROCm/clr commit: 35c8493237a7cfe3c8620e1fef6387296a9c1e11]
---
 projects/clr/hipamd/rocclr/hip_memory.cpp | 18 ++++++++++++++++++
 1 file changed, 18 insertions(+)

diff --git a/projects/clr/hipamd/rocclr/hip_memory.cpp b/projects/clr/hipamd/rocclr/hip_memory.cpp
index de29366213..e28b9e05d8 100755
--- a/projects/clr/hipamd/rocclr/hip_memory.cpp
+++ b/projects/clr/hipamd/rocclr/hip_memory.cpp
@@ -204,6 +204,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
     return hipErrorInvalidValue;
   }
 
+  amd::HostQueue* newQueue = nullptr;
   amd::Command* command = nullptr;
   amd::Command::EventWaitList waitList;
 
@@ -292,6 +293,9 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
       }
       command = new amd::CopyMemoryCommand(*pQueue, CL_COMMAND_COPY_BUFFER, waitList,
           *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
+      if (pQueue != &queue) {
+        newQueue = pQueue;
+      }
     }
   }
   if (command == nullptr) {
@@ -301,6 +305,20 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   command->enqueue();
   if (!isAsync) {
     command->awaitCompletion();
+  } else {
+    if (newQueue != nullptr) {
+      amd::Command::EventWaitList waitList;
+      amd::Command* cmd = newQueue->getLastQueuedCommand(true);
+      if (cmd != nullptr) {
+        waitList.push_back(cmd);
+        amd::Command* depdentMarker = new amd::Marker(queue, true, waitList);
+        if (depdentMarker != nullptr) {
+          depdentMarker->enqueue();
+          depdentMarker->release();
+        }
+        cmd->release();
+      }
+    }
   }
   command->release();