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