diff --git a/projects/clr/hipamd/api/hip/hip_internal.hpp b/projects/clr/hipamd/api/hip/hip_internal.hpp
index eae35ecae9..5d32f1049f 100644
--- a/projects/clr/hipamd/api/hip/hip_internal.hpp
+++ b/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) {
diff --git a/projects/clr/hipamd/api/hip/hip_stream.cpp b/projects/clr/hipamd/api/hip/hip_stream.cpp
index 59456303fe..cbba339efb 100644
--- a/projects/clr/hipamd/api/hip/hip_stream.cpp
+++ b/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,