From 7d0ff387e9cbe9b2cb8efec402dc365187030cd7 Mon Sep 17 00:00:00 2001
From: Marko Arandjelovic <marandje@amd.com>
Date: Wed, 15 May 2024 15:19:29 +0200
Subject: [PATCH] SWDEV-461791 make memcpy synchronous for D2D if src&dst ptrs
 have SYNC_MEMOPS attribute

Change-Id: I603081d21e5eb3c73111845e350d8fa2ba5a7733
---
 hipamd/src/hip_memory.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index f8bb441a34..90b01f43a3 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -537,7 +537,8 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
     hipMemoryType dstMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) &
         dstMemory->getMemFlags())? hipMemoryTypeHost : hipMemoryTypeDevice;
     // Device to Device copies do not need to host side synchronization.
-    if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeDevice)) {
+    if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeDevice) &&
+        (!srcMemory->getUserData().sync_mem_ops_ || !dstMemory->getUserData().sync_mem_ops_)) {
       isHostAsync = true;
     }
   }