From 7bd577696b9ee01b988b87e468bde59ad420a5fd Mon Sep 17 00:00:00 2001 From: Jason Tang Date: Tue, 21 Jul 2020 12:10:43 -0400 Subject: [PATCH] SWDEV-243576 - Fix hipMemcpy regression. If the queue device doesn't match the device on any memory object, use the queue device from the memory object. Change-Id: I5fdcf00494f8391574f4443332c01788b8da44ef [ROCm/hip commit: 678b3eeed5289afc2d6d2d5f9b8d55742cb63682] --- projects/hip/rocclr/hip_memory.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/projects/hip/rocclr/hip_memory.cpp b/projects/hip/rocclr/hip_memory.cpp index b7596e8b0f..c909928a05 100755 --- a/projects/hip/rocclr/hip_memory.cpp +++ b/projects/hip/rocclr/hip_memory.cpp @@ -196,7 +196,16 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin return hipErrorInvalidValue; } } else { - command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList, + amd::HostQueue* pQueue = &queue; + if (queueDevice != srcMemory->getContext().devices()[0]) { + pQueue = hip::getNullStream(srcMemory->getContext()); + amd::Command* cmd = queue.getLastQueuedCommand(true); + if (cmd != nullptr) { + waitList.push_back(cmd); + } + } + + command = new amd::CopyMemoryCommand(*pQueue, CL_COMMAND_COPY_BUFFER, waitList, *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes); } }