diff --git a/api/hip/hip_context.cpp b/api/hip/hip_context.cpp
index a4971d50d4..7d7fdb6743 100644
--- a/api/hip/hip_context.cpp
+++ b/api/hip/hip_context.cpp
@@ -69,7 +69,8 @@ amd::HostQueue* getNullStream() {
   auto stream = g_nullStreams.find(getCurrentContext());
   if (stream == g_nullStreams.end()) {
     amd::Device* device = getCurrentContext()->devices()[0];
-    amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, 0,
+    cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE;
+    amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, properties,
                                                amd::CommandQueue::RealTimeDisabled,
                                                amd::CommandQueue::Priority::Normal);
     g_nullStreams[getCurrentContext()] = queue;
@@ -159,6 +160,9 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
+  // Release last tracked command
+  hip::getNullStream()->setLastQueuedCommand(nullptr);
+
   // Need to remove the ctx of calling thread if its the top one
   if (!g_ctxtStack.empty() && g_ctxtStack.top() == amdContext) {
     g_ctxtStack.pop();
diff --git a/api/hip/hip_event.cpp b/api/hip/hip_event.cpp
index d01273cdb5..0bde2150f8 100644
--- a/api/hip/hip_event.cpp
+++ b/api/hip/hip_event.cpp
@@ -139,9 +139,8 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
   } else {
     e->stream_ = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
-  amd::Command* command = (e->flags & hipEventDisableTiming)? new amd::Marker(*e->stream_, true) :
-    new hip::TimerMarker(*e->stream_);
-  command->enqueue();
+
+  amd::Command* command = e->stream_->getLastQueuedCommand(true);
 
   if (e->event_ != nullptr) {
     e->event_->release();
diff --git a/api/hip/hip_stream.cpp b/api/hip/hip_stream.cpp
index d7cd42bb51..af00b0f104 100644
--- a/api/hip/hip_stream.cpp
+++ b/api/hip/hip_stream.cpp
@@ -43,7 +43,8 @@ void syncStreams() {
 static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
   amd::Device* device = hip::getCurrentContext()->devices()[0];
 
-  amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, 0,
+  cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE;
+  amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, properties,
                                              amd::CommandQueue::RealTimeDisabled,
                                              amd::CommandQueue::Priority::Normal);
 
@@ -135,6 +136,10 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
   amd::ScopedLock lock(streamSetLock);
 
   amd::HostQueue* hostQueue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
+
+  // Release last tracked command
+  hostQueue->setLastQueuedCommand(nullptr);
+
   hostQueue->release();
   streamSet.erase(hostQueue);