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