From b432ced4240431c4042d9cec3846cc12cd5bde9f Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Mon, 7 Aug 2023 13:46:02 -0400 Subject: [PATCH] SWDEV-408473 - Add wait time of 10 us if the waiting signal copy was < 24K. Change-Id: I438ec9eb07e5034042a4a9a5e6e51d74daba2c83 [ROCm/clr commit: 6f5277c70184298baa5fa1639d245285365386b4] --- projects/clr/rocclr/device/rocm/rocblit.cpp | 28 +++++++++++++++++-- projects/clr/rocclr/device/rocm/rocblit.hpp | 2 ++ projects/clr/rocclr/device/rocm/rocdevice.hpp | 20 +++++++++++-- .../clr/rocclr/device/rocm/rocvirtual.cpp | 20 ++++++------- .../clr/rocclr/device/rocm/rocvirtual.hpp | 8 +++--- 5 files changed, 59 insertions(+), 19 deletions(-) diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index befb7b6226..d470a1f997 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -696,8 +696,12 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, copyMask = kUseRegularCopyApi ? 0 : dev().fetchSDMAMask(this, true); } + // Check if host wait has to be forced + bool forceHostWait = forceHostWaitFunc(size[0]); + auto wait_events = gpu().Barriers().WaitingSignal(engine); - hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp()); + hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp(), + forceHostWait); if (!kUseRegularCopyApi && engine != HwQueueEngine::Unknown) { if (copyMask == 0) { @@ -2670,6 +2674,26 @@ amd::Memory* DmaBlitManager::pinHostMemory(const void* hostMem, size_t pinSize, return amdMemory; } +bool DmaBlitManager::forceHostWaitFunc(size_t copy_size) const { + // 10us wait is true for all other targets. + bool forceHostWait = true; + // Based on the profiled results, do not wait for copy size > 24 KB. + static constexpr size_t kGfx90aCopyThreshold = 24; + + if ((dev().isa().versionMajor() == 9 && dev().isa().versionMinor() == 0 + && dev().isa().versionStepping() == 10) && (copy_size >= kGfx90aCopyThreshold * Ki)) { + // Check if this is gfx90a and restrict small copy to 24K. + forceHostWait = false; + } else if ((dev().isa().versionMajor() == 9) && (dev().isa().versionMinor() == 4) + && (dev().isa().versionStepping() == 0 || dev().isa().versionStepping() == 1 + || dev().isa().versionStepping() == 2)) { + // for gfx940, gfx941, gfx942, dependency signal resolution is fast, so no Host wait at all. + forceHostWait = false; + } + + return forceHostWait; +} + Memory* KernelBlitManager::createView(const Memory& parent, cl_image_format format, cl_mem_flags flags) const { assert((parent.owner()->asBuffer() == nullptr) && "View supports images only"); @@ -2781,4 +2805,4 @@ bool KernelBlitManager::RunGwsInit( return result; } -} // namespace pal +} // namespace roc diff --git a/projects/clr/rocclr/device/rocm/rocblit.hpp b/projects/clr/rocclr/device/rocm/rocblit.hpp index d8a143a3e2..1eff72e0b4 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.hpp +++ b/projects/clr/rocclr/device/rocm/rocblit.hpp @@ -271,6 +271,8 @@ class DmaBlitManager : public device::HostBlitManager { address staging, //!< Staging resource bool hostToDev //!< True if data is copied from Host To Device ) const; + + bool forceHostWaitFunc(size_t copy_size) const; }; //! Kernel Blit Manager diff --git a/projects/clr/rocclr/device/rocm/rocdevice.hpp b/projects/clr/rocclr/device/rocm/rocdevice.hpp index edce714a78..763d941b4b 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.hpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.hpp @@ -81,14 +81,28 @@ public: hsa_signal_t signal_; //!< HSA signal to track profiling information Timestamp* ts_; //!< Timestamp object associated with the signal HwQueueEngine engine_; //!< Engine used with this signal - bool done_; //!< True if signal is done amd::Monitor lock_; //!< Signal lock for update + + typedef union { + struct { + uint32_t done_ : 1; //!< True if signal is done + uint32_t forceHostWait_ : 1; //!< Force Host Wait for dependency signals + uint32_t reserved_ : 30; + }; + uint32_t data_; + } Flags; + + Flags flags_; + ProfilingSignal() : ts_(nullptr) , engine_(HwQueueEngine::Compute) - , done_(true) , lock_("Signal Ops Lock", true) - { signal_.handle = 0; } + { + signal_.handle = 0; + flags_.done_ = true; + flags_.forceHostWait_ = true; + } virtual ~ProfilingSignal(); amd::Monitor& LockSignalOps() { return lock_; } diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 60415175e9..129e7466b2 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -151,7 +151,7 @@ void Timestamp::checkGpuTime() { ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Signal = (0x%lx), start = %ld, " "end = %ld time taken= %ld ns", it->signal_.handle, start, end, end - start); } - it->done_ = true; + it->flags_.done_ = true; } signals_.clear(); if (end != 0) { @@ -361,7 +361,7 @@ bool VirtualGPU::HwQueueTracker::Create() { // ================================================================================================ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal( - hsa_signal_value_t init_val, Timestamp* ts) { + hsa_signal_value_t init_val, Timestamp* ts, bool forceHostWait) { bool new_signal = false; // Peep signal +2 ahead to see if its done @@ -422,8 +422,9 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal( ProfilingSignal* prof_signal = signal_list_[current_id_]; // Reset the signal and return hsa_signal_silent_store_relaxed(prof_signal->signal_, init_val); - prof_signal->done_ = false; + prof_signal->flags_.done_ = false; prof_signal->engine_ = engine_; + prof_signal->flags_.forceHostWait_ = forceHostWait; if (ts != 0) { // Save HSA signal earlier to make sure the possible callback will have a valid // value for processing @@ -473,13 +474,9 @@ 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 @@ -498,6 +495,7 @@ std::vector& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi } } } + // Check if a wait is required if (explicit_wait) { bool skip_internal_signal = false; @@ -520,8 +518,10 @@ 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_, false, sdma_wait)) { + // Actively wait on CPU to avoid extra overheads of signal tracking on GPU. + // For small copies set forced wait + if (!WaitForSignal(external_signals_[i]->signal_, false, + external_signals_[i]->flags_.forceHostWait_)) { if (settings.cpu_wait_for_signal_) { // Wait on CPU for completion if requested CpuWaitForSignal(external_signals_[i]); @@ -556,7 +556,7 @@ bool VirtualGPU::HwQueueTracker::CpuWaitForSignal(ProfilingSignal* signal) { LogPrintfError("Failed signal [0x%lx] wait", signal->signal_); return false; } - signal->done_ = true; + signal->flags_.done_ = true; } return true; } diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index e34c075625..107dd56976 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.hpp @@ -46,7 +46,7 @@ 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 kForcedTimeout = 10; +constexpr static uint64_t kForcedTimeout10us = 10; template inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool forced_wait = false) { @@ -56,8 +56,8 @@ inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool fo timeout = kUnlimitedWait; } if (active_wait_timeout) { - // If diff engine, wait to 10 ms. Otherwise no wait - timeout = (forced_wait ? kForcedTimeout : ROC_ACTIVE_WAIT_TIMEOUT) * K; + // If forced wait is set, then wait for 10us, else dont wait. (ns * K = us) + timeout = (forced_wait ? kForcedTimeout10us : ROC_ACTIVE_WAIT_TIMEOUT) * K; if (timeout == 0) { return false; } @@ -232,7 +232,7 @@ class VirtualGPU : public device::VirtualDevice { //! Finds a free signal for the upcomming operation hsa_signal_t ActiveSignal(hsa_signal_value_t init_val = kInitSignalValueOne, - Timestamp* ts = nullptr); + Timestamp* ts = nullptr, bool forceHostWait = true); //! Wait for the curent active signal. Can idle the queue bool WaitCurrent() {