From 74269eb547a66ee49269dfb81600f1e0aaeedd95 Mon Sep 17 00:00:00 2001
From: Sarbojit Sarkar <sarbojit.sarkar@amd.com>
Date: Wed, 22 Dec 2021 11:52:24 +0000
Subject: [PATCH] SWDEV-316155 - Fix for P2P sample failure

Change-Id: I95a70d44949c59319a9a09f84be771740c3f0521


[ROCm/clr commit: 845b923507eb842e679500f5453484ab18dcda23]
---
 projects/clr/hipamd/src/hip_memory.cpp | 17 +++++++++++++++++
 1 file changed, 17 insertions(+)

diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index 04127fef4d..5e30be0a06 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -357,6 +357,23 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
     if ((srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) &&
         ((srcMemory->getContext().devices().size() == 1) &&
          (dstMemory->getContext().devices().size() == 1))) {
+      amd::HostQueue* currNullStream = hip::getNullStream();
+      if (currNullStream == &queue) {
+        if (queueDevice != srcMemory->getContext().devices()[0]) {
+          amd::HostQueue* pQueue = hip::getNullStream(srcMemory->getContext());
+          amd::Command* cmd = pQueue->getLastQueuedCommand(true);
+          if (cmd != nullptr) {
+            waitList.push_back(cmd);
+          }
+        }
+        if (queueDevice != dstMemory->getContext().devices()[0]) {
+          amd::HostQueue* pQueue = hip::getNullStream(dstMemory->getContext());
+          amd::Command* cmd = pQueue->getLastQueuedCommand(true);
+          if (cmd != nullptr) {
+            waitList.push_back(cmd);
+          }
+        }
+      }
       command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
           *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
       if (command == nullptr) {