From fae2466d946e86a10be19f050ff6aff5e028ca4a Mon Sep 17 00:00:00 2001 From: foreman Date: Thu, 8 Aug 2019 19:03:46 -0400 Subject: [PATCH] P4 to Git Change 1980493 by cpaquot@cpaquot-ocl-lc-lnx on 2019/08/08 18:57:21 SWDEV-198556 - [HIP] Override queue if src/dst's device doesn't match the queue's device. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#66 edit [ROCm/hip commit: 8fdb248e424a49af3ffb58aa6d908607d853b489] --- projects/hip/api/hip/hip_memory.cpp | 14 +++++++++++--- 1 file changed, 11 insertions(+), 3 deletions(-) diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp index 807874dc34..177199d061 100644 --- a/projects/hip/api/hip/hip_memory.cpp +++ b/projects/hip/api/hip/hip_memory.cpp @@ -91,17 +91,26 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin amd::Memory *srcMemory = getMemoryObject(src, sOffset); size_t dOffset = 0; amd::Memory *dstMemory = getMemoryObject(dst, dOffset); + amd::Device* queueDevice = &queue.device(); if (((srcMemory == nullptr) && (dstMemory == nullptr)) || (kind == hipMemcpyHostToHost)) { memcpy(dst, src, sizeBytes); return hipSuccess; } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - command = new amd::WriteMemoryCommand(queue, CL_COMMAND_WRITE_BUFFER, waitList, + amd::HostQueue* pQueue = &queue; + if (queueDevice != dstMemory->getContext().devices()[0]) { + pQueue = hip::getNullStream(*hip::host_context); + } + command = new amd::WriteMemoryCommand(*pQueue, CL_COMMAND_WRITE_BUFFER, waitList, *dstMemory->asBuffer(), dOffset, sizeBytes, src); isAsync = false; } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - command = new amd::ReadMemoryCommand(queue, CL_COMMAND_READ_BUFFER, waitList, + amd::HostQueue* pQueue = &queue; + if (queueDevice != srcMemory->getContext().devices()[0]) { + pQueue = hip::getNullStream(*hip::host_context); + } + command = new amd::ReadMemoryCommand(*pQueue, CL_COMMAND_READ_BUFFER, waitList, *srcMemory->asBuffer(), sOffset, sizeBytes, dst); isAsync = false; } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { @@ -110,7 +119,6 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin kind == hipMemcpyDefault) && ((srcMemory->getMemFlags() & hostMem) == 0) && ((dstMemory->getMemFlags() & hostMem) == 0)) { - amd::Device* queueDevice = &queue.device(); if (queueDevice != srcMemory->getContext().devices()[0]) { amd::Coord3D srcOffset(sOffset, 0, 0); amd::Coord3D dstOffset(dOffset, 0, 0);