diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 048ccf35fa..e5c08233f3 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -394,8 +394,8 @@ hsa_signal_t* VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngine engine) { // Early signal status check if (hsa_signal_load_relaxed(prof_signal->signal_) > 0) { const Settings& settings = gpu_.dev().settings(); - // Actively wait on CPU for 50 us to avoid extra overheads of signal tracking on GPU - if (!WaitForSignal(prof_signal->signal_)) { + // Actively wait on CPU for 750 us to avoid extra overheads of signal tracking on GPU + if (!WaitForSignal(prof_signal->signal_)) { if (settings.cpu_wait_for_signal_) { // Wait on CPU for completion if requested CpuWaitForSignal(prof_signal); diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index 3799713e1b..dff3d841b4 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.hpp @@ -52,8 +52,8 @@ struct ProfilingSignal : public amd::HeapObject { constexpr static hsa_signal_value_t kInitSignalValueOne = 1; // Timeouts for HSA signal wait -constexpr static uint64_t kTimeout30us = 30000; -constexpr static uint64_t kTimeout50us = 50000; +constexpr static uint64_t kTimeout100us = 100000; +constexpr static uint64_t kTimeout750us = 750000; constexpr static uint64_t kUnlimitedWait = std::numeric_limits::max(); template @@ -64,7 +64,7 @@ inline bool WaitForSignal(hsa_signal_t signal) { return false; } } else { - uint64_t timeout = (ROC_ACTIVE_WAIT) ? kUnlimitedWait : kTimeout30us; + uint64_t timeout = (ROC_ACTIVE_WAIT) ? kUnlimitedWait : kTimeout100us; // Active wait with a timeout if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,