From 0d993a5cfc7096f60aa9810d47e484a58f8a197a Mon Sep 17 00:00:00 2001 From: jujiang Date: Mon, 20 Jul 2020 12:03:04 -0400 Subject: [PATCH] SWDEV-242344-hipMemcpyAll failed on MGPU Change-Id: I01671e322c84213964d3d48bd6b5dc704dba731a [ROCm/hip commit: 448008348e78f3d3a3015f6575f24d39272f9280] --- projects/hip/rocclr/hip_memory.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/projects/hip/rocclr/hip_memory.cpp b/projects/hip/rocclr/hip_memory.cpp index cf3d17b574..b7596e8b0f 100755 --- a/projects/hip/rocclr/hip_memory.cpp +++ b/projects/hip/rocclr/hip_memory.cpp @@ -178,9 +178,12 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin *srcMemory->asBuffer(), sOffset, sizeBytes, dst); isAsync = false; } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - // Check if the queue device doesn't match the device on any memory object. Hence - // it's a P2P transfer, because the app has requested access to another GPU - if (srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) { + // Check if the queue device doesn't match the device on any memory object. + // And any of them are not host allocation. + // Hence it's a P2P transfer, because the app has requested access to another GPU + if ((srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) && + ((srcMemory->getContext().devices().size() == 1) && + (dstMemory->getContext().devices().size() == 1))) { command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList, *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes); if (command == nullptr) {