diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index c85a4851c0..3adcf81155 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -475,9 +475,13 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal( // ================================================================================================ std::vector& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngine engine) { bool explicit_wait = false; + bool sdma_wait = false; // Reset all current waiting signals waiting_signals_.clear(); + if(engine != HwQueueEngine::Compute) + sdma_wait = true; + // Does runtime switch the active engine? if (engine != engine_) { // Yes, return the signal from the previous operation for a wait @@ -518,7 +522,7 @@ std::vector& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) { const Settings& settings = gpu_.dev().settings(); // Actively wait on CPU to avoid extra overheads of signal tracking on GPU - if (!WaitForSignal(external_signals_[i]->signal_)) { + if (!WaitForSignal(external_signals_[i]->signal_, false, sdma_wait)) { if (settings.cpu_wait_for_signal_) { // Wait on CPU for completion if requested CpuWaitForSignal(external_signals_[i]); diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index cffc4a329e..512419a3d0 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.hpp @@ -44,15 +44,20 @@ constexpr static hsa_signal_value_t kInitSignalValueOne = 1; constexpr static uint64_t kTimeout100us = 100 * K; constexpr static uint64_t kUnlimitedWait = std::numeric_limits::max(); +// Active wait time out incase same sdma engine is used again, +// then just wait instead of adding dependency wait signal. +constexpr static uint64_t kSDMAEngineTimeout = 10; + template -inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false) { +inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool sdma_wait = false) { if (hsa_signal_load_relaxed(signal) > 0) { uint64_t timeout = kTimeout100us; if (active_wait) { timeout = kUnlimitedWait; } if (active_wait_timeout) { - timeout = ROC_ACTIVE_WAIT_TIMEOUT * K; + // If diff engine, wait to 10 ms. Otherwise no wait + timeout = (sdma_wait ? kSDMAEngineTimeout : ROC_ACTIVE_WAIT_TIMEOUT) * K; if (timeout == 0) { return false; }