From 4b3f70a6b44d1f4e7fee2c3c58ea8af80d19efd1 Mon Sep 17 00:00:00 2001
From: Jaydeep Patel <jaydeepkumar.patel@amd.com>
Date: Mon, 6 Feb 2023 17:27:31 +0000
Subject: [PATCH] SWDEV-380405 - Make hipMemcpyPeer async w.r.t. host and let
 dst device wait till copy finishes by src device.

Change-Id: Idbe25404d0c31b93436e190b548f12f327f47a04


[ROCm/clr commit: 5f231c55e3e4924f76ee8b3ed64f67c499977a38]
---
 projects/clr/hipamd/src/hip_memory.cpp | 15 +++++++++++++++
 1 file changed, 15 insertions(+)

diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index f3fcf68f1e..3eb2ad1e0e 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -464,6 +464,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   if (src == dst && kind == hipMemcpyDefault) {
     return hipSuccess;
   }
+  bool isP2P = false;
   size_t sOffset = 0;
   amd::Memory* srcMemory = getMemoryObject(src, sOffset);
   size_t dOffset = 0;
@@ -475,6 +476,11 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
     isAsync = false;
   } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
     isAsync = false;
+  } else if ((srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) &&
+             (srcMemory->getContext().devices().size() == 1) &&
+             (dstMemory->getContext().devices().size() == 1)) {
+    isAsync = true;
+    isP2P = true;
   }
   amd::Command* command = nullptr;
   status = ihipMemcpyCommand(command, dst, src, sizeBytes, kind, queue, isAsync);
@@ -484,6 +490,15 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   command->enqueue();
   if (!isAsync) {
     command->awaitCompletion();
+  } else if (isP2P) {
+    amd::HostQueue* pQueue = hip::getNullStream(dstMemory->getContext());
+    amd::Command::EventWaitList waitList;
+    waitList.push_back(command);
+    amd::Command* depdentMarker = new amd::Marker(*pQueue, false, waitList);
+    if (depdentMarker != nullptr) {
+      depdentMarker->enqueue();
+      depdentMarker->release();
+    }
   } else {
     amd::HostQueue* newQueue = command->queue();
     if (newQueue != &queue) {