From 2bdfc73649c44288abe83852d3d5e68d8d265a5e Mon Sep 17 00:00:00 2001
From: Christophe Paquot <christophe.paquot@amd.com>
Date: Mon, 24 Feb 2020 13:28:08 -0800
Subject: [PATCH] Fixed a few multithreaded potential issues

Also make D2H and H2D keep track of the chain of events
when we need to use a different HostQueue.

Change-Id: I1c5da6ea6104b37ad7aac00f0eb8ea9371e6ba1c
---
 hipamd/vdi/hip_memory.cpp | 6 ++++++
 hipamd/vdi/hip_stream.cpp | 6 ++++--
 2 files changed, 10 insertions(+), 2 deletions(-)

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<hip::Stream*>(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,