From bbb635bc32a88bd6fe949687631c26d27e33afaf Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Mon, 1 Nov 2021 09:44:33 -0400 Subject: [PATCH] SWDEV-305016 - Add a timeout wait into IsHwEventReady() Just signal check will still submit the marker and then later runtime will have a timeout, but the barrier packet is still generated. Hence early timeout will allow to skip the marker. Change-Id: Ieb7d89becbcff43a4f4c46715354ca65ab4a80b9 --- rocclr/device/rocm/rocdevice.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index dc609b45f9..5577167822 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -2565,10 +2565,10 @@ bool Device::IsHwEventReady(const amd::Event& event, bool wait) const { ClPrint(amd::LOG_INFO, amd::LOG_SIG, "No HW event"); return false; } else if (wait) { - WaitForSignal(reinterpret_cast(hw_event)->signal_, ActiveWait()); - return true; + return WaitForSignal(reinterpret_cast(hw_event)->signal_, ActiveWait()); } - return (hsa_signal_load_relaxed(reinterpret_cast(hw_event)->signal_) <= 0); + static constexpr bool Timeout = true; + return WaitForSignal(reinterpret_cast(hw_event)->signal_); } // ================================================================================================