From 56daa6c4891b43ec233e9c63f755e3f7b45842b4 Mon Sep 17 00:00:00 2001
From: taosang2 <tao.sang@amd.com>
Date: Thu, 29 Jun 2023 19:01:52 -0400
Subject: [PATCH] SWDEV-408687 - Prevent updating dstXInBytes/srcXInBytes twice

 Update pCopy->srcMemoryType and pCopy->dstMemoryType so
 that ihipGetMemcpyParam3DCommand() won't redetect memory
 type. This will prevent dstXInBytes and srcXInBytes getting
 updated twice with offset in hipMemoryTypeUnified type.

Change-Id: Ief4ea81941e25d22bcafd213b3b84a65eb23cf94
---
 hipamd/src/hip_memory.cpp | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index 46e73f96d2..1a4269c222 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -2200,6 +2200,8 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool
       // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system memory.
       const_cast<HIP_MEMCPY3D*>(pCopy)->srcHost = pCopy->srcDevice;
       const_cast<HIP_MEMCPY3D*>(pCopy)->srcXInBytes += offset;
+      // We don't need detect memory type again for hipMemoryTypeUnified
+      const_cast<HIP_MEMCPY3D*>(pCopy)->srcMemoryType = srcMemoryType;
     }
   }
   offset = 0;
@@ -2216,6 +2218,8 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool
     if (dstMemoryType == hipMemoryTypeHost) {
       const_cast<HIP_MEMCPY3D*>(pCopy)->dstHost = pCopy->dstDevice;
       const_cast<HIP_MEMCPY3D*>(pCopy)->dstXInBytes += offset;
+      // We don't need detect memory type again for hipMemoryTypeUnified
+      const_cast<HIP_MEMCPY3D*>(pCopy)->dstMemoryType = dstMemoryType;
     }
   }
   // If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned.