From e7c316bc92eadd019f9e4e98e52f65f5986b0929 Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Fri, 9 Aug 2019 16:34:38 -0400
Subject: [PATCH] P4 to Git Change 1981065 by cpaquot@cpaquot-ocl-lc-lnx on
 2019/08/09 16:31:44

	SWDEV-198556 - [HIP] Use src/dstMemory->getContext instead of host_context.
	Also relax the check for P2P copies in case of hipMemcpy(hostMalloced, hipMalloced(dev1), dev0)

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#67 edit


[ROCm/clr commit: 7bbbdabfb0e57a2975f653fa432b19731f8c2a50]
---
 projects/clr/hipamd/api/hip/hip_memory.cpp | 56 ++++++++++------------
 1 file changed, 25 insertions(+), 31 deletions(-)

diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp
index 177199d061..8ae5b36d3d 100644
--- a/projects/clr/hipamd/api/hip/hip_memory.cpp
+++ b/projects/clr/hipamd/api/hip/hip_memory.cpp
@@ -100,7 +100,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
     amd::HostQueue* pQueue = &queue;
     if (queueDevice != dstMemory->getContext().devices()[0]) {
-      pQueue = hip::getNullStream(*hip::host_context);
+      pQueue = hip::getNullStream(dstMemory->getContext());
     }
     command = new amd::WriteMemoryCommand(*pQueue, CL_COMMAND_WRITE_BUFFER, waitList,
               *dstMemory->asBuffer(), dOffset, sizeBytes, src);
@@ -108,43 +108,37 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
     amd::HostQueue* pQueue = &queue;
     if (queueDevice != srcMemory->getContext().devices()[0]) {
-      pQueue = hip::getNullStream(*hip::host_context);
+      pQueue = hip::getNullStream(srcMemory->getContext());
     }
     command = new amd::ReadMemoryCommand(*pQueue, CL_COMMAND_READ_BUFFER, waitList,
               *srcMemory->asBuffer(), sOffset, sizeBytes, dst);
     isAsync = false;
   } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
-    static const uint hostMem = CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR;
-    if ((kind == hipMemcpyDeviceToDevice ||
-         kind == hipMemcpyDefault) &&
-        ((srcMemory->getMemFlags() & hostMem) == 0) &&
-        ((dstMemory->getMemFlags() & hostMem) == 0)) {
-      if (queueDevice != srcMemory->getContext().devices()[0]) {
-        amd::Coord3D srcOffset(sOffset, 0, 0);
-        amd::Coord3D dstOffset(dOffset, 0, 0);
-        amd::Coord3D copySize(sizeBytes, 1, 1);
-        command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
-                  *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
-        command->enqueue();
-        if (!isAsync) {
-          command->awaitCompletion();
-        }
-        command->release();
-        return hipSuccess;
+    if (queueDevice != srcMemory->getContext().devices()[0]) {
+      amd::Coord3D srcOffset(sOffset, 0, 0);
+      amd::Coord3D dstOffset(dOffset, 0, 0);
+      amd::Coord3D copySize(sizeBytes, 1, 1);
+      command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
+                *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
+      command->enqueue();
+      if (!isAsync) {
+        command->awaitCompletion();
       }
-      if (queueDevice != dstMemory->getContext().devices()[0]) {
-        amd::Coord3D srcOffset(sOffset, 0, 0);
-        amd::Coord3D dstOffset(dOffset, 0, 0);
-        amd::Coord3D copySize(sizeBytes, 1, 1);
-        command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
-                  *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
-        command->enqueue();
-        if (!isAsync) {
-          command->awaitCompletion();
-        }
-        command->release();
-        return hipSuccess;
+      command->release();
+      return hipSuccess;
+    }
+    if (queueDevice != dstMemory->getContext().devices()[0]) {
+      amd::Coord3D srcOffset(sOffset, 0, 0);
+      amd::Coord3D dstOffset(dOffset, 0, 0);
+      amd::Coord3D copySize(sizeBytes, 1, 1);
+      command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
+                *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
+      command->enqueue();
+      if (!isAsync) {
+        command->awaitCompletion();
       }
+      command->release();
+      return hipSuccess;
     }
     command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
               *srcMemory->asBuffer(),*dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);