From 5a8946190aaeb4e402667329d69309f2d6d3c015 Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Tue, 19 Jan 2021 18:06:08 -0500 Subject: [PATCH] SWDEV-268381 - Enable wait on CPU before SDMA transfer - There is a performance regression with a HW wait for HSA signal on ROCr async operation. For now move the logic back to CPU wait. - Fix profiling issue with multiple HSA signal per single timestamp object. Some copies require multiple ROCR calls and if profiling is required, then the execution time is derived from all used signals. Change-Id: Id003e4abb8c2de378eedc152a7e389500fc6f4ce --- rocclr/device/rocm/rocblit.cpp | 14 +++- rocclr/device/rocm/rocdefs.hpp | 7 ++ rocclr/device/rocm/rocvirtual.cpp | 9 ++- rocclr/device/rocm/rocvirtual.hpp | 110 ++++++++++++------------------ 4 files changed, 66 insertions(+), 74 deletions(-) diff --git a/rocclr/device/rocm/rocblit.cpp b/rocclr/device/rocm/rocblit.cpp index f5e1d5882a..f0fabd64ce 100644 --- a/rocclr/device/rocm/rocblit.cpp +++ b/rocclr/device/rocm/rocblit.cpp @@ -443,10 +443,12 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d if (isSubwindowRectCopy ) { hsa_signal_t wait = gpu().Barriers().WaitSignal(); hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp()); + uint32_t num_wait_events = (wait.handle == 0) ? 0 : 1; + hsa_signal_t* wait_event = (wait.handle == 0) ? nullptr : &wait; // Copy memory line by line hsa_status_t status = hsa_amd_memory_async_copy_rect(&dstMem, &offset, - &srcMem, &offset, &dim, agent, direction, 1, &wait, active); + &srcMem, &offset, &dim, agent, direction, num_wait_events, wait_event, active); if (status != HSA_STATUS_SUCCESS) { LogPrintfError("DMA buffer failed with code %d", status); return false; @@ -456,6 +458,8 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d const hsa_signal_value_t kInitVal = size[2] * size[1]; hsa_signal_t wait = gpu().Barriers().WaitSignal(); hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitVal, gpu().timestamp()); + uint32_t num_wait_events = (wait.handle == 0) ? 0 : 1; + hsa_signal_t* wait_event = (wait.handle == 0) ? nullptr : &wait; for (size_t z = 0; z < size[2]; ++z) { for (size_t y = 0; y < size[1]; ++y) { @@ -466,7 +470,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d hsa_status_t status = hsa_amd_memory_async_copy( (reinterpret_cast
(dst) + dstOffset), dstAgent, (reinterpret_cast(src) + srcOffset), srcAgent, - size[0], 1, &wait, active); + size[0], num_wait_events, wait_event, active); gpu().setLastCommandSDMA(true) ; if (status != HSA_STATUS_SUCCESS) { LogPrintfError("DMA buffer failed with code %d", status); @@ -640,8 +644,12 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, hsa_signal_t wait = gpu().Barriers().WaitSignal(); hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp()); + uint32_t num_wait_events = (wait.handle == 0) ? 0 : 1; + hsa_signal_t* wait_event = (wait.handle == 0) ? nullptr : &wait; + // Use SDMA to transfer the data - status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, size[0], 1, &wait, active); + status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, + size[0], num_wait_events, wait_event, active); gpu().setLastCommandSDMA(true); // Explicit wait for now, until runtime could distinguish compute and sdma operations gpu().Barriers().WaitCurrent(); diff --git a/rocclr/device/rocm/rocdefs.hpp b/rocclr/device/rocm/rocdefs.hpp index 889d95fc4a..28284d76e0 100644 --- a/rocclr/device/rocm/rocdefs.hpp +++ b/rocclr/device/rocm/rocdefs.hpp @@ -33,6 +33,13 @@ static constexpr uint DeviceQueueMaskSize = 32; //! Set to match the number of pipes, which is 8. static constexpr uint kMaxAsyncQueues = 8; +enum HwQueueEngine : uint32_t { + Compute = 0, + SdmaRead = 1, + SdmaWrite = 2, + Unknown = 3 +}; + } // namespace roc #endif diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 6cb93b208e..65b5d83fb9 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -871,7 +871,7 @@ void VirtualGPU::profilingBegin(amd::Command& command, bool drmProfiling) { */ void VirtualGPU::profilingEnd(amd::Command& command) { if (command.profilingInfo().enabled_) { - if (timestamp_->getProfilingSignal() == nullptr) { + if (!timestamp_->HwProfiling()) { timestamp_->end(); } command.setData(reinterpret_cast(timestamp_)); @@ -1186,6 +1186,8 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { // Initialize signal for the barrier hsa_signal_t wait = Barriers().WaitSignal(); hsa_signal_t active = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_); + uint32_t num_wait_events = (wait.handle == 0) ? 0 : 1; + hsa_signal_t* wait_event = (wait.handle == 0) ? nullptr : &wait; // Find the requested agent for the transfer hsa_agent_t agent = (cmd.cpu_access() || @@ -1194,7 +1196,7 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { // Initiate a prefetch command hsa_status_t status = hsa_amd_svm_prefetch_async( - const_cast(cmd.dev_ptr()), cmd.count(), agent, 1, &wait, active); + const_cast(cmd.dev_ptr()), cmd.count(), agent, num_wait_events, wait_event, active); // Wait for the prefetch. Should skip wait, but may require extra tracking for kernel execution. if ((status != HSA_STATUS_SUCCESS) || !Barriers().WaitCurrent()) { @@ -2134,9 +2136,6 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const dim = i; iteration = sizes.global()[i] / 0xC0000000 + ((sizes.global()[i] % 0xC0000000) ? 1 : 0); globalStep = (sizes.global()[i] / sizes.local()[i]) / iteration * sizes.local()[dim]; - if (timestamp_ != nullptr) { - timestamp_->setSplittedDispatch(); - } break; } } diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index d8e75cda2e..1c5146e93a 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -21,6 +21,7 @@ #pragma once #include "platform/commandqueue.hpp" +#include "rocdefs.hpp" #include "rocdevice.hpp" #include "utils/util.hpp" #include "hsa.h" @@ -36,11 +37,15 @@ class Memory; class Timestamp; struct ProfilingSignal : public amd::HeapObject { - hsa_signal_t signal_; //!< HSA signal to track profiling information - Timestamp* ts_; //!< Timestamp object associated with the signal - bool done_; //!< True if signal is done - - ProfilingSignal() : ts_(nullptr), done_(true) { signal_.handle = 0; } + 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 + ProfilingSignal() + : ts_(nullptr) + , engine_(HwQueueEngine::Compute) + , done_(true) + { signal_.handle = 0; } }; // Initial HSA signal value @@ -67,13 +72,12 @@ inline bool WaitForSignal(hsa_signal_t signal) { // including EnqueueNDRangeKernel and clEnqueueCopyBuffer. class Timestamp { private: - uint64_t start_; - uint64_t end_; - ProfilingSignal* profilingSignal_; - hsa_agent_t agent_; static double ticksToTime_; - bool splittedDispatch_; - std::vector splittedSignals_; + + uint64_t start_; + uint64_t end_; + hsa_agent_t agent_; + std::vector signals_; public: uint64_t getStart() { @@ -86,21 +90,15 @@ class Timestamp { return end_; } - void setProfilingSignal(ProfilingSignal* signal) { - profilingSignal_ = signal; - if (splittedDispatch_) { - splittedSignals_.push_back(profilingSignal_->signal_); - } - } - const ProfilingSignal* getProfilingSignal() const { return profilingSignal_; } + void AddProfilingSignal(ProfilingSignal* signal) { signals_.push_back(signal); } + + const bool HwProfiling() const { return (signals_.size() > 0) ? true : false; } void setAgent(hsa_agent_t agent) { agent_ = agent; } Timestamp() - : start_(0) - , end_(0) - , profilingSignal_(nullptr) - , splittedDispatch_(false) { + : start_(std::numeric_limits::max()) + , end_(0) { agent_.handle = 0; } @@ -108,51 +106,30 @@ class Timestamp { //! Finds execution ticks on GPU void checkGpuTime() { - if (profilingSignal_ != nullptr) { - hsa_amd_profiling_dispatch_time_t time; + if (HwProfiling()) { + hsa_amd_profiling_dispatch_time_t time = {}; - if (splittedDispatch_) { - uint64_t start = std::numeric_limits::max(); - uint64_t end = 0; - for (auto it = splittedSignals_.begin(); it < splittedSignals_.end(); it++) { - if (hsa_signal_load_relaxed(profilingSignal_->signal_) > 0) { - WaitForSignal(*it); - } - hsa_amd_profiling_get_dispatch_time(agent_, *it, &time); - if ((time.end - time.start) == 0) { - hsa_amd_profiling_async_copy_time_t time_sdma = {}; - hsa_amd_profiling_get_async_copy_time(profilingSignal_->signal_, &time_sdma); - time.start = time_sdma.start; - time.end = time_sdma.end; - } - if (time.start < start) { - start = time.start; - } - if (time.end > end) { - end = time.end; - } + uint64_t start = std::numeric_limits::max(); + uint64_t end = 0; + for (auto it : signals_) { + if (hsa_signal_load_relaxed(it->signal_) > 0) { + WaitForSignal(it->signal_); } - start_ = start * ticksToTime_; - end_ = end * ticksToTime_; - } else { - // If the signalValue is the same as initial set value, it means its not written to - if (hsa_signal_load_relaxed(profilingSignal_->signal_) > 0) { - WaitForSignal(profilingSignal_->signal_); - } - hsa_amd_profiling_get_dispatch_time(agent_, profilingSignal_->signal_, &time); + hsa_amd_profiling_get_dispatch_time(agent_, it->signal_, &time); if ((time.end - time.start) == 0) { hsa_amd_profiling_async_copy_time_t time_sdma = {}; - hsa_amd_profiling_get_async_copy_time(profilingSignal_->signal_, &time_sdma); - start_ = time_sdma.start * ticksToTime_; - end_ = time_sdma.end * ticksToTime_; - } else { - start_ = time.start * ticksToTime_; - end_ = time.end * ticksToTime_; + hsa_amd_profiling_get_async_copy_time(it->signal_, &time_sdma); + time.start = time_sdma.start; + time.end = time_sdma.end; } + start = std::min(time.start, start); + end = std::max(time.end, end); + it->ts_ = nullptr; + it->done_ = true; } - profilingSignal_->ts_ = nullptr; - profilingSignal_->done_ = true; - profilingSignal_ = nullptr; + signals_.clear(); + start_ = start * ticksToTime_; + end_ = end * ticksToTime_; } } @@ -162,9 +139,6 @@ class Timestamp { // End a timestamp (get timestamp from OS) void end() { end_ = amd::Os::timeNanos(); } - bool isSplittedDispatch() const { return splittedDispatch_; } - void setSplittedDispatch() { splittedDispatch_ = true; } - static void setGpuTicksToTime(double ticksToTime) { ticksToTime_ = ticksToTime; } static double getGpuTicksToTime() { return ticksToTime_; } }; @@ -268,7 +242,7 @@ class VirtualGPU : public device::VirtualDevice { sdma_profiling_ = true; } signal_list_[current_id_]->ts_ = ts; - ts->setProfilingSignal(signal_list_[current_id_]); + ts->AddProfilingSignal(signal_list_[current_id_]); ts->setAgent(agent_); } return signal_list_[current_id_]->signal_; @@ -278,7 +252,11 @@ class VirtualGPU : public device::VirtualDevice { bool WaitCurrent() { return WaitIndex(current_id_); } //! Returns the last submitted signal for a wait - hsa_signal_t WaitSignal() const { return signal_list_[current_id_]->signal_; } + hsa_signal_t WaitSignal() { + //! @note Currently wait on CPU unconditionally to avoid a negative performance impact + WaitCurrent(); + return hsa_signal_t{}; + } private: //! Wait for the next active signal