From 133447a412996eb0d32cdb7d04cb8a3dbf6e5ec6 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Thu, 18 Oct 2018 13:41:32 -0400
Subject: [PATCH] P4 to Git Change 1620765 by
 skudchad@skudchad_test2_win_opencl on 2018/10/17 16:58:04

	SWDEV-145570 - [HIP] Track last used event and use last enqueued command in a stream rather than creating a new event

	ReviewBoardURL = http://ocltc.amd.com/reviews/r/15996/diff/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#15 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#7 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#14 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/command.cpp#90 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/commandqueue.cpp#28 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/commandqueue.hpp#20 edit
---
 api/hip/hip_context.cpp | 6 +++++-
 api/hip/hip_event.cpp   | 5 ++---
 api/hip/hip_stream.cpp  | 7 ++++++-
 3 files changed, 13 insertions(+), 5 deletions(-)

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);