From 0dd2e2b010c26a0e3f64c099b4aea1aad33afe84 Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Fri, 19 Jun 2020 18:04:25 -0400
Subject: [PATCH] Disable flush on marker

Change-Id: I04d0fd995fb4b29a6bc372e1543902a053017d6a


[ROCm/clr commit: 265d84bab9f4475b22eec4881ca5cab4c5bec108]
---
 projects/clr/hipamd/rocclr/hip_event.cpp    | 6 +++---
 projects/clr/hipamd/rocclr/hip_internal.hpp | 3 +++
 projects/clr/hipamd/rocclr/hip_stream.cpp   | 4 ++--
 3 files changed, 8 insertions(+), 5 deletions(-)

diff --git a/projects/clr/hipamd/rocclr/hip_event.cpp b/projects/clr/hipamd/rocclr/hip_event.cpp
index 7fd84c0fb7..d4aca089ed 100644
--- a/projects/clr/hipamd/rocclr/hip_event.cpp
+++ b/projects/clr/hipamd/rocclr/hip_event.cpp
@@ -99,7 +99,7 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) {
     // Events are the same, which indicates the stream is empty and likely
     // eventRecord is called on another stream. For such cases insert and measure a
     // marker.
-    amd::Command* command = new amd::Marker(*event_->command().queue(), false);
+    amd::Command* command = new amd::Marker(*event_->command().queue(), kMarkerDisableFlush);
     command->enqueue();
     command->awaitCompletion();
     ms = static_cast<float>(static_cast<int64_t>(command->event().profilingInfo().end_ -
@@ -127,7 +127,7 @@ hipError_t Event::streamWait(amd::HostQueue* hostQueue, uint flags) {
   amd::Command::EventWaitList eventWaitList;
   eventWaitList.push_back(event_);
 
-  amd::Command* command = new amd::Marker(*hostQueue, false, eventWaitList);
+  amd::Command* command = new amd::Marker(*hostQueue, kMarkerDisableFlush, eventWaitList);
   if (command == NULL) {
     return hipErrorOutOfMemory;
   }
@@ -240,7 +240,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
   amd::HostQueue* queue = hip::getQueue(stream);
   amd::Command* command = queue->getLastQueuedCommand(true);
   if (command == nullptr) {
-    command = new amd::Marker(*queue, false);
+    command = new amd::Marker(*queue, kMarkerDisableFlush);
     command->enqueue();
   }
 
diff --git a/projects/clr/hipamd/rocclr/hip_internal.hpp b/projects/clr/hipamd/rocclr/hip_internal.hpp
index de41120f79..6dbf77baed 100755
--- a/projects/clr/hipamd/rocclr/hip_internal.hpp
+++ b/projects/clr/hipamd/rocclr/hip_internal.hpp
@@ -242,4 +242,7 @@ extern bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, vo
 
 constexpr bool kOptionChangeable = true;
 constexpr bool kNewDevProg = false;
+
+constexpr bool kMarkerDisableFlush = true;   //!< Avoids command batch flush in ROCclr
+
 #endif // HIP_SRC_HIP_INTERNAL_H
diff --git a/projects/clr/hipamd/rocclr/hip_stream.cpp b/projects/clr/hipamd/rocclr/hip_stream.cpp
index d764e1a5fe..8d3d3e4340 100755
--- a/projects/clr/hipamd/rocclr/hip_stream.cpp
+++ b/projects/clr/hipamd/rocclr/hip_stream.cpp
@@ -161,7 +161,7 @@ void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream
 
   // Check if we have to wait anything
   if (eventWaitList.size() > 0) {
-    amd::Command* command = new amd::Marker(*blocking_queue, false, eventWaitList);
+    amd::Command* command = new amd::Marker(*blocking_queue, kMarkerDisableFlush, eventWaitList);
     if (command != nullptr) {
       command->enqueue();
       command->release();
@@ -322,7 +322,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
   amd::Command* command = hostQueue->getLastQueuedCommand(true);
   if (command == nullptr) {
     amd::Command::EventWaitList eventWaitList;
-    command = new amd::Marker(*hostQueue, false, eventWaitList);
+    command = new amd::Marker(*hostQueue, kMarkerDisableFlush, eventWaitList);
     command->enqueue();
   }
   amd::Event& event = command->event();