From 27596aef0c95314009d5c9eab632aceb5a87c187 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Sun, 15 May 2022 22:04:45 -0500 Subject: [PATCH] Track size of pending operations in blits. Track and report the size, in bytes, of pending unexecuted blit commands. To be used in copy ganging. Change-Id: Ia7453ff88571e927df771c6c819b73c17e67708e --- .../hsa-runtime/core/inc/amd_blit_kernel.h | 28 ++++++ runtime/hsa-runtime/core/inc/amd_blit_sdma.h | 32 ++++++- runtime/hsa-runtime/core/inc/blit.h | 4 + .../core/runtime/amd_blit_kernel.cpp | 60 +++++++++++- .../core/runtime/amd_blit_sdma.cpp | 91 +++++++++++++++---- 5 files changed, 195 insertions(+), 20 deletions(-) diff --git a/runtime/hsa-runtime/core/inc/amd_blit_kernel.h b/runtime/hsa-runtime/core/inc/amd_blit_kernel.h index f261ffc131..ceecfcd15c 100644 --- a/runtime/hsa-runtime/core/inc/amd_blit_kernel.h +++ b/runtime/hsa-runtime/core/inc/amd_blit_kernel.h @@ -45,6 +45,8 @@ #include #include +#include +#include #include #include "core/inc/blit.h" @@ -109,6 +111,8 @@ class BlitKernel : public core::Blit { virtual hsa_status_t EnableProfiling(bool enable) override; + virtual uint64_t PendingBytes() override; + private: union KernelArgs { struct __ALIGNED__(16) { @@ -144,6 +148,12 @@ class BlitKernel : public core::Blit { } fill; }; + // Index after which bytes will have been written. + struct BytesWritten { + uint64_t index; + uint64_t bytes; + }; + /// Reserve a slot in the queue buffer. The call will wait until the queue /// buffer has a room. uint64_t AcquireWriteIndex(uint32_t num_packet); @@ -158,6 +168,8 @@ class BlitKernel : public core::Blit { KernelArgs* ObtainAsyncKernelCopyArg(); + void RecordBlitHistory(uint64_t size, uint64_t index); + /// AQL code object and size for each kernel. enum class KernelType { CopyAligned, @@ -184,6 +196,22 @@ class BlitKernel : public core::Blit { /// Completion signal for every kernel dispatched. hsa_signal_t completion_signal_; + /// Bytes moved by commands < index. + /// Any record's byte value may be inexact by the size of concurrently issued operations. + std::vector bytes_written_; + + /// Total bytes written by all commands issued. + uint64_t bytes_queued_; + + /// Index where most recent blit operation queued. + uint64_t last_queued_; + + /// Orders command indices and bytes_queued_ updates + std::mutex reservation_lock_; + + /// Search resume index + std::atomic pending_search_index_; + /// Lock to synchronize access to kernarg_ and completion_signal_ std::mutex lock_; diff --git a/runtime/hsa-runtime/core/inc/amd_blit_sdma.h b/runtime/hsa-runtime/core/inc/amd_blit_sdma.h index 5c8528054b..7e9c5e9ae3 100644 --- a/runtime/hsa-runtime/core/inc/amd_blit_sdma.h +++ b/runtime/hsa-runtime/core/inc/amd_blit_sdma.h @@ -143,6 +143,8 @@ class BlitSdma : public BlitSdmaBase { virtual hsa_status_t EnableProfiling(bool enable) override; + virtual uint64_t PendingBytes() override; + private: /// @brief Acquires the address into queue buffer where a new command /// packet of specified size could be written. The address that is @@ -212,11 +214,11 @@ class BlitSdma : public BlitSdmaBase { void BuildGCRCommand(char* cmd_addr, bool invalidate); - hsa_status_t SubmitCommand(const void* cmds, size_t cmd_size, + hsa_status_t SubmitCommand(const void* cmds, size_t cmd_size, uint64_t size, const std::vector& dep_signals, core::Signal& out_signal); - hsa_status_t SubmitBlockingCommand(const void* cmds, size_t cmd_size); + hsa_status_t SubmitBlockingCommand(const void* cmds, size_t cmd_size, uint64_t size); // Agent object owning the SDMA engine. GpuAgent* agent_; @@ -224,6 +226,32 @@ class BlitSdma : public BlitSdmaBase { /// Base address of the Queue buffer at construction time. char* queue_start_addr_; + // Pending bytes tracking + // bytes_written_ is indexed with wrapped command queue indices (which are in bytes). + // The data_ index corresponding to a command queue index is the first uint64_t index which begins + // in the packet area. All packets have a header & at least one address so must be larger than 12 + // bytes, thus this index always exists. + std::mutex reservation_lock_; + uint64_t bytes_queued_; + class { + public: + // Indexed by wrapped command queue indices (offsets). + uint64_t& operator[](uint32_t index) { return data_[convert(index)]; } + + void resize(size_t size) { data_.resize(convert(size)); } + + void fill(uint32_t start, uint32_t stop, uint64_t value) { + for (uint32_t i = convert(start); i < convert(stop); i++) { + data_[i] = value; + } + } + + private: + uint32_t convert(uint32_t index) { return (index + sizeof(uint64_t) - 1) / sizeof(uint64_t); } + + std::vector data_; + } bytes_written_; + // Internal signals for blocking APIs core::unique_signal_ptr signals_[2]; KernelMutex lock_; diff --git a/runtime/hsa-runtime/core/inc/blit.h b/runtime/hsa-runtime/core/inc/blit.h index ad663c36dc..9e927755bf 100644 --- a/runtime/hsa-runtime/core/inc/blit.h +++ b/runtime/hsa-runtime/core/inc/blit.h @@ -110,6 +110,10 @@ class Blit { /// @brief Blit operations use SDMA. virtual bool isSDMA() const { return false; } + + /// @Brief Reports the approximate number of remaining bytes to copy or fill. Any return of zero + /// must be exact. + virtual uint64_t PendingBytes() = 0; }; } // namespace core } // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp b/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp index 2461e05f46..681b310a19 100644 --- a/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp @@ -522,6 +522,9 @@ BlitKernel::BlitKernel(core::Queue* queue) kernarg_async_(NULL), kernarg_async_mask_(0), kernarg_async_counter_(0), + bytes_queued_(0), + last_queued_(0), + pending_search_index_(0), num_cus_(0) { completion_signal_.handle = 0; } @@ -531,6 +534,9 @@ BlitKernel::~BlitKernel() {} hsa_status_t BlitKernel::Initialize(const core::Agent& agent) { queue_bitmask_ = queue_->public_handle()->size - 1; + bytes_written_.resize(queue_->public_handle()->size); + memset(&bytes_written_[0], -1, bytes_written_.size() * sizeof(BytesWritten)); + hsa_status_t status = HSA::hsa_signal_create(1, 0, NULL, &completion_signal_); if (HSA_STATUS_SUCCESS != status) { return status; @@ -619,7 +625,13 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand( const uint32_t num_barrier_packet = uint32_t((dep_signals.size() + 4) / 5); const uint32_t total_num_packet = num_barrier_packet + 1; - uint64_t write_index = AcquireWriteIndex(total_num_packet); + uint64_t write_index; + { + std::lock_guard lock(reservation_lock_); + write_index = AcquireWriteIndex(total_num_packet); + RecordBlitHistory(size, write_index + total_num_packet - 1); + } + uint64_t write_index_temp = write_index; // Insert barrier packets to handle dependent signals. @@ -758,9 +770,16 @@ hsa_status_t BlitKernel::SubmitLinearFillCommand(void* ptr, uint32_t value, // Submit dispatch packet. HSA::hsa_signal_store_relaxed(completion_signal_, 1); - uint64_t write_index = AcquireWriteIndex(1); + uint64_t write_index; + { + std::lock_guard lock(reservation_lock_); + write_index = AcquireWriteIndex(1); + RecordBlitHistory(fill_size, write_index); + } + PopulateQueue(write_index, uintptr_t(kernels_[KernelType::Fill].code_buf_), args, num_workitems, completion_signal_); + ReleaseWriteIndex(write_index, 1); // Wait for the packet to finish. @@ -842,5 +861,42 @@ BlitKernel::KernelArgs* BlitKernel::ObtainAsyncKernelCopyArg() { return arg; } +void BlitKernel::RecordBlitHistory(uint64_t size, uint64_t index) { + uint64_t queued = bytes_queued_; + bytes_queued_ += size; + bytes_written_[index & queue_bitmask_].bytes = queued; + bytes_written_[index & queue_bitmask_].index = index; + last_queued_ = index; +} + +uint64_t BlitKernel::PendingBytes() { + uint64_t read = queue_->LoadReadIndexRelaxed(); + uint64_t index = pending_search_index_.load(); + uint64_t last = last_queued_; + // If the last blit command has been run then the blit is empty. + if (read > last) return 0; + + index = Max(index, read); + while (index <= last) { + // Ensure any record we use was not wrapped. + if (index == bytes_written_[index & queue_bitmask_].index) { + uint64_t ret = bytes_queued_ - bytes_written_[index & queue_bitmask_].bytes; + + // Store max search index. + uint64_t old = pending_search_index_.load(); + while (old < index) { + if (pending_search_index_.compare_exchange_strong(old, index)) break; + } + + return ret; + } + index++; + } + debug_warning(false && "Race between PendingBytes and blit submission detected."); + // Zero is a valid return in this case since the command which was last when the search started is + // now complete. + return 0; +} + } // namespace amd } // namespace rocr diff --git a/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp b/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp index 29131aa1c0..2a4688df67 100644 --- a/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp @@ -116,6 +116,7 @@ template ::BlitSdma() : agent_(NULL), queue_start_addr_(NULL), + bytes_queued_(0), parity_(false), cached_reserve_index_(0), cached_commit_index_(0), @@ -172,6 +173,8 @@ hsa_status_t BlitSdma: MAKE_NAMED_SCOPE_GUARD(cleanupOnException, [&]() { Destroy(agent); };); std::memset(queue_start_addr_, 0, kQueueSize); + bytes_written_.resize(kQueueSize); + // Access kernel driver to initialize the queue control block // This call binds user mode queue object to underlying compute // device. ROCr creates queues that are of two kinds: PCIe optimized @@ -223,7 +226,8 @@ hsa_status_t BlitSdma: template hsa_status_t BlitSdma::SubmitBlockingCommand(const void* cmd, size_t cmd_size) { + useGCR>::SubmitBlockingCommand(const void* cmd, size_t cmd_size, + uint64_t size) { ScopedAcquire lock(&lock_); // Alternate between completion signals @@ -244,14 +248,15 @@ hsa_status_t BlitSdma(), *completionSignal); + hsa_status_t ret = + SubmitCommand(cmd, cmd_size, size, std::vector(), *completionSignal); completionSignal->WaitRelaxed(HSA_SIGNAL_CONDITION_EQ, 1, -1, HSA_WAIT_STATE_BLOCKED); return ret; } template hsa_status_t BlitSdma::SubmitCommand( - const void* cmd, size_t cmd_size, const std::vector& dep_signals, + const void* cmd, size_t cmd_size, uint64_t size, const std::vector& dep_signals, core::Signal& out_signal) { // The signal is 64 bit value, and poll checks for 32 bit value. So we // need to use two poll operations per dependent signal. @@ -307,11 +312,19 @@ hsa_status_t BlitSdma: total_timestamp_command_size + interrupt_command_size + flush_cmd_size; RingIndexTy curr_index; - char* command_addr = AcquireWriteAddress(total_command_size, curr_index); - - if (command_addr == NULL) { - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + char* command_addr; + uint64_t prior_bytes, post_bytes; + { + std::lock_guard lock(reservation_lock_); + command_addr = AcquireWriteAddress(total_command_size, curr_index); + if (command_addr == nullptr) { + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } + prior_bytes = bytes_queued_; + bytes_queued_ += size; + post_bytes = bytes_queued_; } + uint32_t wrapped_index = WrapIntoRing(curr_index); for (size_t i = 0; i < dep_signals.size(); ++i) { uint32_t* signal_addr = @@ -319,14 +332,20 @@ hsa_status_t BlitSdma: // Wait for the higher 64 bit to 0. BuildPollCommand(command_addr, &signal_addr[1], 0); command_addr += poll_command_size_; + bytes_written_[wrapped_index] = prior_bytes; + wrapped_index += poll_command_size_; // Then wait for the lower 64 bit to 0. BuildPollCommand(command_addr, &signal_addr[0], 0); command_addr += poll_command_size_; + bytes_written_[wrapped_index] = prior_bytes; + wrapped_index += poll_command_size_; } if (profiling_enabled) { BuildGetGlobalTimestampCommand(command_addr, reinterpret_cast(start_ts_addr)); command_addr += timestamp_command_size_; + bytes_written_[wrapped_index] = prior_bytes; + wrapped_index += timestamp_command_size_; } // Issue a Hdp flush cmd @@ -334,6 +353,8 @@ hsa_status_t BlitSdma: if ((HwIndexMonotonic) && (hdp_flush_support_)) { BuildHdpFlushCommand(command_addr); command_addr += flush_command_size_; + bytes_written_[wrapped_index] = prior_bytes; + wrapped_index += flush_command_size_; } } @@ -341,16 +362,22 @@ hsa_status_t BlitSdma: if (useGCR) { BuildGCRCommand(command_addr, true); command_addr += gcr_command_size_; + bytes_written_[wrapped_index] = prior_bytes; + wrapped_index += gcr_command_size_; } // Do the command after all polls are satisfied. memcpy(command_addr, cmd, cmd_size); command_addr += cmd_size; + bytes_written_.fill(wrapped_index, wrapped_index + cmd_size, prior_bytes); + wrapped_index += cmd_size; // Issue cache writeback if (useGCR) { BuildGCRCommand(command_addr, false); command_addr += gcr_command_size_; + bytes_written_[wrapped_index] = post_bytes; + wrapped_index += gcr_command_size_; } if (profiling_enabled) { @@ -358,25 +385,31 @@ hsa_status_t BlitSdma: BuildGetGlobalTimestampCommand(command_addr, reinterpret_cast(end_ts_addr)); command_addr += timestamp_command_size_; + bytes_written_[wrapped_index] = post_bytes; + wrapped_index += timestamp_command_size_; } // After transfer is completed, decrement the signal value. if (platform_atomic_support_) { BuildAtomicDecrementCommand(command_addr, out_signal.ValueLocation()); command_addr += atomic_command_size_; - + bytes_written_[wrapped_index] = post_bytes; + wrapped_index += atomic_command_size_; } else { uint32_t* signal_value_location = reinterpret_cast(out_signal.ValueLocation()); if (completion_signal_value > UINT32_MAX) { BuildFenceCommand(command_addr, signal_value_location + 1, static_cast(completion_signal_value >> 32)); command_addr += fence_command_size_; + bytes_written_[wrapped_index] = post_bytes; + wrapped_index += fence_command_size_; } BuildFenceCommand(command_addr, signal_value_location, static_cast(completion_signal_value)); - command_addr += fence_command_size_; + bytes_written_[wrapped_index] = post_bytes; + wrapped_index += fence_command_size_; } // Update mailbox event and send interrupt to IH. @@ -385,8 +418,12 @@ hsa_status_t BlitSdma: reinterpret_cast(out_signal.signal_.event_mailbox_ptr), static_cast(out_signal.signal_.event_id)); command_addr += fence_command_size_; + bytes_written_[wrapped_index] = post_bytes; + wrapped_index += fence_command_size_; BuildTrapCommand(command_addr, out_signal.signal_.event_id); + bytes_written_[wrapped_index] = post_bytes; + wrapped_index += trap_command_size_; } ReleaseWriteAddress(curr_index, total_command_size); @@ -404,7 +441,7 @@ hsa_status_t BlitSdma buff(num_copy_command); BuildCopyCommand(reinterpret_cast(&buff[0]), num_copy_command, dst, src, size); - return SubmitBlockingCommand(&buff[0], buff.size() * sizeof(SDMA_PKT_COPY_LINEAR)); + return SubmitBlockingCommand(&buff[0], buff.size() * sizeof(SDMA_PKT_COPY_LINEAR), size); } template @@ -420,7 +457,7 @@ hsa_status_t BlitSdma buff(num_copy_command); BuildCopyCommand(reinterpret_cast(&buff[0]), num_copy_command, dst, src, size); - return SubmitCommand(&buff[0], buff.size() * sizeof(SDMA_PKT_COPY_LINEAR), dep_signals, + return SubmitCommand(&buff[0], buff.size() * sizeof(SDMA_PKT_COPY_LINEAR), size, dep_signals, out_signal); } @@ -452,6 +489,7 @@ BlitSdma::SubmitCopyRe const uint max_pitch = 1 << SDMA_PKT_COPY_LINEAR_RECT::pitch_bits; std::vector pkts; + std::vector bytes_moved; auto append = [&](size_t size) { assert(size == sizeof(SDMA_PKT_COPY_LINEAR_RECT) && "SDMA packet size missmatch"); pkts.emplace_back(SDMA_PKT_COPY_LINEAR_RECT()); @@ -484,7 +522,9 @@ BlitSdma::SubmitCopyRe BuildCopyRectCommand(append, dst, dst_offset, src, src_offset, range); } - return SubmitCommand(&pkts[0], pkts.size() * sizeof(SDMA_PKT_COPY_LINEAR_RECT), dep_signals, + uint64_t size = range->x * range->y * range->z; + + return SubmitCommand(&pkts[0], pkts.size() * sizeof(SDMA_PKT_COPY_LINEAR_RECT), size, dep_signals, out_signal); } @@ -498,7 +538,7 @@ hsa_status_t BlitSdma buff(num_fill_command); BuildFillCommand(reinterpret_cast(&buff[0]), num_fill_command, ptr, value, count); - return SubmitBlockingCommand(&buff[0], buff.size() * sizeof(SDMA_PKT_CONSTANT_FILL)); + return SubmitBlockingCommand(&buff[0], buff.size() * sizeof(SDMA_PKT_CONSTANT_FILL), size); } template @@ -512,7 +552,7 @@ char* BlitSdma::Acquir uint32_t cmd_size, RingIndexTy& curr_index) { // Ring is full when all but one byte is written. if (cmd_size >= kQueueSize) { - return NULL; + return nullptr; } while (true) { @@ -545,7 +585,7 @@ char* BlitSdma::Acquir os::YieldThread(); } - return NULL; + return nullptr; } template @@ -566,7 +606,7 @@ void BlitSdma(queue_resource_.Queue_write_ptr) = (HwIndexMonotonic ? new_index : WrapIntoRing(new_index)); @@ -614,6 +654,9 @@ void BlitSdma::PadRing char* nop_address = queue_start_addr_ + WrapIntoRing(curr_index); memset(nop_address, 0, new_index - curr_index); + // Pad pending bytes tracking + bytes_written_.fill(WrapIntoRing(curr_index), WrapIntoRing(new_index), bytes_queued_); + UpdateWriteAndDoorbellRegister(curr_index, new_index); } } @@ -952,6 +995,22 @@ void BlitSdma::BuildGC addr->WORD2_UNION.GCR_CONTROL_GL2_RANGE = 0; } +template +uint64_t BlitSdma::PendingBytes() { + RingIndexTy commit = atomic::Load(&cached_commit_index_, std::memory_order_acquire); + RingIndexTy hw_read_index = *reinterpret_cast(queue_resource_.Queue_read_ptr); + RingIndexTy read; + if (HwIndexMonotonic) { + read = hw_read_index; + } else { + RingIndexTy dist_to_read_index = WrapIntoRing(commit - hw_read_index); + read = commit - dist_to_read_index; + } + + if (commit == read) return 0; + return bytes_queued_ - bytes_written_[WrapIntoRing(read)]; +} + template class BlitSdma; template class BlitSdma; template class BlitSdma;