Disable flush on marker

Change-Id: I04d0fd995fb4b29a6bc372e1543902a053017d6a


[ROCm/clr commit: 265d84bab9]

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

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

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