diff --git a/hipamd/vdi/hip_memory.cpp b/hipamd/vdi/hip_memory.cpp index 63dd2c712e..58c27cb89d 100644 --- a/hipamd/vdi/hip_memory.cpp +++ b/hipamd/vdi/hip_memory.cpp @@ -102,6 +102,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin amd::HostQueue* pQueue = &queue; if (queueDevice != dstMemory->getContext().devices()[0]) { pQueue = hip::getNullStream(dstMemory->getContext()); + waitList.push_back(queue.getLastQueuedCommand(true)); } command = new amd::WriteMemoryCommand(*pQueue, CL_COMMAND_WRITE_BUFFER, waitList, *dstMemory->asBuffer(), dOffset, sizeBytes, src); @@ -110,6 +111,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin amd::HostQueue* pQueue = &queue; if (queueDevice != srcMemory->getContext().devices()[0]) { pQueue = hip::getNullStream(srcMemory->getContext()); + waitList.push_back(queue.getLastQueuedCommand(true)); } command = new amd::ReadMemoryCommand(*pQueue, CL_COMMAND_READ_BUFFER, waitList, *srcMemory->asBuffer(), sOffset, sizeBytes, dst); @@ -155,6 +157,10 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin } command->release(); + if (waitList.size() > 0) { + waitList[0]->release(); + } + return hipSuccess; } diff --git a/hipamd/vdi/hip_stream.cpp b/hipamd/vdi/hip_stream.cpp index 7b32801aa8..55a975cc5c 100644 --- a/hipamd/vdi/hip_stream.cpp +++ b/hipamd/vdi/hip_stream.cpp @@ -224,7 +224,7 @@ hipError_t hipStreamQuery(hipStream_t stream) { hostQueue = reinterpret_cast(stream)->asHostQueue(); } - amd::Command* command = hostQueue->getLastQueuedCommand(false); + amd::Command* command = hostQueue->getLastQueuedCommand(true); if (command == nullptr) { HIP_RETURN(hipSuccess); } @@ -233,7 +233,9 @@ hipError_t hipStreamQuery(hipStream_t stream) { if (command->type() != 0) { event.notifyCmdQueue(); } - HIP_RETURN((command->status() == CL_COMPLETE) ? hipSuccess : hipErrorNotReady); + hipError_t status = (command->status() == CL_COMPLETE) ? hipSuccess : hipErrorNotReady; + command->release(); + HIP_RETURN(status); } hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,