From e430506f706c71fe7fc93115079aff4cade4542f 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 --- hipamd/rocclr/hip_memory.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/hipamd/rocclr/hip_memory.cpp b/hipamd/rocclr/hip_memory.cpp index b7596e8b0f..c909928a05 100755 --- a/hipamd/rocclr/hip_memory.cpp +++ b/hipamd/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); } }