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