P4 to Git Change 1972922 by cpaquot@cpaquot-ocl-lc-lnx on 2019/07/24 11:59:16

SWDEV-196881 - [HIP] Fix hipStreamQuery implementation. Using isEmpty isn't correct.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#32 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#20 edit


[ROCm/clr commit: 47ac573d70]

projects/clr/hipamd/api/hip/hip_internal.hpp

@@ -159,6 +159,7 @@ extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset);
 
 #define HIP_RETURN(ret)          \
         hip::g_lastError = ret;  \
+        LogPrintfInfo("[%zx] %s: Returned %s", std::this_thread::get_id(), __func__, hipGetErrorName(ret)); \
         return hip::g_lastError; \
 
 inline std::ostream& operator<<(std::ostream& os, const dim3& s) {

projects/clr/hipamd/api/hip/hip_stream.cpp

@@ -210,7 +210,18 @@ hipError_t hipStreamQuery(hipStream_t stream) {
   } else {
     hostQueue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
-  HIP_RETURN(hostQueue->isEmpty() ? hipSuccess : hipErrorNotReady);
+
+  amd::Command* command = hostQueue->getLastQueuedCommand(false);
+  if (command == nullptr) {
+    HIP_RETURN(hipSuccess);
+  }
+
+  amd::Event& event = command->event();
+
+  if (command->type() != CL_COMMAND_MARKER) {
+    event.notifyCmdQueue();
+  }
+  HIP_RETURN((command->status() == CL_COMPLETE) ? hipSuccess : hipErrorNotReady);
 }
 
 hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,