From b416ad7b9d396c25120f713fce74da9cde7e500b Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Wed, 9 Jun 2021 00:08:02 -0700 Subject: [PATCH] SWDEV-247372 - Active wait timeout env var - Create an env var ROC_ACTIVE_WAIT_TIMEOUT to set active wait timeout - Record profiling informaion if marker_ts_ property is valid. Change-Id: If0d8aec8d9b0715027cf0f7c3dc8a4c722a6bae6 --- rocclr/device/rocm/rocvirtual.cpp | 9 ++++-- rocclr/device/rocm/rocvirtual.hpp | 46 +++++++++++++++++++------------ rocclr/utils/flags.hpp | 4 +-- 3 files changed, 37 insertions(+), 22 deletions(-) diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index e8e21a5c7f..e486c43bec 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -122,7 +122,7 @@ void Timestamp::checkGpuTime() { } // Avoid profiling data for the sync barrier, in tiny performance tests the first call // to ROCr is very slow and that also affects the overall performance of the callback thread - if (command().GetBatchHead() == nullptr) { + if (command().GetBatchHead() == nullptr || command().profilingInfo().marker_ts_) { hsa_amd_profiling_dispatch_time_t time = {}; if (it->engine_ == HwQueueEngine::Compute) { hsa_amd_profiling_get_dispatch_time(gpu()->gpu_device(), it->signal_, &time); @@ -132,8 +132,11 @@ void Timestamp::checkGpuTime() { time.start = time_sdma.start; time.end = time_sdma.end; } + start = std::min(time.start, start); end = std::max(time.end, end); + ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Signal = (0x%lx), start = %ld, " + "end = %ld", it->signal_.handle, start, end); } it->ts_ = nullptr; it->done_ = true; @@ -452,8 +455,8 @@ std::vector& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi // Early signal status check if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) { const Settings& settings = gpu_.dev().settings(); - // Actively wait on CPU for 750 us to avoid extra overheads of signal tracking on GPU - if (!WaitForSignal(external_signals_[i]->signal_)) { + // Actively wait on CPU to avoid extra overheads of signal tracking on GPU + if (!WaitForSignal(external_signals_[i]->signal_)) { if (settings.cpu_wait_for_signal_) { // Wait on CPU for completion if requested CpuWaitForSignal(external_signals_[i]); diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index f39d901bfb..4d40023d94 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -52,30 +52,42 @@ struct ProfilingSignal : public amd::HeapObject { constexpr static hsa_signal_value_t kInitSignalValueOne = 1; // Timeouts for HSA signal wait -constexpr static uint64_t kTimeout100us = 100000; -constexpr static uint64_t kTimeout750us = 750000; +constexpr static uint64_t kTimeout100us = 100 * K; constexpr static uint64_t kUnlimitedWait = std::numeric_limits::max(); -template +template inline bool WaitForSignal(hsa_signal_t signal) { - if (wait_time != 0) { - if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, - wait_time, HSA_WAIT_STATE_ACTIVE) != 0) { - return false; - } - } else { - uint64_t timeout = (ROC_ACTIVE_WAIT) ? kUnlimitedWait : kTimeout100us; - - // Active wait with a timeout - if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, - timeout, HSA_WAIT_STATE_ACTIVE) != 0) { - // Wait until the completion with CPU suspend - if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, - kUnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) { + if (hsa_signal_load_relaxed(signal) > 0) { + if (active_wait_timeout) { + uint64_t timeout = ROC_ACTIVE_WAIT_TIMEOUT * K; + if (timeout == 0) { return false; } + ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Host active wait for Signal = (0x%lx) for %d us", + signal.handle, ROC_ACTIVE_WAIT_TIMEOUT); + + if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, + timeout, HSA_WAIT_STATE_ACTIVE) != 0) { + return false; + } + } else { + + uint64_t timeout = kTimeout100us; + ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Host wait until Signal = (0x%lx) decremented", + signal.handle); + + // Active wait with a timeout + if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, + timeout, HSA_WAIT_STATE_ACTIVE) != 0) { + // Wait until the completion with CPU suspend + if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, + kUnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) { + return false; + } + } } } + return true; } diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp index 50c0b23f55..a9d7edf008 100644 --- a/rocclr/utils/flags.hpp +++ b/rocclr/utils/flags.hpp @@ -233,8 +233,8 @@ release(uint, HIP_HIDDEN_FREE_MEM, 0, \ "0 = Disable") \ release(size_t, GPU_FORCE_BLIT_COPY_SIZE, 0, \ "Size in KB of the threshold below which to force blit instead for sdma") \ -release(bool, ROC_ACTIVE_WAIT, false, \ - "Forces unconditional active wait for GPU") \ +release(uint, ROC_ACTIVE_WAIT_TIMEOUT, 750, \ + "Forces active wait of GPU interrup for the timeout(us)") \ release(bool, ROC_ENABLE_LARGE_BAR, true, \ "Enable Large Bar if supported by the device") \ release(bool, ROC_CPU_WAIT_FOR_SIGNAL, true, \