From 9864d6e56c7cf12e45b77e3268a6c4a47e38208a Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Thu, 11 Jan 2024 07:03:40 +0000 Subject: [PATCH] SWDEV-301667 - Missing marker for hipExtLaunch* optimization - Use isFenceDirty() check in hipStreamQuery. In case if there is a stopEvent for hipExtLaunch* API we avoid system scope if there are consecutive hipExtLaunch* and replace them with agent scope, but then if one uses hipStreamQuery it should detect we need a barrier with system scope. Change-Id: I539810f659d4e1c49104e3a0da6ff2803c6b3dee [ROCm/clr commit: 4b3c4709ffa94f6bd13e95ad365059a134af00c0] --- projects/clr/hipamd/src/hip_stream.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index fad39e1692..c8ff4bd269 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -592,6 +592,14 @@ hipError_t hipStreamQuery_common(hipStream_t stream) { bool wait = (stream == nullptr) ? true : false; hip::Stream* hip_stream = hip::getStream(stream, wait); + if (hip_stream->vdev()->isFenceDirty()) { + amd::Command* command = new amd::Marker(*hip_stream, kMarkerDisableFlush); + if (command != nullptr) { + command->enqueue(); + command->release(); + } + } + amd::Command* command = hip_stream->getLastQueuedCommand(true); if (command == nullptr) { // Nothing was submitted to the queue @@ -602,6 +610,7 @@ hipError_t hipStreamQuery_common(hipStream_t stream) { if (command->type() != 0) { event.notifyCmdQueue(); } + // Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status bool ready = command->queue()->device().IsHwEventReady(event); if (!ready) {