From c3eb5ccc8e10b7612244a6e45ee6ba60cc4cab19 Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <Satyanvesh.Dittakavi@amd.com>
Date: Mon, 28 Dec 2020 11:27:09 -0500
Subject: [PATCH] SWDEV-263678 - fix rocHPCG mgpu perf drop

Change-Id: I33269fe4f6165d1ed172b19033ab9fcb92f338c3
---
 rocclr/hip_memory.cpp | 22 +++++++++++++++++++---
 1 file changed, 19 insertions(+), 3 deletions(-)

diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp
index 65e1c1682c..af4e368d2f 100755
--- a/rocclr/hip_memory.cpp
+++ b/rocclr/hip_memory.cpp
@@ -211,15 +211,31 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
       }
     } else {
       amd::HostQueue* pQueue = &queue;
-      if (((srcMemory->getContext().devices()[0] == dstMemory->getContext().devices()[0]) ||
-          (srcMemory->getContext().devices().size() != 1)  ||
-          (dstMemory->getContext().devices().size() != 1)) &&
+      if ((srcMemory->getContext().devices()[0] == dstMemory->getContext().devices()[0]) &&
           (queueDevice != srcMemory->getContext().devices()[0])) {
         pQueue = hip::getNullStream(srcMemory->getContext());
         amd::Command* cmd = queue.getLastQueuedCommand(true);
         if (cmd != nullptr) {
           waitList.push_back(cmd);
         }
+      } else if (srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) {
+        // Scenarios such as DtoH where dst is pinned memory
+        if ((queueDevice != srcMemory->getContext().devices()[0]) &&
+            (dstMemory->getContext().devices().size() != 1)) {
+          pQueue = hip::getNullStream(srcMemory->getContext());
+          amd::Command* cmd = queue.getLastQueuedCommand(true);
+          if (cmd != nullptr) {
+            waitList.push_back(cmd);
+          }
+        // Scenarios such as HtoD where src is pinned memory
+        } else if ((queueDevice != dstMemory->getContext().devices()[0]) &&
+                   (srcMemory->getContext().devices().size() != 1)) {
+          pQueue = hip::getNullStream(dstMemory->getContext());
+          amd::Command* cmd = queue.getLastQueuedCommand(true);
+          if (cmd != nullptr) {
+            waitList.push_back(cmd);
+          }
+        }
       }
       command = new amd::CopyMemoryCommand(*pQueue, CL_COMMAND_COPY_BUFFER, waitList,
           *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);