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

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

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

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