From bb636d4fcabdce3ebf1f7e8bfdfb67d1dee1f0c8 Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Fri, 16 Apr 2021 17:47:57 -0400 Subject: [PATCH] SWDEV-278896 - Increase thresholds to match MT behavior MT doesn't use GPU waits, but CPU for sync between engines. Change the threshold values for CPU waits for direct dispatch. That will bring behavior closer to MT. Change-Id: Ia41c3cb812614962aff2746b6cf858f1bf77dda2 [ROCm/clr commit: ca2ea70a6c25c27ae65d43fcf10c2a0b64b590a2] --- projects/clr/rocclr/device/rocm/rocvirtual.cpp | 4 ++-- projects/clr/rocclr/device/rocm/rocvirtual.hpp | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) 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,