From be27d78b913a25c42dbecf5acb50a69aa4eebdb5 Mon Sep 17 00:00:00 2001 From: Shweta Khatri <155576586+shwetagkhatri@users.noreply.github.com> Date: Fri, 25 Sep 2026 14:37:44 -0400 Subject: [PATCH 1/4] rocr: serve swap copies with a shader blit where the DMA engine cannot A swap exchanges the contents of two buffers. Only some DMA engines carry a swap sub-op, and until now an agent whose engine did not was left rejecting the request outright, even though the same exchange is expressible as a compute dispatch the agent can run. Add a swap blit shader and route to it by capability. The DMA engine keeps the op on the parts that already reach it; an agent whose engine does not carry the sub-op falls through to the shader rather than returning an error for an operation the agent can complete. The capability claim is deliberately narrower than the packet definition: a part may encode a swap sub-op without the engine retiring it, and an engine that accepts a swap it never completes stalls the queue rather than returning a status the caller can act on. Routing for parts already served by the native packet is left exactly as it was. This commit only adds a path for agents that previously had none, so no agent moves off an engine it was using before. The blit entry point takes a batch of address pairs because that is the shape the batch copy API and its caller use; the single-pair form forwards to it and is non-virtual so an implementation cannot support one shape and reject the other. Batches larger than one queue reservation are split across reservations, and a split batch waits for its dependent signals off-queue first: barriers at the head of the first reservation would otherwise wait on signals only the caller can satisfy, inside the call the caller must return from to satisfy them. The shader addresses memory per 16-byte chunk plus a byte tail, so the fallback has no alignment requirement of its own and does not inherit the DMA engine's. rocrtst coverage for both the single-pair and batch forms, including the multi-reservation split and its deferred-dependency path, lands in a companion commit in this series rather than here, so that this commit stays runtime-only. Co-Authored-By: Claude Opus 5 --- .../hsa-runtime/core/inc/amd_blit_kernel.h | 76 ++++ .../hsa-runtime/core/inc/amd_gpu_agent.h | 9 + .../runtime/hsa-runtime/core/inc/blit.h | 47 +++ .../core/runtime/amd_blit_kernel.cpp | 336 +++++++++++++++--- .../core/runtime/amd_blit_sdma.cpp | 16 +- .../core/runtime/amd_gpu_agent.cpp | 192 +++++++--- .../core/runtime/blit_shaders/CMakeLists.txt | 2 +- .../core/runtime/blit_shaders/blit_swapCopy.s | 330 +++++++++++++++++ 8 files changed, 914 insertions(+), 94 deletions(-) create mode 100644 projects/rocr-runtime/runtime/hsa-runtime/core/runtime/blit_shaders/blit_swapCopy.s diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h index 06cb1d0094bf..026d28a163ac 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h @@ -43,6 +43,7 @@ #ifndef HSA_RUNTIME_CORE_INC_AMD_BLIT_KERNEL_H_ #define HSA_RUNTIME_CORE_INC_AMD_BLIT_KERNEL_H_ +#include #include #include #include @@ -115,8 +116,35 @@ class BlitKernel : public core::Blit { virtual void GangLeader(bool gang_leader) override {} virtual bool GangLeader() const override { return false; } + /// @brief The compute-shader blit is the swap fallback wherever the matching + /// shader is built for the agent's ISA. Initialize() registers a kernel only + /// when its shader exists, so presence is the capability. + bool SupportsSwapCopyCommand() const override { + return kernels_.count(KernelType::SwapCopy) != 0; + } + + /// @brief Submit a batch of swap copy commands. Entries and oversized + /// entries are split across dispatches; all of them share @p out_signal. + virtual hsa_status_t SubmitSwapCopyCommand( + uint32_t num_entries, void* const* addr_a_list, void* const* addr_b_list, + const size_t* size_list, std::vector& dep_signals, + core::Signal& out_signal) override; + + /// Keep the base class' single-pair convenience overload visible; declaring + /// the N-entry override above would otherwise hide it for callers holding a + /// BlitKernel rather than a core::Blit. + using core::Blit::SubmitSwapCopyCommand; + const uint16_t kInvalidPacketHeader = HSA_PACKET_TYPE_INVALID; + /// Bytes one swap dispatch exchanges. KernelArgs::swap::swap_size is a + /// uint32_t, and the bound is rounded down to a multiple of the shader's + /// 16-byte chunk so every chunk boundary stays 16-byte aligned - the shader + /// splits its grid into a 16-byte phase and a trailing-byte phase, and a + /// boundary that was not a multiple of 16 would give the tail of one chunk + /// and the head of the next inconsistent phase assignments. + static constexpr size_t kMaxSwapBytesPerDispatch = size_t(UINT32_MAX) & ~size_t(15); + private: union KernelArgs { struct __ALIGNED__(16) { @@ -150,6 +178,12 @@ class BlitKernel : public core::Blit { uint32_t fill_value; uint32_t num_workitems; } fill; + + struct __ALIGNED__(16) { + uint64_t addr_a; // First buffer address + uint64_t addr_b; // Second buffer address + uint32_t swap_size; // Size to swap in bytes + } swap; }; // Index after which bytes will have been written. @@ -170,6 +204,47 @@ class BlitKernel : public core::Blit { void PopulateQueue(uint64_t index, uint64_t code_handle, void* args, uint32_t grid_size_x, hsa_signal_t completion_signal); + /// Emit BARRIER_AND packets to gate the following dispatch on @p dep_signals. + /// Packets are written starting at @p barrier_start_index (one packet per up + /// to five dependent signals). Shared by the linear and swap copy submit + /// paths. @p completion_signal is attached to the last barrier packet; + /// pass a zero handle when a following dispatch carries the signal instead. + void SubmitDependentSignalBarriers(uint64_t barrier_start_index, + std::vector& dep_signals, + hsa_signal_t completion_signal = {0}); + + /// Complete a zero-byte copy. Still honors @p dep_signals: @p out_signal is + /// only released once they have fired. + hsa_status_t SubmitEmptyCopyCommand(std::vector& dep_signals, + core::Signal& out_signal); + + /// Dispatch packets that still fit in one queue reservation once + /// @p num_barrier_packet barrier packets have taken their place in it. + /// AcquireWriteIndex asserts the whole reservation fits in the queue, so a + /// request needing more dispatches than this is split across reservations. + /// Never returns zero. + uint32_t MaxDispatchesPerReservation(uint32_t num_barrier_packet) const; + + /// Run @p submit on a runtime handler thread once every signal in + /// @p pending_deps has reached zero, and return immediately. + /// + /// AcquireWriteIndex blocks the calling thread until the packet processor + /// retires enough earlier packets to free queue space. That is only safe + /// while those packets can retire unaided. Dependency barriers break that: + /// they sit at the head of the first reservation waiting on signals that + /// only the caller can satisfy, so a batch large enough to need a second + /// reservation would spin forever inside the very call the caller has to + /// return from in order to satisfy them. + /// + /// Deferring keeps the two mutually exclusive. A batch that fits in one + /// reservation carries its barriers in it as before; one that does not waits + /// off-queue instead and is then submitted with nothing left to gate it, so + /// its reservations drain on GPU progress alone. @p submit must therefore + /// not carry dependencies, and must already be past every failure path - + /// there is no caller left to return a status to. + hsa_status_t DeferUntilDepsResolved(const std::vector& pending_deps, + std::function submit); + KernelArgs* ObtainAsyncKernelCopyArg(); void RecordBlitHistory(uint64_t size, uint64_t index); @@ -179,6 +254,7 @@ class BlitKernel : public core::Blit { CopyAligned, CopyMisaligned, Fill, + SwapCopy, }; struct KernelCode { diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h index 10a52105ad57..ff2f06d94389 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h @@ -255,6 +255,15 @@ class GpuAgent : public GpuAgentInt { void AssembleShader(const char* func_name, AssembleTarget assemble_target, void*& code_buf, size_t& code_buf_size) const; + // @brief Whether a precompiled shader named func_name exists for this agent's ISA. + // + // Not every shader is built for every generation. Callers that can operate + // without a given shader must consult this before calling AssembleShader, + // which requires the shader to be present. + // + // @param [in] func_name Name of the shader, as passed to AssembleShader. + bool IsShaderAvailable(const char* func_name) const; + // @brief Frees code object created by AssembleShader. // // @param [in] code_buf Code object buffer. diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/blit.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/blit.h index b71ea83c01a1..97f4c66abaef 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/blit.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/blit.h @@ -119,6 +119,53 @@ class Blit { virtual void GangLeader(bool gang_leader) = 0; virtual bool GangLeader() const { return false; }; + + /// @brief Whether this blit implements SubmitSwapCopyCommand below. Lets the + /// copy router pick between a capable engine and a per-entry fallback without + /// switching on the concrete type. + /// + /// @note Deliberately not named SwapSupported(): BlitSdmaBase already uses + /// that name for "this SDMA engine has the swap packet", which is a different + /// question - the SDMA engine reaches its swap packet through SubmitBodies, + /// not through SubmitSwapCopyCommand. + virtual bool SupportsSwapCopyCommand() const { return false; } + + /// @brief Submit a batch of swap copy commands, exchanging the contents of + /// each address pair. After execution entry i's buffer A holds the original + /// contents of its buffer B and vice versa. All entries share @p out_signal, + /// which is released once every entry has completed. + /// + /// @note N-entry rather than single-pair because that is the shape the HSA + /// batch-copy API and its HIP caller use; a swap batch must not have to be + /// unrolled into N ops by the caller just to reach a blit. + /// + /// @param num_entries Number of address pairs, >= 1. + /// @param addr_a_list Array of @p num_entries first-buffer addresses. + /// @param addr_b_list Array of @p num_entries second-buffer addresses. + /// @param size_list Array of @p num_entries swap sizes in bytes. + /// @param dep_signals Arrays of dependent signals. + /// @param out_signal Output signal. + virtual hsa_status_t SubmitSwapCopyCommand( + uint32_t num_entries, void* const* addr_a_list, void* const* addr_b_list, + const size_t* size_list, std::vector& dep_signals, + core::Signal& out_signal) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + /// @brief Single-pair form of the above. + /// + /// @note Deliberately non-virtual: the N-entry overload is the only + /// overridable entry point, so an implementation cannot accidentally support + /// one shape and silently reject the other. + hsa_status_t SubmitSwapCopyCommand(void* addr_a, void* addr_b, size_t size, + std::vector& dep_signals, + core::Signal& out_signal) { + void* const addr_a_list[1] = {addr_a}; + void* const addr_b_list[1] = {addr_b}; + const size_t size_list[1] = {size}; + return SubmitSwapCopyCommand(1, addr_a_list, addr_b_list, size_list, + dep_signals, out_signal); + } }; } // namespace core } // namespace rocr diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp index 0e3000c00130..b6eb7c48812b 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp @@ -50,6 +50,7 @@ #include "core/inc/amd_gpu_agent.h" #include "core/inc/hsa_internal.h" +#include "core/inc/runtime.h" #include "core/util/utils.h" namespace rocr { @@ -568,16 +569,27 @@ hsa_status_t BlitKernel::Initialize(const core::Agent& agent) { // Obtain the number of compute units in the underlying agent. num_cus_ = gpuAgent->properties().NumFComputeCores / 4; - // Assemble shaders to AQL code objects. - std::map kernel_names = { - {KernelType::CopyAligned, "CopyAligned"}, - {KernelType::CopyMisaligned, "CopyMisaligned"}, - {KernelType::Fill, "Fill"}}; - - for (auto kernel_name : kernel_names) { - KernelCode& kernel = kernels_[kernel_name.first]; - gpuAgent->AssembleShader(kernel_name.second, AMD::GpuAgent::AssembleTarget::AQL, kernel.code_buf_, - kernel.code_buf_size_); + // Assemble shaders to AQL code objects. Optional shaders are not built for + // every generation; when one is missing it is left unregistered and the + // matching capability query reports it unsupported, rather than failing + // initialization of the whole blit engine. + struct KernelShader { + KernelType type; + const char* name; + bool optional; + }; + static const KernelShader kernel_shaders[] = { + {KernelType::CopyAligned, "CopyAligned", false}, + {KernelType::CopyMisaligned, "CopyMisaligned", false}, + {KernelType::Fill, "Fill", false}, + {KernelType::SwapCopy, "SwapCopy", true}}; + + for (const KernelShader& kernel_shader : kernel_shaders) { + if (kernel_shader.optional && !gpuAgent->IsShaderAvailable(kernel_shader.name)) continue; + + KernelCode& kernel = kernels_[kernel_shader.type]; + gpuAgent->AssembleShader(kernel_shader.name, AMD::GpuAgent::AssembleTarget::AQL, + kernel.code_buf_, kernel.code_buf_size_); // Publish the code object the same way RegionMemory::Freeze() does. gpuAgent->PcieWcFlush(kernel.code_buf_, kernel.code_buf_size_); const_cast(gpuAgent)->InvalidateCodeCaches(kernel.code_buf_, @@ -652,36 +664,10 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand(void* dst, const void* src, return HSA_STATUS_SUCCESS; } -hsa_status_t BlitKernel::SubmitLinearCopyCommand( - void* dst, const void* src, size_t size, - std::vector& dep_signals, core::Signal& out_signal, - std::vector& gang_signals) { - - if (core::Runtime::runtime_singleton_->flag().enable_dtif_fast_copy() && dep_signals.empty()) { - LogPrint(HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS, "[ROCDTIF blit kernel] src = %p, dst = %p, size = 0x%lx", src, dst, size); - memcpy(dst, src, size); - LogPrint(HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS, "[ROCDTIF blit kernel] Fast copy success"); - - hsa_signal_t signal = {(core::Signal::Convert(&out_signal)).handle}; - if (signal.handle) { - LogPrint(HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS, "[ROCDTIF blit kernel] completion_signal = 0x%zx", signal.handle); - hsa_signal_subtract_relaxed(signal, 1); - } - return HSA_STATUS_SUCCESS; - } - - // Reserve write index for barrier(s) + dispatch packet. - 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; - { - 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; +void BlitKernel::SubmitDependentSignalBarriers(uint64_t barrier_start_index, + std::vector& dep_signals, + hsa_signal_t completion_signal) { + if (dep_signals.empty()) return; // Insert barrier packets to handle dependent signals. // Barrier bit keeps signal checking traffic from competing with a copy. @@ -696,15 +682,20 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand( reinterpret_cast( queue_->public_handle()->base_address); + uint64_t barrier_index = barrier_start_index; const size_t dep_signal_count = dep_signals.size(); for (size_t i = 0; i < dep_signal_count; ++i) { const size_t idx = i % 5; barrier_packet.dep_signal[idx] = core::Signal::Convert(dep_signals[i]); if (i == (dep_signal_count - 1) || idx == 4) { + // Only the last barrier signals completion, and only when the caller has + // no dispatch of its own to carry the signal. + if (i == (dep_signal_count - 1)) barrier_packet.completion_signal = completion_signal; + std::atomic_thread_fence(std::memory_order_acquire); - queue_buffer[(write_index)&queue_bitmask_] = barrier_packet; + queue_buffer[(barrier_index)&queue_bitmask_] = barrier_packet; std::atomic_thread_fence(std::memory_order_release); - queue_buffer[(write_index)&queue_bitmask_].header = kBarrierPacketHeader; + queue_buffer[(barrier_index)&queue_bitmask_].header = kBarrierPacketHeader; LogPrint(HSA_AMD_LOG_FLAG_AQL, "HWq=%p, id=%lu, Barrier Header = " @@ -727,14 +718,49 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand( barrier_packet.dep_signal[3].handle, barrier_packet.dep_signal[4].handle, barrier_packet.completion_signal.handle, - queue_->LoadReadIndexRelaxed(), write_index); + queue_->LoadReadIndexRelaxed(), barrier_index); - ++write_index; + ++barrier_index; memset(&barrier_packet, 0, sizeof(hsa_barrier_and_packet_t)); barrier_packet.header = HSA_PACKET_TYPE_INVALID; } } +} + +hsa_status_t BlitKernel::SubmitLinearCopyCommand( + void* dst, const void* src, size_t size, + std::vector& dep_signals, core::Signal& out_signal, + std::vector& gang_signals) { + + if (core::Runtime::runtime_singleton_->flag().enable_dtif_fast_copy() && dep_signals.empty()) { + LogPrint(HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS, "[ROCDTIF blit kernel] src = %p, dst = %p, size = 0x%lx", src, dst, size); + memcpy(dst, src, size); + LogPrint(HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS, "[ROCDTIF blit kernel] Fast copy success"); + + hsa_signal_t signal = {(core::Signal::Convert(&out_signal)).handle}; + if (signal.handle) { + LogPrint(HSA_AMD_LOG_FLAG_BLIT_KERNEL_PKTS, "[ROCDTIF blit kernel] completion_signal = 0x%zx", signal.handle); + hsa_signal_subtract_relaxed(signal, 1); + } + return HSA_STATUS_SUCCESS; + } + + // Reserve write index for barrier(s) + dispatch packet. + 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; + { + 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; + + SubmitDependentSignalBarriers(write_index, dep_signals); + write_index += num_barrier_packet; // Insert dispatch packet for copy kernel. KernelArgs* args = ObtainAsyncKernelCopyArg(); @@ -1001,5 +1027,227 @@ uint64_t BlitKernel::PendingBytes() { return 0; } +hsa_status_t BlitKernel::SubmitEmptyCopyCommand(std::vector& dep_signals, + core::Signal& out_signal) { + // Nothing to move, but the completion signal must still not fire ahead of the + // dependencies. With no dependencies it is just a local decrement; otherwise + // the barrier packets carry the completion signal in place of a dispatch. + if (dep_signals.empty()) { + out_signal.SubRelease(1); + return HSA_STATUS_SUCCESS; + } + + const uint32_t num_barrier_packet = uint32_t((dep_signals.size() + 4) / 5); + + uint64_t write_index; + { + std::lock_guard lock(reservation_lock_); + write_index = AcquireWriteIndex(num_barrier_packet); + RecordBlitHistory(0, write_index + num_barrier_packet - 1); + } + + SubmitDependentSignalBarriers(write_index, dep_signals, core::Signal::Convert(&out_signal)); + + ReleaseWriteIndex(write_index, num_barrier_packet); + + return HSA_STATUS_SUCCESS; +} + +uint32_t BlitKernel::MaxDispatchesPerReservation(uint32_t num_barrier_packet) const { + const uint32_t queue_size = queue_bitmask_ + 1; + // The barriers only ride along in the first reservation; a queue too small to + // hold them plus one dispatch would make no forward progress, but + // AcquireWriteIndex already asserts that case away, so clamp to at least one. + if (queue_size <= num_barrier_packet) return 1; + return queue_size - num_barrier_packet; +} + +hsa_status_t BlitKernel::DeferUntilDepsResolved(const std::vector& pending_deps, + std::function submit) { + assert(!pending_deps.empty() && "deferral with nothing to wait on never fires"); + + struct DeferredSubmit { + std::function submit; + std::atomic remaining; + }; + + DeferredSubmit* op = new DeferredSubmit{std::move(submit), {}}; + op->remaining.store(uint32_t(pending_deps.size()), std::memory_order_release); + + // Each dependency calls this when it reaches zero; the last one in runs the + // submission. Returning false unregisters the handler. + static hsa_amd_signal_handler handler = [](hsa_signal_value_t value, void* arg) { + DeferredSubmit* op = reinterpret_cast(arg); + if (op->remaining.fetch_sub(1, std::memory_order_acq_rel) == 1) { + op->submit(); + delete op; + } + return false; + }; + + // SetAsyncSignalHandler currently always succeeds. If that ever changes, + // failing part way would leave some handlers registered against op and the + // rest never firing, so op could be neither run nor safely deleted here. + for (size_t i = 0; i < pending_deps.size(); ++i) { + core::Runtime::runtime_singleton_->SetAsyncSignalHandler( + core::Signal::Convert(pending_deps[i]), HSA_SIGNAL_CONDITION_EQ, 0, handler, op); + } + + return HSA_STATUS_SUCCESS; +} + +hsa_status_t BlitKernel::SubmitSwapCopyCommand(uint32_t num_entries, void* const* addr_a_list, + void* const* addr_b_list, const size_t* size_list, + std::vector& dep_signals, + core::Signal& out_signal) { + if (!SupportsSwapCopyCommand()) return HSA_STATUS_ERROR_INVALID_ARGUMENT; + + if (num_entries == 0 || addr_a_list == nullptr || addr_b_list == nullptr || + size_list == nullptr) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + for (uint32_t e = 0; e < num_entries; ++e) { + if (addr_a_list[e] == nullptr || addr_b_list[e] == nullptr) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + } + + // Count the dispatches up front so out_signal can be raised in one step and + // so a batch whose entries exceed one dispatch's reach is split rather than + // rejected - the native SDMA swap path chunks oversized transfers too. + uint64_t total_dispatches = 0; + uint64_t total_bytes = 0; + for (uint32_t e = 0; e < num_entries; ++e) { + const size_t sz = size_list[e]; + if (sz == 0) continue; + total_dispatches += (sz + kMaxSwapBytesPerDispatch - 1) / kMaxSwapBytesPerDispatch; + // Swap moves 2x the size (read both, write both). + total_bytes += uint64_t(sz) * 2; + } + + if (total_dispatches == 0) { + // Every entry is zero-length: nothing to move, but the completion signal + // must still wait on the dependencies. + return SubmitEmptyCopyCommand(dep_signals, out_signal); + } + + const uint32_t num_barrier_packet = uint32_t((dep_signals.size() + 4) / 5); + + // A batch that spans reservations must not also carry dependency barriers - + // see DeferUntilDepsResolved for why the combination cannot make progress. + // Signals already at zero gate nothing, so they are dropped first; that + // alone keeps the overwhelming majority of batches on the direct path. Swap + // reaches this bound easily, emitting at least one dispatch per entry, so a + // few dozen entries can already exceed one reservation. + if (total_dispatches > MaxDispatchesPerReservation(num_barrier_packet)) { + std::vector pending_deps; + for (core::Signal* dep : dep_signals) { + if (dep->LoadRelaxed() != 0) pending_deps.push_back(dep); + } + if (!pending_deps.empty()) { + // Copied because the lists outlive this call once it defers. All three + // are validated above, so the resubmission cannot fail. + std::vector a(addr_a_list, addr_a_list + num_entries); + std::vector b(addr_b_list, addr_b_list + num_entries); + std::vector sizes(size_list, size_list + num_entries); + core::Signal* out = &out_signal; + return DeferUntilDepsResolved(pending_deps, [this, a, b, sizes, out]() { + std::vector no_deps; + SubmitSwapCopyCommand(uint32_t(a.size()), a.data(), b.data(), sizes.data(), no_deps, *out); + }); + } + } + + // Every dispatch carries out_signal and so decrements it once on retirement. + // The caller has set it to 1, so raise it by the extra dispatches before any + // packet is made visible to the packet processor. Relaxed is sufficient: the + // release fence in ReleaseWriteIndex orders this ahead of the doorbell, and + // no dispatch exists to observe it before then. + if (total_dispatches > 1) out_signal.AddRelaxed(total_dispatches - 1); + + // Walk the entries as a flat stream of chunks, filling one queue reservation + // at a time. entry / offset carry the position across reservations. + uint32_t entry = 0; + size_t offset = 0; + uint64_t remaining = total_dispatches; + bool barriers_pending = (num_barrier_packet != 0); + bool first_reservation = true; + + while (remaining != 0) { + const uint32_t barriers_now = barriers_pending ? num_barrier_packet : 0; + const uint32_t dispatches_now = uint32_t(std::min( + remaining, MaxDispatchesPerReservation(barriers_now))); + const uint32_t total_num_packet = barriers_now + dispatches_now; + + uint64_t write_index; + { + std::lock_guard lock(reservation_lock_); + write_index = AcquireWriteIndex(total_num_packet); + // Attribute the batch's bytes to the first reservation only; the later + // ones just extend the packet range those same bytes are pending behind, + // so counting them again would double-count against PendingBytes. + RecordBlitHistory(first_reservation ? total_bytes : 0, + write_index + total_num_packet - 1); + first_reservation = false; + } + + if (barriers_now != 0) { + SubmitDependentSignalBarriers(write_index, dep_signals); + barriers_pending = false; + } + + for (uint32_t d = 0; d < dispatches_now; ++d) { + // Skip entries that are empty or already fully dispatched. The bound is + // redundant with the dispatch count computed above, but keeps a miscount + // from walking off the caller's arrays. + while (entry < num_entries && offset >= size_list[entry]) { + ++entry; + offset = 0; + } + + size_t chunk = 0; + uint32_t chunk_entry = num_entries - 1; + if (entry < num_entries) { + chunk_entry = entry; + chunk = std::min(size_list[entry] - offset, kMaxSwapBytesPerDispatch); + } else { + // Unreachable: total_dispatches was derived from these same sizes. If + // it ever were reached, emit an empty dispatch rather than leaving an + // INVALID packet in a slot already reserved and doorbelled, which would + // stall the queue and hang every later copy on it. + assert(false && "swap dispatch count disagrees with the entry sizes"); + offset = 0; + } + + // Dispatch parameters. The shader runs two phases over one flat thread + // index: threads [0, chunk/16) each swap a 16-byte chunk, then threads + // [chunk/16, chunk/16 + chunk%16) each swap one of the trailing bytes. + // The two phases use disjoint thread ranges, so the grid is their sum - + // rounding chunk up to a multiple of 16 instead would leave the last + // bytes of a non-multiple-of-16 swap with no thread assigned to them. + const uint32_t bytes_per_thread = 16; + const uint32_t sz = static_cast(chunk); + const uint32_t total_workitems = (sz / bytes_per_thread) + (sz % bytes_per_thread); + + KernelArgs* args = ObtainAsyncKernelCopyArg(); + args->swap.addr_a = reinterpret_cast(addr_a_list[chunk_entry]) + offset; + args->swap.addr_b = reinterpret_cast(addr_b_list[chunk_entry]) + offset; + args->swap.swap_size = sz; + + PopulateQueue(write_index + barriers_now + d, + uintptr_t(kernels_[KernelType::SwapCopy].code_buf_), args, total_workitems, + core::Signal::Convert(&out_signal)); + + offset += chunk; + } + + ReleaseWriteIndex(write_index, total_num_packet); + remaining -= dispatches_now; + } + + return HSA_STATUS_SUCCESS; +} + } // namespace amd } // namespace rocr diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp index 2a5e617025bb..b10749443c03 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_sdma.cpp @@ -215,9 +215,23 @@ hsa_status_t BlitSdma::Initialize(const core::Agent& agent, broadcast_supported_ = true; } else if (major == 9) { broadcast_supported_ = (minor >= 4) || (minor == 0 && stepping >= 10); - swap_supported_ = (minor >= 4); } + // Swap linear copy: claimed on gfx94x+ only, the parts on which the packet + // has been exercised on hardware. + // + // The claim is deliberately narrower than the packet definition. A part may + // encode a swap sub-op without the engine retiring it, and an engine that + // accepts a swap it never completes stalls the queue instead of returning a + // status the caller can act on - so an over-broad claim costs a copy engine + // rather than an error. + // + // This flag gates the swap sub-op alone. gfx125+ reaches the packet through a + // separate condition in SubmitBodies and so does not depend on it; every other + // part that does not claim it is served by the shader blit (see + // GpuAgent::DmaCopySwap). + swap_supported_ = (major == 9 && minor >= 4); + // Allocate queue buffer. // NonPaged: queue buffer, resolved via amdgpu_vm_bo_lookup_mapping(). queue_start_addr_ = (char*)agent_->system_allocator()( diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index b94bd9fa7e82..cc5e09735231 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -329,9 +329,9 @@ GpuAgent::~GpuAgent() { regions_.clear(); } -void GpuAgent::AssembleShader(const char* func_name, AssembleTarget assemble_target, - void*& code_buf, size_t& code_buf_size) const { - // Select precompiled shader implementation from name/target. +namespace { + + // Precompiled shader implementation, selected by name and agent ISA. struct ASICShader { const void* code; size_t size; @@ -352,7 +352,7 @@ void GpuAgent::AssembleShader(const char* func_name, AssembleTarget assemble_tar ASICShader compute_1250; }; - std::map compiled_shaders = { + const std::map compiled_shaders = { {"TrapHandler", { {NULL, 0, 0, 0}, // gfx7 @@ -418,51 +418,90 @@ void GpuAgent::AssembleShader(const char* func_name, AssembleTarget assemble_tar {kCodeFill11, sizeof(kCodeFill11), 19, 8}, // gfx11 {kCodeFill12, sizeof(kCodeFill12), 19, 8}, // gfx12 {kCodeFill1250, sizeof(kCodeFill1250), 19, 8}, // gfx1250 + }}, + {"SwapCopy", + { + {NULL, 0, 0, 0}, // gfx7 (not supported) + {NULL, 0, 0, 0}, // gfx8 (not supported) + {kCodeSwapCopy9, sizeof(kCodeSwapCopy9), 19, 16}, // gfx9 + {kCodeSwapCopy9, sizeof(kCodeSwapCopy9), 19, 16}, // gfx90a + {kCodeSwapCopy9, sizeof(kCodeSwapCopy9), 19, 16}, // gfx942 + {kCodeSwapCopy1010, sizeof(kCodeSwapCopy1010), 19, 16}, // gfx1010 + {kCodeSwapCopy10, sizeof(kCodeSwapCopy10), 19, 16}, // gfx10 + {kCodeSwapCopy11, sizeof(kCodeSwapCopy11), 19, 16}, // gfx11 + {kCodeSwapCopy12, sizeof(kCodeSwapCopy12), 19, 16}, // gfx12 + {kCodeSwapCopy1250, sizeof(kCodeSwapCopy1250), 19, 16}, // gfx1250 }}}; - auto compiled_shader_it = compiled_shaders.find(func_name); - assert(compiled_shader_it != compiled_shaders.end() && - "Precompiled shader unavailable"); + // @brief Returns the shader named func_name built for isa, or NULL when this + // generation has no build of it. Single source of truth for both + // GpuAgent::AssembleShader and GpuAgent::IsShaderAvailable. + const ASICShader* SelectASICShader(const char* func_name, const core::Isa& isa) { + auto compiled_shader_it = compiled_shaders.find(func_name); + assert(compiled_shader_it != compiled_shaders.end() && + "Precompiled shader unavailable"); - ASICShader* asic_shader = NULL; + const ASICShader* asic_shader = NULL; - switch (supported_isas()[0]->GetMajorVersion()) { - case 7: - asic_shader = &compiled_shader_it->second.compute_7; - break; - case 8: - asic_shader = &compiled_shader_it->second.compute_8; - break; - case 9: - if((supported_isas()[0]->GetMinorVersion() == 0) && (supported_isas()[0]->GetStepping() == 10)) { - asic_shader = &compiled_shader_it->second.compute_90a; - } else if(supported_isas()[0]->GetMinorVersion() == 4 || supported_isas()[0]->GetMinorVersion() == 5) { - asic_shader = &compiled_shader_it->second.compute_942; - } else { - asic_shader = &compiled_shader_it->second.compute_9; - } - break; - case 10: - if(supported_isas()[0]->GetMinorVersion() == 1) - asic_shader = &compiled_shader_it->second.compute_1010; - else - asic_shader = &compiled_shader_it->second.compute_10; - break; - case 11: - asic_shader = &compiled_shader_it->second.compute_11; - break; - case 12: - if(supported_isas()[0]->GetMinorVersion() >= 5) - asic_shader = &compiled_shader_it->second.compute_1250; + switch (isa.GetMajorVersion()) { + case 7: + asic_shader = &compiled_shader_it->second.compute_7; + break; + case 8: + asic_shader = &compiled_shader_it->second.compute_8; + break; + case 9: + if((isa.GetMinorVersion() == 0) && (isa.GetStepping() == 10)) { + asic_shader = &compiled_shader_it->second.compute_90a; + } else if(isa.GetMinorVersion() == 4 || isa.GetMinorVersion() == 5) { + asic_shader = &compiled_shader_it->second.compute_942; + } else { + asic_shader = &compiled_shader_it->second.compute_9; + } + break; + case 10: + if(isa.GetMinorVersion() == 1) + asic_shader = &compiled_shader_it->second.compute_1010; else - asic_shader = &compiled_shader_it->second.compute_12; - break; - default: - assert(false && "Precompiled shader unavailable for target"); + asic_shader = &compiled_shader_it->second.compute_10; + break; + case 11: + asic_shader = &compiled_shader_it->second.compute_11; + break; + case 12: + if(isa.GetMinorVersion() >= 5) + asic_shader = &compiled_shader_it->second.compute_1250; + else + asic_shader = &compiled_shader_it->second.compute_12; + break; + default: + assert(false && "Precompiled shader unavailable for target"); + return NULL; + } + + if (asic_shader->code == NULL || asic_shader->size == 0 || + asic_shader->num_sgprs == 0 || asic_shader->num_vgprs == 0) { + return NULL; + } + + return asic_shader; } - assert((asic_shader->code && asic_shader->size && asic_shader->num_sgprs && asic_shader->num_vgprs) - && "Invalid shader"); +} // namespace + +bool GpuAgent::IsShaderAvailable(const char* func_name) const { + return SelectASICShader(func_name, *supported_isas()[0]) != NULL; +} + +void GpuAgent::AssembleShader(const char* func_name, AssembleTarget assemble_target, + void*& code_buf, size_t& code_buf_size) const { + const ASICShader* asic_shader = SelectASICShader(func_name, *supported_isas()[0]); + assert(asic_shader != NULL && "Invalid shader"); + if (asic_shader == NULL) { + code_buf = NULL; + code_buf_size = 0; + return; + } // Allocate a GPU-visible buffer for the shader. size_t header_size = @@ -1671,7 +1710,7 @@ hsa_status_t GpuAgent::DmaCopyFanOutOp( } if (op == HSA_AMD_MEMORY_COPY_OP_LINEAR_SWAP && - !coordinator->SwapSupported() && !coordinator->IsGfx125Plus()) + !coordinator->SwapSupported()) return HSA_STATUS_ERROR_INVALID_ARGUMENT; // Swap ops alignment validation @@ -2105,6 +2144,21 @@ hsa_status_t GpuAgent::DmaCopySwap( core::Signal* out_signal_obj = core::Signal::Convert(op.completion_signal); core::Signal& out_signal = *out_signal_obj; + // Route by capability: where the SDMA engine has no usable swap sub-op, the + // shader blit serves the op instead of the agent rejecting a request it can + // still satisfy. Swap support is a property of the agent's ISA, so any SDMA + // blit answers for all of them; the H2D one is the same probe DmaCopyBatch + // uses. Both the scalar and the list form route to the fallback: the list + // form is the one a batch-copy caller builds, so restricting the fallback to + // the scalar form would leave a batched swap with no route at all on an + // agent whose SDMA engine lacks the swap sub-op. + lazy_ptr& sdma_blit = GetBlitObject(BlitHostToDev); + const bool sdma_swap = sdma_blit->isSDMA() && + static_cast((*sdma_blit).get())->SwapSupported(); + + if (!sdma_swap) + return DmaCopyBatchFallback(op, dep_signals); + if (op.num_entries == 0) { const void* src_arr[1] = { op.src }; void* dst_arr[1] = { op.dst }; @@ -2187,16 +2241,58 @@ hsa_status_t GpuAgent::DmaCopyBatchFallback( out_signal.SubRelease(1); return HSA_STATUS_SUCCESS; } + case HSA_AMD_MEMORY_COPY_OP_LINEAR_SWAP: { + // Swap shader fallback: bidirectional buffer exchange using BlitKernel. + // Reached both where the SDMA engine has no usable swap sub-op and where + // SDMA is disabled outright. + // + // No alignment gate here, unlike the native path, which rejects addresses + // below the swap packet's granularity (see the SWAP alignment note in + // hsa_ext_amd.h). The shader addresses memory per 16-byte chunk plus a + // byte tail, so it is correct on any alignment. Rejecting input this path + // can serve, purely to match a restriction of a different engine, would + // fail a request for no reason; the documented alignment requirement still + // governs what a portable caller may rely on. + // + // Asymmetric swap is not supported: both regions must be the same size + // (mirrors the native DmaCopySwap path). The list form carries one size + // per entry and so is symmetric by construction. + if (op.num_entries == 0 && op.src_size != op.dst_size) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + SetCopyRequestRefCount(true); + MAKE_SCOPE_GUARD([&]() { SetCopyRequestRefCount(false); }); + + lazy_ptr& blit_dev = GetBlitObject(BlitDevToDev); + + // Route by capability rather than assuming the fallback blit implements + // swap; if it does not, reject rather than silently no-op. + if (!blit_dev->SupportsSwapCopyCommand()) + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + + if (profiling_enabled()) + out_signal.async_copy_agent(core::Agent::Convert(this->public_handle())); + + if (op.num_entries == 0) { + return blit_dev->SubmitSwapCopyCommand(const_cast(op.src), op.dst, op.src_size, + dep_signals, out_signal); + } + + // List form: src_list / dst_list hold the two sides of each exchange. + return blit_dev->SubmitSwapCopyCommand(op.num_entries, op.src_list, op.dst_list, + op.size_list, dep_signals, out_signal); + } case HSA_AMD_MEMORY_COPY_OP_LINEAR_BROADCAST: - case HSA_AMD_MEMORY_COPY_OP_LINEAR_SWAP: case HSA_AMD_MEMORY_COPY_OP_LINEAR_INDIRECT_SRC: case HSA_AMD_MEMORY_COPY_OP_LINEAR_INDIRECT_DST: case HSA_AMD_MEMORY_COPY_OP_LINEAR_INDIRECT_SRCDST: case HSA_AMD_MEMORY_COPY_OP_RECT: - // No shader-blit equivalent for broadcast/swap/indirect/rect yet; these are the - // slots for the 1-to-N / swap / indirect / rect blit shaders once added. Until - // then, reject under SDMA=0 (same as the SDMA fan-out path would), leaving - // the completion signal untouched as above. + // No shader-blit equivalent for broadcast, indirect or rect yet; broadcast + // is the slot for the 1-to-N blit shader once added, indirect copy is served + // only by the native SDMA wait/signal-indirect path, and rect has no shader + // form. Until then, reject under SDMA=0 (same as the SDMA fan-out path would), + // leaving the completion signal untouched as above. return HSA_STATUS_ERROR_INVALID_ARGUMENT; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/blit_shaders/CMakeLists.txt b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/blit_shaders/CMakeLists.txt index fbecee5441df..3dabe5932fa9 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/blit_shaders/CMakeLists.txt +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/blit_shaders/CMakeLists.txt @@ -156,7 +156,7 @@ endfunction(generate_bytecodeStrm) # Build kernels for deviceodeCopyAligned -build_kernels_for_devices("kCodeCopyAligned;kCodeCopyMisaligned;kCodeFill" "blit_copyAligned.s;blit_copyMisaligned.s;blit_fill.s") +build_kernels_for_devices("kCodeCopyAligned;kCodeCopyMisaligned;kCodeFill;kCodeSwapCopy" "blit_copyAligned.s;blit_copyMisaligned.s;blit_fill.s;blit_swapCopy.s") # Generate bytecode stream generate_bytecodeStrm("amd_blit_shaders_v2") diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/blit_shaders/blit_swapCopy.s b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/blit_shaders/blit_swapCopy.s new file mode 100644 index 000000000000..5a4d0937f97a --- /dev/null +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/blit_shaders/blit_swapCopy.s @@ -0,0 +1,330 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// +// SPDX-License-Identifier: MIT +// +//////////////////////////////////////////////////////////////////////////////// + +// SwapCopy Shader +// Exchanges contents of two buffers A and B using compute shader. +// After execution: A contains original B, B contains original A. +// All threads in workgroup participate in parallel swap. +// +// Algorithm: +// 1. Vectorized phase: Each thread swaps 16 bytes at global_id*16 +// - Bounds checked against swap_size & ~0xF +// 2. Tail phase: Threads with global_id < (swap_size % 16) swap 1 byte each +// - Handles non-16-byte-aligned sizes +// +// Kernel Arguments (20 bytes): +// [DW 0-1] Address A (64-bit) +// [DW 2-3] Address B (64-bit) +// [DW 4 ] Swap size in bytes (32-bit) +// +// Register allocation: +// s[0:1] = kernarg_segment_ptr +// s2 = workgroup_id in (tgid_x SGPR; GFX12 sources it from ttmp9), +// then workgroup_id * 64 +// s3 = saved EXEC (GFX10+) / s[2:3] saved EXEC (GFX9) +// s[4:5] = addr_a +// s[6:7] = addr_b +// s8 = swap_size +// s9 = workgroup size (64) +// s10 = stride (workgroup_size * 16) +// s11 = tail: aligned_size / 16 (number of full chunks) +// s12 = tail: remainder = swap_size & 0xF +// +// v0 = global thread offset (global_id) +// v1 = saved global_id +// v2 = vectorized: addr_a + global_id*16 / tail: byte from A +// v3 = vectorized: carry / tail: byte from B +// v4 = vectorized: addr_b + global_id*16 / tail: tail_thread_index +// v5 = vectorized: carry +// v[6:7] = vectorized: end address / tail: addr_a tail pointer +// v[8:9] = tail: addr_b tail pointer +// v[8:11] = vectorized: data from buffer A (16 bytes) +// v[12:15]= vectorized: data from buffer B (16 bytes) + +.text + +.macro V_ADD_CO_U32 vdst, src0, vsrc1 + .if (.amdgcn.gfx_generation_number >= 10) + v_add_co_u32 \vdst, vcc_lo, \src0, \vsrc1 + .elseif (.amdgcn.gfx_generation_number >= 9) + v_add_co_u32 \vdst, vcc, \src0, \vsrc1 + .else + v_add_u32 \vdst, vcc, \src0, \vsrc1 + .endif +.endm + +.macro V_ADD_CO_CI_U32 vdst, src0, vsrc1 + .if (.amdgcn.gfx_generation_number >= 10) + v_add_co_ci_u32 \vdst, vcc_lo, \src0, \vsrc1, vcc_lo + .elseif (.amdgcn.gfx_generation_number >= 9) + v_addc_co_u32 \vdst, vcc, \src0, \vsrc1, vcc + .else + v_addc_u32 \vdst, vcc, \src0, \vsrc1, vcc + .endif +.endm + +.macro V_CMP_LT_U64 src0, vsrc1 + .if (.amdgcn.gfx_generation_number >= 10) + v_cmp_lt_u64 vcc_lo, \src0, \vsrc1 + .else + v_cmp_lt_u64 vcc, \src0, \vsrc1 + .endif +.endm + +.macro V_CMP_LT_U32 src0, vsrc1 + .if (.amdgcn.gfx_generation_number >= 10) + v_cmp_lt_u32 vcc_lo, \src0, \vsrc1 + .else + v_cmp_lt_u32 vcc, \src0, \vsrc1 + .endif +.endm + +.macro S_AND_B64 sdst, src0, ssrc1 + .if (.amdgcn.gfx_generation_number >= 10) + s_and_b32 \sdst, \src0, \ssrc1 + .else + s_and_b64 \sdst, \src0, \ssrc1 + .endif +.endm + +// Non-carry add - used when we don't need carry propagation +.macro V_ADD_NC_U32 vdst, src0, vsrc1 + .if (.amdgcn.gfx_generation_number >= 10) + v_add_nc_u32 \vdst, \src0, \vsrc1 + .else + v_add_u32 \vdst, \src0, \vsrc1 + .endif +.endm + +// Non-carry subtract (reverse): vdst = vsrc1 - src0 +.macro V_SUBREV_NC_U32 vdst, src0, vsrc1 + .if (.amdgcn.gfx_generation_number >= 10) + v_sub_nc_u32 \vdst, \vsrc1, \src0 + .else + v_subrev_u32 \vdst, \src0, \vsrc1 + .endif +.endm + +// GFX12.5 (gfx1250) uses new wait instructions +.macro S_WAITCNT_KMCNT + .if (.amdgcn.gfx_generation_number == 12 && .amdgcn.gfx_generation_minor >= 5) + s_wait_kmcnt 0 + .else + s_waitcnt lgkmcnt(0) + .endif +.endm + +.macro S_WAITCNT_LOADCNT + .if (.amdgcn.gfx_generation_number == 12 && .amdgcn.gfx_generation_minor >= 5) + s_wait_loadcnt 0 + .else + s_waitcnt vmcnt(0) + .endif +.endm + +// GFX12.5 uses s_wait_storecnt instead of s_waitcnt expcnt(0) +.macro S_WAITCNT_STORECNT + .if (.amdgcn.gfx_generation_number == 12 && .amdgcn.gfx_generation_minor >= 5) + s_wait_storecnt 0 + .else + s_waitcnt expcnt(0) + .endif +.endm + +.set kSwapCopyVecWidth, 4 +.set kSwapCopyUnroll, 2 + +// The shader occupies s0-s12 (13 SGPRs) and v0-v15 (16 VGPRs). +// +// Only the .text section of this file is extracted, so the compute_pgm_rsrc1 +// assignments below never reach the GPU: the runtime builds the kernel +// descriptor itself from its own copy of these counts. Keep the two in sync. +// +// The SGPR count must leave room above the highest one used for the registers +// the hardware aliases into the top of the wave's allocation: on GFX9 that is +// FLAT_SCRATCH, XNACK_MASK and VCC, six in total, hence 13 + 6. Declaring only +// the 13 used would alias XNACK_MASK onto s[12:13] and corrupt the tail +// remainder held in s12 - on exactly the XNACK-enabled configuration this +// shader is the fallback for. +.set kSwapCopyNumSGPRs, 19 +.set kSwapCopyNumVGPRs, 16 + +.set SwapCopyRsrc1SGPRs, (kSwapCopyNumSGPRs - 1) / 8 + .if SwapCopyRsrc1SGPRs < 0 + .set SwapCopyRsrc1SGPRs, 0 + .endif + +.set SwapCopyRsrc1VGPRs, (kSwapCopyNumVGPRs - 1) / 4 + .if SwapCopyRsrc1VGPRs < 0 + .set SwapCopyRsrc1VGPRs, 0 + .endif + +.p2align 8 + +SwapCopy: + + compute_pgm_rsrc1_sgprs = SwapCopyRsrc1SGPRs + compute_pgm_rsrc1_vgprs = SwapCopyRsrc1VGPRs + compute_pgm_rsrc2_user_sgpr = 2 + compute_pgm_rsrc2_tgid_x_en = 1 + enable_sgpr_kernarg_segment_ptr = 1 + + // Load kernel arguments + // s[4:5] = addr_a, s[6:7] = addr_b, s8 = swap_size + s_load_dwordx4 s[4:7], s[0:1], 0x0 + s_load_dword s8, s[0:1], 0x10 + S_WAITCNT_KMCNT + + // Compute global thread ID = workgroup_id * 64 + local_id + // Workgroup-ID source per architecture (must match the other blit shaders): + // GFX12: ttmp9 holds workgroup_id_x + // others: s2, the architected tgid_x SGPR enabled by + // compute_pgm_rsrc2_tgid_x_en above + .if (.amdgcn.gfx_generation_number == 12) + s_lshl_b32 s2, ttmp9, 0x6 // s2 = workgroup_id * 64 + .else + s_lshl_b32 s2, s2, 0x6 // s2 = workgroup_id * 64 + .endif + V_ADD_CO_U32 v0, s2, v0 // v0 = global_id + + // Save global_id for later use + v_mov_b32 v1, v0 + + // s9 = workgroup size (64 threads) - unused now + s_mov_b32 s9, 64 + + // s10 = stride (unused in single-iteration version) + s_lshl_b32 s10, s9, 4 + + // v[2:3] = addr_a + global_id*16 + v_lshlrev_b32 v2, 4, v0 + v_mov_b32 v3, s5 + V_ADD_CO_U32 v2, v2, s4 + V_ADD_CO_CI_U32 v3, v3, 0x0 + + // v[4:5] = addr_b + global_id*16 + v_lshlrev_b32 v4, 4, v0 + v_mov_b32 v5, s7 + V_ADD_CO_U32 v4, v4, s6 + V_ADD_CO_CI_U32 v5, v5, 0x0 + + // Compute end address for bounds checking: v[6:7] = addr_a + (swap_size & ~0xF) + // Align swap_size down to 16 bytes since we operate on 16-byte chunks + v_and_b32 v6, 0xFFFFFFF0, s8 + v_mov_b32 v7, s5 + V_ADD_CO_U32 v6, v6, s4 + V_ADD_CO_CI_U32 v7, v7, 0x0 + + // Bounds check: skip vectorized swap if v[2:3] >= v[6:7] + // This handles dispatch grid rounding (grid rounded up to multiple of 64) + V_CMP_LT_U64 v[2:3], v[6:7] + s_cbranch_vccz L_SWAP_VEC_DONE + + // Mask off threads that are out of bounds for vectorized phase + .if (.amdgcn.gfx_generation_number >= 10) + s_and_saveexec_b32 s3, vcc_lo + .else + s_and_saveexec_b64 s[2:3], vcc + .endif + + // Each thread swaps its 16-byte chunk + flat_load_dwordx4 v[8:11], v[2:3] + flat_load_dwordx4 v[12:15], v[4:5] + S_WAITCNT_LOADCNT + flat_store_dwordx4 v[2:3], v[12:15] + flat_store_dwordx4 v[4:5], v[8:11] + S_WAITCNT_STORECNT + + // Restore EXEC for tail phase + .if (.amdgcn.gfx_generation_number >= 10) + s_mov_b32 exec_lo, s3 + .else + s_mov_b64 exec, s[2:3] + .endif + +L_SWAP_VEC_DONE: + // Tail Loop: Byte-by-byte swap for remaining bytes (swap_size % 16) + // The tail region is bytes [aligned_size, swap_size). + // + // The vectorized phase handles threads 0 to (aligned_size/16 - 1). + // The tail phase needs threads starting at global_id = aligned_size/16. + // + // For a tail thread: + // tail_thread_index = global_id - (aligned_size / 16) + // byte_offset = aligned_size + tail_thread_index + // + // Only threads with tail_thread_index < remainder participate. + // Where remainder = swap_size - aligned_size = swap_size & 0xF + + // v1 = saved global_id from earlier + // s11 = aligned_size / 16 (number of full 16-byte chunks) + s_lshr_b32 s11, s8, 4 // s11 = swap_size / 16 + s_and_b32 s12, s8, 0xF // s12 = remainder = swap_size & 0xF + + // Skip tail if no remainder + s_cmp_eq_u32 s12, 0 + s_cbranch_scc1 L_SWAP_DONE + + // v4 = tail_thread_index = global_id - (aligned_size / 16) + // = global_id - s11 + V_SUBREV_NC_U32 v4, s11, v1 + + // Only threads with 0 <= tail_thread_index < remainder participate + // First check tail_thread_index >= 0 (i.e., global_id >= aligned_size/16) + // Since v_subrev_u32 does unsigned subtraction, if global_id < s11, + // the result wraps to a large value > remainder, so the comparison handles it. + + // v2 = remainder (broadcast to VGPR for comparison) + v_mov_b32 v2, s12 + + // Compare: tail_thread_index < remainder + V_CMP_LT_U32 v4, v2 + s_cbranch_vccz L_SWAP_DONE + + // Mask off threads that don't participate in tail + .if (.amdgcn.gfx_generation_number >= 10) + s_and_b32 exec_lo, exec_lo, vcc_lo + .else + s_and_b64 exec, exec, vcc + .endif + + // Compute byte address: + // byte_offset = aligned_size + tail_thread_index + // = (swap_size & ~0xF) + (global_id - swap_size/16) + // Simplify: aligned_size = (swap_size/16) * 16 + // byte_offset = (swap_size/16)*16 + global_id - swap_size/16 + // = global_id*16 - swap_size/16*(16-1) ... actually simpler: + // + // byte_offset = global_id - (swap_size/16) + (swap_size & ~0xF) + // = v4 + (s8 & ~0xF) + + // v[6:7] = addr_a + aligned_size + tail_thread_index + v_and_b32 v6, 0xFFFFFFF0, s8 // aligned_size = swap_size & ~0xF + V_ADD_NC_U32 v6, v6, v4 // + tail_thread_index + v_mov_b32 v7, s5 + V_ADD_CO_U32 v6, v6, s4 // + addr_a base + V_ADD_CO_CI_U32 v7, v7, 0x0 + + // v[8:9] = addr_b + aligned_size + tail_thread_index + v_and_b32 v8, 0xFFFFFFF0, s8 // aligned_size + V_ADD_NC_U32 v8, v8, v4 // + tail_thread_index + v_mov_b32 v9, s7 + V_ADD_CO_U32 v8, v8, s6 // + addr_b base + V_ADD_CO_CI_U32 v9, v9, 0x0 + + // Byte swap: load from both, store swapped + // v2 = byte from A, v3 = byte from B (reusing registers) + flat_load_ubyte v2, v[6:7] + flat_load_ubyte v3, v[8:9] + S_WAITCNT_LOADCNT + flat_store_byte v[6:7], v3 + flat_store_byte v[8:9], v2 + S_WAITCNT_STORECNT + +L_SWAP_DONE: + s_endpgm From add3d6810fc280ac0fd5ffbd5a51710f03b05fe9 Mon Sep 17 00:00:00 2001 From: Shweta Khatri <155576586+shwetagkhatri@users.noreply.github.com> Date: Sat, 3 Oct 2026 02:16:06 -0400 Subject: [PATCH 2/4] rocr: reserve blit ring space before publishing the write index The blit kernel advanced the queue write index by the full reservation and only then waited for room. Between those two steps the write index sat more than one ring beyond the read index, which the queue's producer contract does not permit, and the slots it named aliased packets the packet processor had not consumed yet. A submission that reserved the whole ring while a whole ring was already outstanding then waited for a drain that could not happen, because its own reservation had aliased the work it was waiting on. Wait for space first and take the slots with a compare-and-swap, so the write index is never published more than one ring beyond the read index. A full ring stays a legal reservation, which is the bound the AQL queue's own producer already uses. This also narrows the window in which the packet processor can observe a slot a reservation has claimed but not yet filled. It does not close that window: the packet body is still written after the index is published, exactly as before. A reservation larger than the ring can never be satisfied at all. Add a predicate for that and have every submit path reject it with an out of resources status instead of waiting forever, checked before the completion signal is raised so a rejected submission leaves no counted but unsent work behind. Reproduced and fixed on GFX12.5, where multi entry swap copies make whole ring reservations; the same reservation sizing is used on earlier GFX families. Co-Authored-By: Claude Opus 5 --- .../hsa-runtime/core/inc/amd_blit_kernel.h | 8 ++- .../core/runtime/amd_blit_kernel.cpp | 55 ++++++++++++++++--- 2 files changed, 55 insertions(+), 8 deletions(-) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h index 026d28a163ac..df6bd3fff4a0 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h @@ -192,8 +192,14 @@ class BlitKernel : public core::Blit { uint64_t bytes; }; + /// True when the ring is large enough to hold num_packet packets at once. + /// A reservation that fails this can never be satisfied, so the submission + /// has to be rejected rather than waited on. + bool ReservationFits(uint32_t num_packet) const; + /// Reserve a slot in the queue buffer. The call will wait until the queue - /// buffer has a room. + /// buffer has a room. Must only be called for a reservation that satisfies + /// ReservationFits. uint64_t AcquireWriteIndex(uint32_t num_packet); /// Update the queue doorbell register with ::write_index. This diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp index b6eb7c48812b..1ef3aa6fc6a4 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp @@ -750,6 +750,11 @@ 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; + // Enough dependencies to need more barrier packets than the ring holds is a + // request the queue cannot express. Report it instead of waiting for room + // that will never appear. + if (!ReservationFits(total_num_packet)) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + uint64_t write_index; { std::lock_guard lock(reservation_lock_); @@ -895,16 +900,44 @@ hsa_status_t BlitKernel::EnableProfiling(bool enable) { return HSA_STATUS_SUCCESS; } -uint64_t BlitKernel::AcquireWriteIndex(uint32_t num_packet) { - assert(queue_->public_handle()->size >= num_packet); +bool BlitKernel::ReservationFits(uint32_t num_packet) const { + return uint64_t(num_packet) <= uint64_t(queue_->public_handle()->size); +} - uint64_t write_index = queue_->AddWriteIndexAcqRel(num_packet); +uint64_t BlitKernel::AcquireWriteIndex(uint32_t num_packet) { + // A reservation larger than the ring can never be satisfied. Callers check + // this with ReservationFits and fail the submission, because there is no + // value this function could return that would be correct. + assert(ReservationFits(num_packet)); + + const uint64_t queue_size = queue_->public_handle()->size; + uint64_t write_index = queue_->LoadWriteIndexRelaxed(); + + while (true) { + // Wait for room BEFORE publishing the reservation, not after. Advancing the + // write index first leaves it pointing past the end of the ring, at slots + // this call has not written yet and which still hold the previous wrap's + // packets, headers and all. Two things go wrong there. The packet processor + // may consume those stale packets, and - when the write index ends up an + // exact multiple of the ring size ahead of the read index - the two indices + // are congruent modulo the ring size, which reads as an empty queue. The + // packet processor then goes idle with a full ring of work outstanding, so + // nothing retires and the drain this thread is waiting for can never + // happen. Reserving only space that already exists keeps the write index + // within one ring of the read index at all times, which is what makes the + // two unambiguous. + while (write_index + num_packet - queue_->LoadReadIndexRelaxed() > queue_size) { + os::YieldThread(); + write_index = queue_->LoadWriteIndexRelaxed(); + } - while (write_index + num_packet - queue_->LoadReadIndexRelaxed() > queue_->public_handle()->size) { - os::YieldThread(); + // Publish the reservation only once it fits. Every caller holds + // reservation_lock_, so this succeeds first time; the retry keeps the + // function correct for any future caller that does not. + const uint64_t prev = queue_->CasWriteIndexAcqRel(write_index, write_index + num_packet); + if (prev == write_index) return write_index; + write_index = prev; } - - return write_index; } void BlitKernel::ReleaseWriteIndex(uint64_t write_index, uint32_t num_packet) { @@ -1039,6 +1072,8 @@ hsa_status_t BlitKernel::SubmitEmptyCopyCommand(std::vector& dep_ const uint32_t num_barrier_packet = uint32_t((dep_signals.size() + 4) / 5); + if (!ReservationFits(num_barrier_packet)) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + uint64_t write_index; { std::lock_guard lock(reservation_lock_); @@ -1159,6 +1194,12 @@ hsa_status_t BlitKernel::SubmitSwapCopyCommand(uint32_t num_entries, void* const } } + // No reservation the loop takes can be larger than this. Checking it once up + // front means the loop cannot fail part way and leave the completion signal + // raised for a dispatch that was never submitted. + if (!ReservationFits(num_barrier_packet + MaxDispatchesPerReservation(num_barrier_packet))) + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + // Every dispatch carries out_signal and so decrements it once on retirement. // The caller has set it to 1, so raise it by the extra dispatches before any // packet is made visible to the packet processor. Relaxed is sufficient: the From 89bc4231ca18706612a0bf78cc5f42ca16659e11 Mon Sep 17 00:00:00 2001 From: Shweta Khatri <155576586+shwetagkhatri@users.noreply.github.com> Date: Fri, 16 Oct 2026 23:46:54 -0400 Subject: [PATCH 3/4] test: cover the swap shader blit fallback and its batch form The swap shader fallback and the ring reservation that serves it landed without any test of their own: the runtime commits beneath this one change only runtime sources, so nothing in the suite exercises the path they add and a regression in it would reach a release unobserved. Add two suites. The single-pair suite states the operation's contract - after a swap each buffer holds what the other held - across sizes from one byte to a megabyte, over aligned and unaligned lengths, and through a double swap that must restore the original contents. The batch suite drives the list form, where several thousand entries exceed what one queue reservation holds, so the split across reservations and the completion-signal accounting that split requires are what is under test, together with the shader's trailing-byte phase for lengths that are not a multiple of sixteen. The dependent-signal case asserts the completion signal is unresolved rather than that it still holds the value it was created with: the fan-out routing raises the count once per engine group before submitting, so a value above one is legitimate mid-flight and only zero means done. The same case states the requirement a second way, engine-independently, by checking that no byte has moved while the dependency is still blocked. The submission wrapper screens no argument of its own. A wrapper that rejected a null pointer or a zero length before the call would answer the negative cases itself, leaving the runtime's own validation unexercised and the suite passing whether or not that validation exists. For the same reason every negative case supplies a real completion signal, since the entry point validates the signal before it reaches any per-type field and a zero handle would be refused by the signal check rather than by the argument the case names. Nothing here skips. A skipped test is reported as passing, so an environment that supplied no device or could not obtain its buffers would otherwise be indistinguishable from a clean run; a device that cannot satisfy the test is reported as a failure to test, with the allocation status and the entry it failed on. Both suites are registered in the standard, comprehensive and full tiers. An unregistered suite is built and never run, which looks the same as a suite that passes. Co-Authored-By: Claude Opus 5 --- .../rocrtst/common/copy_test_context.h | 166 +++++ .../functional/swap_copy_multi_entry.cc | 396 +++++++++++ .../suites/functional/swap_copy_tests.cc | 634 ++++++++++++++++++ .../suites/test_common/test_categories.yaml | 6 + 4 files changed, 1202 insertions(+) create mode 100644 projects/rocr-runtime/rocrtst/common/copy_test_context.h create mode 100644 projects/rocr-runtime/rocrtst/suites/functional/swap_copy_multi_entry.cc create mode 100644 projects/rocr-runtime/rocrtst/suites/functional/swap_copy_tests.cc diff --git a/projects/rocr-runtime/rocrtst/common/copy_test_context.h b/projects/rocr-runtime/rocrtst/common/copy_test_context.h new file mode 100644 index 000000000000..ec07f6358996 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/common/copy_test_context.h @@ -0,0 +1,166 @@ +/* + * Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + * + * SPDX-License-Identifier: MIT + */ + +#ifndef ROCRTST_COMMON_COPY_TEST_CONTEXT_H_ +#define ROCRTST_COMMON_COPY_TEST_CONTEXT_H_ + +// ============================================================================= +// HsaTestContext: minimal HSA context manager for copy-operation tests. +// +// Stands apart from BaseRocR because these tests exercise the async batch copy +// API directly and need nothing the harness fixture sets up beyond an agent, a +// pool, and buffers reachable from both sides. +// ============================================================================= + +#include +#include + +#include +#include +#include + +#include "common/common.h" + +// A completion signal that releases itself. The batch copy API validates +// op.completion_signal before it reaches any of the per-type fields, so a test +// that hands it a zero-handle signal gets INVALID_ARGUMENT from the signal +// check no matter what else the op says. A negative test built that way passes +// without ever reaching the argument it names. +class ScopedTestSignal { + public: + ScopedTestSignal() { + if (hsa_signal_create(1, 0, nullptr, &signal_) != HSA_STATUS_SUCCESS) { + signal_.handle = 0; + } + } + + ~ScopedTestSignal() { + if (signal_.handle != 0) hsa_signal_destroy(signal_); + } + + ScopedTestSignal(const ScopedTestSignal&) = delete; + ScopedTestSignal& operator=(const ScopedTestSignal&) = delete; + + bool valid() const { return signal_.handle != 0; } + hsa_signal_t get() const { return signal_; } + + private: + hsa_signal_t signal_{0}; +}; + +class HsaTestContext { + public: + hsa_agent_t gpu_agent; + hsa_agent_t cpu_agent; + hsa_amd_memory_pool_t gpu_pool; + hsa_amd_memory_pool_t cpu_pool; + std::vector all_gpu_agents; + + HsaTestContext() : gpu_agent{0}, cpu_agent{0}, gpu_pool{0}, cpu_pool{0} { + hsa_status_t status = hsa_init(); + if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) { + std::cerr << "Warning: hsa_init failed" << std::endl; + return; + } + initialized_ = true; + + // Agent discovery is the shared rocrtst callbacks, not a local copy of + // them: a re-implementation here would drift from whatever the harness + // considers a GPU the next time that definition changes. + hsa_iterate_agents(rocrtst::IterateGPUAgents, &all_gpu_agents); + + if (!all_gpu_agents.empty()) gpu_agent = all_gpu_agents[0]; + + hsa_iterate_agents(rocrtst::FindCPUDevice, &cpu_agent); + + if (gpu_agent.handle != 0) gpu_pool = FindCoarseGlobalPool(gpu_agent); + if (cpu_agent.handle != 0) cpu_pool = FindCoarseGlobalPool(cpu_agent); + } + + ~HsaTestContext() { + // hsa_init/hsa_shut_down are reference counted, so releasing the reference + // this context took cannot tear the runtime down under a concurrently live + // user. Leaving it unbalanced is what breaks the harness: CommonCleanUp() + // asserts that a second hsa_shut_down() returns NOT_INITIALIZED, which can + // never hold while leaked references are outstanding. + if (initialized_) hsa_shut_down(); + } + + // Each instance owns exactly one hsa_init reference; copying would release it + // more than once. + HsaTestContext(const HsaTestContext&) = delete; + HsaTestContext& operator=(const HsaTestContext&) = delete; + + bool HasGPUAgent() const { return gpu_agent.handle != 0 && gpu_pool.handle != 0; } + + bool HasCPUAgent() const { return cpu_agent.handle != 0 && cpu_pool.handle != 0; } + + // Allocate a buffer the GPU can address and the CPU can still read and write, + // so a test can fill a pattern, run the copy, and verify the result without a + // staging step of its own. Allocated from system memory for that reason, not + // from device-local memory. Returns nullptr on allocation failure. + void* AllocateGPUBuffer(size_t size) { + void* ptr = nullptr; + if (hsa_amd_memory_pool_allocate(cpu_pool, size, 0, &ptr) != HSA_STATUS_SUCCESS) { + return nullptr; + } + if (gpu_agent.handle != 0) hsa_amd_agents_allow_access(1, &gpu_agent, nullptr, ptr); + return ptr; + } + + // Allocate from the agent's own device-local pool. Use this where the point + // of the test is the device-local copy path itself, since AllocateGPUBuffer + // hands back system memory and would exercise a different route through the + // runtime. The result is not guaranteed to be CPU-addressable, so a test that + // verifies contents has to stage through a buffer from AllocateGPUBuffer. + // + // Returns nullptr on allocation failure, or if the agent exposes no + // device-local pool. + void* AllocateDeviceLocalBuffer(size_t size) { + if (gpu_pool.handle == 0) return nullptr; + + void* ptr = nullptr; + if (hsa_amd_memory_pool_allocate(gpu_pool, size, 0, &ptr) != HSA_STATUS_SUCCESS) { + return nullptr; + } + return ptr; + } + + void Free(void* ptr) { + if (ptr != nullptr) hsa_amd_memory_pool_free(ptr); + } + + private: + // First coarse-grained global pool owned by agent, or a zero handle if it has + // none. + static hsa_amd_memory_pool_t FindCoarseGlobalPool(hsa_agent_t agent) { + hsa_amd_memory_pool_t pool = {0}; + hsa_amd_agent_iterate_memory_pools( + agent, + [](hsa_amd_memory_pool_t candidate, void* data) -> hsa_status_t { + hsa_amd_segment_t segment; + hsa_amd_memory_pool_get_info(candidate, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); + if (segment != HSA_AMD_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS; + + uint32_t flags; + hsa_amd_memory_pool_get_info(candidate, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags); + if ((flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) == 0) { + return HSA_STATUS_SUCCESS; + } + + *static_cast(data) = candidate; + return HSA_STATUS_INFO_BREAK; + }, + &pool); + return pool; + } + + // True only if this context's hsa_init() succeeded and therefore owes a + // matching hsa_shut_down(). + bool initialized_ = false; +}; + +#endif // ROCRTST_COMMON_COPY_TEST_CONTEXT_H_ diff --git a/projects/rocr-runtime/rocrtst/suites/functional/swap_copy_multi_entry.cc b/projects/rocr-runtime/rocrtst/suites/functional/swap_copy_multi_entry.cc new file mode 100644 index 000000000000..6a815d79f5ec --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/functional/swap_copy_multi_entry.cc @@ -0,0 +1,396 @@ +/* + * Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + * + * SPDX-License-Identifier: MIT + */ + +// ============================================================================= +// Multi-entry Swap Copy Tests (AILIROCR-24) +// ============================================================================= +// Coverage for the N-entry form of HSA_AMD_MEMORY_COPY_OP_LINEAR_SWAP, i.e. +// op.num_entries > 0 with src_list / dst_list / dst_agent_list / size_list. +// +// swap_copy_tests.cc only ever sets op.num_entries = 0, so the entire list +// form was untested: the per-entry walk, the split of one batch across more +// than one queue reservation, the completion-signal accounting that split +// requires, and the shader tail-byte phase for sizes that are not a multiple +// of 16. Those are exercised here. +// +// Note on zero-length entries: the API layer rejects size_list[d] == 0 for a +// multi-entry swap, so the zero-size skip inside the blit implementation is +// not reachable from the public API. TC_SWAP_ME_005 pins that rejection +// rather than asserting the skip. +// ============================================================================= + +#include +#include + +#include +#include +#include +#include +#include +#include + +#include "gtest/gtest.h" +#include "common/base_rocr_utils.h" +#include "common/common.h" +#include "common/helper_funcs.h" + +#include "common/copy_test_context.h" + +namespace { + +// Deterministic, entry-distinguishing patterns. Mixing the entry index in +// means a batch that swaps the right number of bytes but pairs the wrong +// entries still fails, which a single shared pattern would not catch. +uint8_t ByteA(uint32_t entry, size_t i) { + return static_cast((i * 7 + entry * 31 + 0xAA) & 0xFF); +} + +uint8_t ByteB(uint32_t entry, size_t i) { + return static_cast((i * 13 + entry * 17 + 0x55) & 0xFF); +} + +void FillA(void* buf, uint32_t entry, size_t size) { + uint8_t* p = static_cast(buf); + for (size_t i = 0; i < size; ++i) p[i] = ByteA(entry, i); +} + +void FillB(void* buf, uint32_t entry, size_t size) { + uint8_t* p = static_cast(buf); + for (size_t i = 0; i < size; ++i) p[i] = ByteB(entry, i); +} + +// Returns the offset of the first mismatching byte, or size if the buffer +// matches. Reporting the offset makes a partial-chunk bug legible: a failure +// at a multiple of kMaxSwapBytesPerDispatch or at a 16-byte boundary points +// straight at the chunk walk or the tail phase respectively. +size_t FirstMismatchA(const void* buf, uint32_t entry, size_t size) { + const uint8_t* p = static_cast(buf); + for (size_t i = 0; i < size; ++i) { + if (p[i] != ByteA(entry, i)) return i; + } + return size; +} + +size_t FirstMismatchB(const void* buf, uint32_t entry, size_t size) { + const uint8_t* p = static_cast(buf); + for (size_t i = 0; i < size; ++i) { + if (p[i] != ByteB(entry, i)) return i; + } + return size; +} + +// Submit one multi-entry swap op and wait for it. Mirrors the single-entry +// wrapper in swap_copy_tests.cc, but fills the list members of the op unions +// instead of the scalar ones. +hsa_status_t SubmitMultiSwap(const std::vector& a_list, const std::vector& b_list, + const std::vector& size_list, + std::vector& dst_agents, hsa_agent_t src_agent, + uint32_t num_dep_signals, const hsa_signal_t* dep_signals, + hsa_signal_t out_signal) { + hsa_amd_memory_copy_op_t op; + memset(&op, 0, sizeof(op)); + + op.version = HSA_AMD_MEMORY_COPY_OP_VERSION; + op.type = HSA_AMD_MEMORY_COPY_OP_LINEAR_SWAP; + op.num_entries = static_cast(a_list.size()); + op.traffic_class = 0; + op.completion_signal = out_signal; + op.src_list = const_cast(a_list.data()); + op.src_agent = src_agent; + op.dst_list = const_cast(b_list.data()); + op.dst_agent_list = dst_agents.data(); + op.size_list = const_cast(size_list.data()); + + return hsa_amd_memory_async_batch_copy(&op, 1, num_dep_signals, dep_signals); +} + +// Owns the two buffer sets for a batch so an ASSERT_* mid-test cannot leak +// them. The tests allocate thousands of buffers; a leak per failure would +// exhaust the pool for every later test in the same process. +class SwapBatch { + public: + SwapBatch(HsaTestContext* ctx, const std::vector& sizes) : ctx_(ctx), sizes_(sizes) {} + + ~SwapBatch() { + for (void* p : a_) ctx_->Free(p); + for (void* p : b_) ctx_->Free(p); + } + + SwapBatch(const SwapBatch&) = delete; + SwapBatch& operator=(const SwapBatch&) = delete; + + // Allocates and fills every entry. Returns false if any allocation failed. + // + // A failure here must never be turned into a skip by the caller: a skip that + // stands in for "the buffers could not be obtained" reports the suite green + // while testing nothing. entries_allocated() says how far + // it got so the caller can report a real failure with a usable message. + bool Allocate() { + // Allocate directly rather than through AllocateGPUBuffer, which discards + // the hsa_status. Keeping the status is what turns "allocation failed" + // from an unactionable message into a diagnosis. + auto alloc = [this](size_t size, void** out) { + *out = nullptr; + hsa_status_t st = hsa_amd_memory_pool_allocate(ctx_->cpu_pool, size, 0, out); + if (st != HSA_STATUS_SUCCESS) { + alloc_status_ = st; + return false; + } + if (ctx_->gpu_agent.handle != 0) { + hsa_amd_agents_allow_access(1, &ctx_->gpu_agent, nullptr, *out); + } + return true; + }; + + for (size_t e = 0; e < sizes_.size(); ++e) { + void* pa = nullptr; + void* pb = nullptr; + if (!alloc(sizes_[e], &pa) || !alloc(sizes_[e], &pb)) { + if (pa != nullptr) ctx_->Free(pa); + if (pb != nullptr) ctx_->Free(pb); + failed_entry_ = e; + return false; + } + FillA(pa, static_cast(e), sizes_[e]); + FillB(pb, static_cast(e), sizes_[e]); + a_.push_back(pa); + b_.push_back(pb); + } + dst_agents_.assign(sizes_.size(), ctx_->gpu_agent); + return true; + } + + std::vector& a() { return a_; } + std::vector& b() { return b_; } + std::vector& dst_agents() { return dst_agents_; } + const std::vector& sizes() const { return sizes_; } + + // Index of the entry whose allocation failed, and how many pairs were + // obtained before it. Reported on failure so "could not allocate" is + // distinguishable from "could not allocate this many". + size_t failed_entry() const { return failed_entry_; } + size_t entries_allocated() const { return a_.size(); } + hsa_status_t alloc_status() const { return alloc_status_; } + + private: + HsaTestContext* ctx_; + std::vector sizes_; + std::vector a_; + std::vector b_; + std::vector dst_agents_; + size_t failed_entry_ = 0; + hsa_status_t alloc_status_ = HSA_STATUS_SUCCESS; +}; + +std::string StatusHex(hsa_status_t st) { + char buf[16]; + snprintf(buf, sizeof(buf), "%x", static_cast(st)); + return std::string(buf); +} + +// Explains an allocation failure in terms the reader can act on. Used with +// ASSERT_TRUE rather than GTEST_SKIP so a device that cannot supply the +// buffers is reported as a failure to test, not as a pass. +std::string AllocFailMsg(const SwapBatch& batch) { + return "buffer allocation failed at entry " + std::to_string(batch.failed_entry()) + " of " + + std::to_string(batch.sizes().size()) + " (" + + std::to_string(batch.entries_allocated()) + " pairs obtained, requested " + + std::to_string(batch.sizes()[batch.failed_entry()]) + " bytes, status=0x" + + StatusHex(batch.alloc_status()) + ") - no swap was exercised"; +} + +} // namespace + +// ============================================================================= +// Test Fixture +// ============================================================================= + +class SwapCopyMultiEntry : public ::testing::Test { + protected: + void SetUp() override { ASSERT_TRUE(ctx_.HasGPUAgent()) << "No GPU agent found"; } + + // Runs the batch and verifies that every entry was exchanged. Verification + // covers every byte of every entry, not a sample: the defects this file + // exists to catch (a dropped final chunk, a mis-stepped entry walk, an + // unassigned tail byte) all show up at the edges, which sampling misses. + void RunAndVerify(SwapBatch* batch) { + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + hsa_status_t status = + SubmitMultiSwap(batch->a(), batch->b(), batch->sizes(), batch->dst_agents(), + ctx_.gpu_agent, 0, nullptr, signal); + if (status != HSA_STATUS_SUCCESS) { + hsa_signal_destroy(signal); + FAIL() << "multi-entry swap submission failed, status=0x" << std::hex << status; + } + + hsa_signal_value_t left = hsa_signal_wait_scacquire( + signal, HSA_SIGNAL_CONDITION_LT, 1, kTimeoutNs, HSA_WAIT_STATE_BLOCKED); + hsa_signal_destroy(signal); + // A non-zero residue here is the signature of the accounting bug this + // form is most at risk of: out_signal raised by the dispatch count but + // decremented fewer times. + ASSERT_LT(left, 1) << "completion signal never reached 0 (residual " << left + << ") - dispatch count and signal decrements disagree"; + + for (size_t e = 0; e < batch->sizes().size(); ++e) { + const size_t sz = batch->sizes()[e]; + const uint32_t ent = static_cast(e); + // After the swap, A holds B's pattern and B holds A's. + ASSERT_EQ(sz, FirstMismatchB(batch->a()[e], ent, sz)) + << "entry " << e << " (size " << sz << "): buffer A did not receive B's contents"; + ASSERT_EQ(sz, FirstMismatchA(batch->b()[e], ent, sz)) + << "entry " << e << " (size " << sz << "): buffer B did not receive A's contents"; + } + } + + static constexpr uint64_t kTimeoutNs = 60ull * 1000 * 1000 * 1000; + + HsaTestContext ctx_; +}; + +// ============================================================================= +// TC_SWAP_ME_001 - several entries of different sizes in one batch +// ============================================================================= +TEST_F(SwapCopyMultiEntry, TC_SWAP_ME_001_ThreeEntries_VariedSizes) { + std::cout << "\n=== Multi-entry swap, 3 entries of differing sizes ===" << std::endl; + + // 1031 is deliberately neither a multiple of 16 nor a power of two, so the + // shader's trailing-byte phase has to cover the last 7 bytes. + SwapBatch batch(&ctx_, {4096, 1031, 65536}); + ASSERT_TRUE(batch.Allocate()) << AllocFailMsg(batch); + + RunAndVerify(&batch); +} + +// ============================================================================= +// TC_SWAP_ME_002 - sizes that are entirely or almost entirely tail bytes +// ============================================================================= +TEST_F(SwapCopyMultiEntry, TC_SWAP_ME_002_UnalignedSmallEntries) { + std::cout << "\n=== Multi-entry swap, sub-16-byte and unaligned entries ===" << std::endl; + + // 1 and 15 have no whole 16-byte chunk at all; 17 and 31 have exactly one + // plus a tail. These are the sizes where a grid computed by rounding up to + // a multiple of 16, instead of summing the two phases, drops bytes. + SwapBatch batch(&ctx_, {1, 15, 16, 17, 31, 33}); + ASSERT_TRUE(batch.Allocate()) << AllocFailMsg(batch); + + RunAndVerify(&batch); +} + +// ============================================================================= +// TC_SWAP_ME_003 - more entries than one queue reservation can hold +// ============================================================================= +TEST_F(SwapCopyMultiEntry, TC_SWAP_ME_003_ManyEntries_SpansReservations) { + std::cout << "\n=== Multi-entry swap, 4096 entries spanning queue reservations ===" + << std::endl; + + // Each entry under kMaxSwapBytesPerDispatch costs exactly one dispatch, so + // 4096 entries is 4096 dispatches. The internal blit queue is smaller than + // that on every device tested, which forces the batch to be split across + // several reservations and wraps the kernarg ring. 4096 * 4KB * 2 is 32MB, + // small enough to run anywhere. + const uint32_t kNumEntries = 4096; + SwapBatch batch(&ctx_, std::vector(kNumEntries, 4096)); + ASSERT_TRUE(batch.Allocate()) << AllocFailMsg(batch); + + RunAndVerify(&batch); +} + +// ============================================================================= +// TC_SWAP_ME_004 - dependent signals gate a batch that spans reservations +// ============================================================================= +TEST_F(SwapCopyMultiEntry, TC_SWAP_ME_004_DependentSignal_SpansReservations) { + std::cout << "\n=== Multi-entry swap gated on a dependent signal ===" << std::endl; + + // The barrier packets ride along in the first reservation only. Using a + // batch large enough to need several reservations checks that the later + // ones do not re-emit them and, more importantly, that the dependency + // still gates the whole batch and not just its first slice. + const uint32_t kNumEntries = 2048; + SwapBatch batch(&ctx_, std::vector(kNumEntries, 4096)); + ASSERT_TRUE(batch.Allocate()) << AllocFailMsg(batch); + + hsa_signal_t dep; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &dep)); + hsa_signal_t done; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &done)); + + hsa_status_t status = SubmitMultiSwap(batch.a(), batch.b(), batch.sizes(), batch.dst_agents(), + ctx_.gpu_agent, 1, &dep, done); + if (status != HSA_STATUS_SUCCESS) { + hsa_signal_destroy(dep); + hsa_signal_destroy(done); + FAIL() << "gated multi-entry swap submission failed, status=0x" << std::hex << status; + } + + // With the dependency still held, the batch must not complete. A short + // bounded wait that times out is the expected outcome. + hsa_signal_value_t early = hsa_signal_wait_scacquire( + done, HSA_SIGNAL_CONDITION_LT, 1, 200ull * 1000 * 1000, HSA_WAIT_STATE_BLOCKED); + EXPECT_GE(early, 1) << "batch completed while its dependency was still unsatisfied"; + + hsa_signal_store_screlease(dep, 0); + + hsa_signal_value_t left = hsa_signal_wait_scacquire(done, HSA_SIGNAL_CONDITION_LT, 1, + kTimeoutNs, HSA_WAIT_STATE_BLOCKED); + hsa_signal_destroy(dep); + hsa_signal_destroy(done); + ASSERT_LT(left, 1) << "batch did not complete after its dependency was released (residual " + << left << ")"; + + for (size_t e = 0; e < batch.sizes().size(); ++e) { + const size_t sz = batch.sizes()[e]; + const uint32_t ent = static_cast(e); + ASSERT_EQ(sz, FirstMismatchB(batch.a()[e], ent, sz)) << "entry " << e << ": A not swapped"; + ASSERT_EQ(sz, FirstMismatchA(batch.b()[e], ent, sz)) << "entry " << e << ": B not swapped"; + } +} + +// ============================================================================= +// TC_SWAP_ME_005 - a zero-length entry is rejected +// ============================================================================= +TEST_F(SwapCopyMultiEntry, TC_SWAP_ME_005_ZeroLengthEntry_ShouldFail) { + std::cout << "\n=== Multi-entry swap with a zero-length entry ===" << std::endl; + + SwapBatch batch(&ctx_, {4096, 4096, 4096}); + ASSERT_TRUE(batch.Allocate()) << AllocFailMsg(batch); + + std::vector sizes = batch.sizes(); + sizes[1] = 0; + + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + hsa_status_t status = SubmitMultiSwap(batch.a(), batch.b(), sizes, batch.dst_agents(), + ctx_.gpu_agent, 0, nullptr, signal); + hsa_signal_destroy(signal); + + ASSERT_NE(HSA_STATUS_SUCCESS, status) << "a zero-length entry should be rejected"; +} + +// ============================================================================= +// TC_SWAP_ME_006 - a null entry pointer is rejected +// ============================================================================= +TEST_F(SwapCopyMultiEntry, TC_SWAP_ME_006_NullEntry_ShouldFail) { + std::cout << "\n=== Multi-entry swap with a null entry pointer ===" << std::endl; + + SwapBatch batch(&ctx_, {4096, 4096, 4096}); + ASSERT_TRUE(batch.Allocate()) << AllocFailMsg(batch); + + // Copy the list so the batch still owns every real pointer and can free + // them; only the copy carries the null. + std::vector a_list = batch.a(); + a_list[2] = nullptr; + + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + hsa_status_t status = SubmitMultiSwap(a_list, batch.b(), batch.sizes(), batch.dst_agents(), + ctx_.gpu_agent, 0, nullptr, signal); + hsa_signal_destroy(signal); + + ASSERT_NE(HSA_STATUS_SUCCESS, status) << "a null entry pointer should be rejected"; +} diff --git a/projects/rocr-runtime/rocrtst/suites/functional/swap_copy_tests.cc b/projects/rocr-runtime/rocrtst/suites/functional/swap_copy_tests.cc new file mode 100644 index 000000000000..afdaa57294bf --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/functional/swap_copy_tests.cc @@ -0,0 +1,634 @@ +/* + * Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + * + * SPDX-License-Identifier: MIT + */ + +// ============================================================================= +// Swap Copy Tests (AILIROCR-24) +// ============================================================================= +// Tests for HSA_AMD_MEMORY_COPY_OP_LINEAR_SWAP operation. +// Swap copy atomically exchanges contents of two buffers: A <-> B +// +// Test Coverage: +// ============================================================================= + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "gtest/gtest.h" +#include "common/base_rocr_utils.h" +#include "common/common.h" +#include "common/helper_funcs.h" +#include "common/hsatimer.h" + +// HsaTestContext: minimal standalone HSA setup, see the header. +#include "common/copy_test_context.h" + +// ============================================================================= +// SwapCopyTestUtils: Helper functions for swap copy tests +// ============================================================================= + +class SwapCopyTestUtils { + public: + // Fill buffer with a specific pattern based on buffer identity + static void FillPatternA(void* buffer, size_t size) { + uint8_t* ptr = static_cast(buffer); + for (size_t i = 0; i < size; ++i) { + ptr[i] = static_cast((i * 7 + 0xAA) & 0xFF); + } + } + + static void FillPatternB(void* buffer, size_t size) { + uint8_t* ptr = static_cast(buffer); + for (size_t i = 0; i < size; ++i) { + ptr[i] = static_cast((i * 13 + 0x55) & 0xFF); + } + } + + // Create reference patterns for verification + static std::vector CreatePatternA(size_t size) { + std::vector pattern(size); + for (size_t i = 0; i < size; ++i) { + pattern[i] = static_cast((i * 7 + 0xAA) & 0xFF); + } + return pattern; + } + + static std::vector CreatePatternB(size_t size) { + std::vector pattern(size); + for (size_t i = 0; i < size; ++i) { + pattern[i] = static_cast((i * 13 + 0x55) & 0xFF); + } + return pattern; + } + + // Verify buffer matches expected pattern + static bool VerifyPattern(const void* buffer, const std::vector& expected, + const char* buffer_name) { + const uint8_t* ptr = static_cast(buffer); + size_t size = expected.size(); + bool match = true; + + for (size_t i = 0; i < size; ++i) { + if (ptr[i] != expected[i]) { + if (match) { + std::cout << " [" << buffer_name << "] First mismatch at offset " << i << ": expected=0x" + << std::hex << static_cast(expected[i]) << ", actual=0x" + << static_cast(ptr[i]) << std::dec << std::endl; + } + match = false; + break; + } + } + + if (match) { + std::cout << " [" << buffer_name << "] [PASS] Pattern verified (" << size << " bytes)" + << std::endl; + } + return match; + } +}; + +// ============================================================================= +// Convenience wrapper: hsa_amd_memory_swap_copy +// ============================================================================= + +// Deliberately no null or zero-size screen of its own. A wrapper that rejects +// those arguments before the call answers the negative tests itself, so the +// runtime's own validation is never reached and the suite passes whether or +// not that validation exists. Every argument is forwarded as given and the +// status comes back from hsa_amd_memory_async_batch_copy. +static inline hsa_status_t hsa_amd_memory_swap_copy(void* buf_a, hsa_agent_t agent_a, void* buf_b, + hsa_agent_t agent_b, size_t size, + uint32_t num_dep_signals, + const hsa_signal_t* dep_signals, + hsa_signal_t out_signal) { + hsa_amd_memory_copy_op_t op; + memset(&op, 0, sizeof(op)); + + op.version = HSA_AMD_MEMORY_COPY_OP_VERSION; + op.type = HSA_AMD_MEMORY_COPY_OP_LINEAR_SWAP; + op.num_entries = 0; // Single swap (not multi-swap) + op.traffic_class = 0; + op.completion_signal = out_signal; + op.src = buf_a; + op.src_agent = agent_a; + op.dst = buf_b; + op.dst_agent = agent_b; + // For SWAP: use src_size/dst_size union members (both must be non-zero) + op.src_size = size; + op.dst_size = size; // For symmetric swap, both sizes are equal + + return hsa_amd_memory_async_batch_copy(&op, 1, num_dep_signals, dep_signals); +} + +// ============================================================================= +// Test Fixture +// ============================================================================= + +class SwapCopyTest : public ::testing::Test { + protected: + void SetUp() override { + // HsaTestContext initializes HSA and finds agents/pools in constructor + ASSERT_TRUE(ctx_.HasGPUAgent()) << "No GPU agent found"; + } + + void TearDown() override { + // HsaTestContext destructor handles cleanup + } + + HsaTestContext ctx_; +}; + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_001_BasicSwapSmall) { + std::cout << "\n=== Basic Swap Copy (4KB) ===" << std::endl; + + const size_t kSize = 4096; + + // Allocate buffers + void* buf_a = ctx_.AllocateGPUBuffer(kSize); + void* buf_b = ctx_.AllocateGPUBuffer(kSize); + ASSERT_NE(nullptr, buf_a); + ASSERT_NE(nullptr, buf_b); + + // Initialize with distinct patterns + SwapCopyTestUtils::FillPatternA(buf_a, kSize); + SwapCopyTestUtils::FillPatternB(buf_b, kSize); + + // Save expected patterns (swapped) + auto expected_a = SwapCopyTestUtils::CreatePatternB(kSize); // A should have B's pattern + auto expected_b = SwapCopyTestUtils::CreatePatternA(kSize); // B should have A's pattern + + // Create completion signal + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + // Perform swap + hsa_status_t status = hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, + kSize, 0, nullptr, signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status) << "Swap copy submission failed"; + + // Wait for completion + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + // Verify swap occurred correctly + EXPECT_TRUE( + SwapCopyTestUtils::VerifyPattern(buf_a, expected_a, "Buffer A (should have B's data)")); + EXPECT_TRUE( + SwapCopyTestUtils::VerifyPattern(buf_b, expected_b, "Buffer B (should have A's data)")); + + // Cleanup + hsa_signal_destroy(signal); + ctx_.Free(buf_a); + ctx_.Free(buf_b); +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_002_VariousSizes) { + std::cout << "\n=== Various Buffer Sizes ===" << std::endl; + + std::vector sizes = {1, 16, 256, 4096, 65536, 1024 * 1024}; + + for (size_t size : sizes) { + std::cout << " Testing size: " << size << " bytes" << std::endl; + + void* buf_a = ctx_.AllocateGPUBuffer(size); + void* buf_b = ctx_.AllocateGPUBuffer(size); + ASSERT_NE(nullptr, buf_a); + ASSERT_NE(nullptr, buf_b); + + SwapCopyTestUtils::FillPatternA(buf_a, size); + SwapCopyTestUtils::FillPatternB(buf_b, size); + + auto expected_a = SwapCopyTestUtils::CreatePatternB(size); + auto expected_b = SwapCopyTestUtils::CreatePatternA(size); + + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + hsa_status_t status = hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, + size, 0, nullptr, signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_a, expected_a, "A")); + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_b, expected_b, "B")); + + hsa_signal_destroy(signal); + ctx_.Free(buf_a); + ctx_.Free(buf_b); + } +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_003_UnalignedSizes) { + std::cout << "\n=== Unaligned Buffer Sizes ===" << std::endl; + + // Test sizes that are not multiples of 16 (shader vectorization boundary) + std::vector sizes = {7, 13, 33, 127, 1000, 4097, 65537}; + + for (size_t size : sizes) { + std::cout << " Testing unaligned size: " << size << " bytes" << std::endl; + + void* buf_a = ctx_.AllocateGPUBuffer(size); + void* buf_b = ctx_.AllocateGPUBuffer(size); + ASSERT_NE(nullptr, buf_a); + ASSERT_NE(nullptr, buf_b); + + SwapCopyTestUtils::FillPatternA(buf_a, size); + SwapCopyTestUtils::FillPatternB(buf_b, size); + + auto expected_a = SwapCopyTestUtils::CreatePatternB(size); + auto expected_b = SwapCopyTestUtils::CreatePatternA(size); + + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + hsa_status_t status = hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, + size, 0, nullptr, signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_a, expected_a, "A")); + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_b, expected_b, "B")); + + hsa_signal_destroy(signal); + ctx_.Free(buf_a); + ctx_.Free(buf_b); + } +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_004_LargeBuffer) { + std::cout << "\n=== Large Buffer Swap (100MB) ===" << std::endl; + + const size_t kSize = 100 * 1024 * 1024; // 100MB + + void* buf_a = ctx_.AllocateGPUBuffer(kSize); + void* buf_b = ctx_.AllocateGPUBuffer(kSize); + ASSERT_NE(nullptr, buf_a); + ASSERT_NE(nullptr, buf_b); + + // Fill with simple patterns for large buffers + memset(buf_a, 0xAA, kSize); + memset(buf_b, 0x55, kSize); + + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + auto start = std::chrono::high_resolution_clock::now(); + + hsa_status_t status = hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, + kSize, 0, nullptr, signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + auto end = std::chrono::high_resolution_clock::now(); + auto duration = std::chrono::duration_cast(end - start); + + std::cout << " Swap completed in " << duration.count() << " ms" << std::endl; + std::cout << " Effective bandwidth: " + << (2.0 * kSize / (1024.0 * 1024.0 * 1024.0)) / (duration.count() / 1000.0) << " GB/s" + << std::endl; + + // Verify swap (spot check) + const uint8_t* ptr_a = static_cast(buf_a); + const uint8_t* ptr_b = static_cast(buf_b); + + bool swap_correct = true; + for (size_t i = 0; i < 1000 && swap_correct; ++i) { + if (ptr_a[i] != 0x55 || ptr_b[i] != 0xAA) { + swap_correct = false; + } + } + // Check end of buffers too + for (size_t i = kSize - 1000; i < kSize && swap_correct; ++i) { + if (ptr_a[i] != 0x55 || ptr_b[i] != 0xAA) { + swap_correct = false; + } + } + + EXPECT_TRUE(swap_correct) << "Large buffer swap data verification failed"; + + hsa_signal_destroy(signal); + ctx_.Free(buf_a); + ctx_.Free(buf_b); +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_005_DoubleSwapRestoresOriginal) { + std::cout << "\n=== Double Swap Restores Original ===" << std::endl; + + const size_t kSize = 8192; + + void* buf_a = ctx_.AllocateGPUBuffer(kSize); + void* buf_b = ctx_.AllocateGPUBuffer(kSize); + ASSERT_NE(nullptr, buf_a); + ASSERT_NE(nullptr, buf_b); + + // Initialize and save original patterns + SwapCopyTestUtils::FillPatternA(buf_a, kSize); + SwapCopyTestUtils::FillPatternB(buf_b, kSize); + + auto original_a = SwapCopyTestUtils::CreatePatternA(kSize); + auto original_b = SwapCopyTestUtils::CreatePatternB(kSize); + + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + // First swap + std::cout << " Performing first swap..." << std::endl; + hsa_status_t status = hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, + kSize, 0, nullptr, signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + // Second swap (should restore original) + std::cout << " Performing second swap..." << std::endl; + hsa_signal_store_relaxed(signal, 1); + status = hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, kSize, 0, nullptr, + signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + // Verify buffers are back to original + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_a, original_a, "A (should be original)")); + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_b, original_b, "B (should be original)")); + + hsa_signal_destroy(signal); + ctx_.Free(buf_a); + ctx_.Free(buf_b); +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_006_SignalDependencies) { + std::cout << "\n=== Signal Dependency Handling ===" << std::endl; + + const size_t kSize = 4096; + + void* buf_a = ctx_.AllocateGPUBuffer(kSize); + void* buf_b = ctx_.AllocateGPUBuffer(kSize); + ASSERT_NE(nullptr, buf_a); + ASSERT_NE(nullptr, buf_b); + + SwapCopyTestUtils::FillPatternA(buf_a, kSize); + SwapCopyTestUtils::FillPatternB(buf_b, kSize); + + auto expected_a = SwapCopyTestUtils::CreatePatternB(kSize); + auto expected_b = SwapCopyTestUtils::CreatePatternA(kSize); + + // Create dependency signal (initially blocked) + hsa_signal_t dep_signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &dep_signal)); + + // Create completion signal + hsa_signal_t out_signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &out_signal)); + + // Submit swap with dependency + std::cout << " Submitting swap with blocked dependency signal..." << std::endl; + hsa_status_t status = hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, + kSize, 1, &dep_signal, out_signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + + // Verify the swap has not completed yet. + // + // The test asserts the signal is unresolved, not that it still holds its + // initial value. How far above zero it sits is a routing detail, not a + // contract: the SDMA fan-out path raises the completion count by one per + // engine group before submitting, so each contributor can decrement it, and + // a caller observing the signal mid-flight legitimately sees a value above + // the one it created the signal with. Only reaching zero means "done". + hsa_signal_value_t val = hsa_signal_load_relaxed(out_signal); + EXPECT_NE(0, val) << "Swap should not have completed yet"; + + // Stronger, engine-independent statement of the same requirement: with the + // dependency still blocked no byte may have moved, so both buffers must + // still hold what they were filled with. + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_a, expected_b, + "A (still original, dep blocked)")); + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_b, expected_a, + "B (still original, dep blocked)")); + + // Release dependency + std::cout << " Releasing dependency signal..." << std::endl; + hsa_signal_store_relaxed(dep_signal, 0); + + // Wait for completion + hsa_signal_wait_relaxed(out_signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, + HSA_WAIT_STATE_BLOCKED); + + std::cout << " Swap completed after dependency released" << std::endl; + + // Verify swap + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_a, expected_a, "A")); + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf_b, expected_b, "B")); + + hsa_signal_destroy(dep_signal); + hsa_signal_destroy(out_signal); + ctx_.Free(buf_a); + ctx_.Free(buf_b); +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_007_NullPointerValidation) { + std::cout << "\n=== Null Pointer Validation ===" << std::endl; + + const size_t kSize = 4096; + void* valid_buf = ctx_.AllocateGPUBuffer(kSize); + ASSERT_NE(nullptr, valid_buf); + + // A real completion signal, not a zero handle. The batch copy entry point + // validates op.completion_signal before it reaches any per-type field, so a + // case built on a zero-handle signal is answered by the signal check and + // never exercises the buffer pointer it names. + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + // Test null buf_a + std::cout << " Testing null buf_a..." << std::endl; + hsa_status_t status = hsa_amd_memory_swap_copy(nullptr, ctx_.gpu_agent, valid_buf, ctx_.gpu_agent, + kSize, 0, nullptr, signal); + EXPECT_EQ(HSA_STATUS_ERROR_INVALID_ARGUMENT, status); + + // Test null buf_b + std::cout << " Testing null buf_b..." << std::endl; + status = hsa_amd_memory_swap_copy(valid_buf, ctx_.gpu_agent, nullptr, ctx_.gpu_agent, kSize, 0, + nullptr, signal); + EXPECT_EQ(HSA_STATUS_ERROR_INVALID_ARGUMENT, status); + + // Test both null + std::cout << " Testing both null..." << std::endl; + status = hsa_amd_memory_swap_copy(nullptr, ctx_.gpu_agent, nullptr, ctx_.gpu_agent, kSize, 0, + nullptr, signal); + EXPECT_EQ(HSA_STATUS_ERROR_INVALID_ARGUMENT, status); + + std::cout << " [PASS] Null pointer validation passed" << std::endl; + + hsa_signal_destroy(signal); + ctx_.Free(valid_buf); +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_008_ZeroSizeHandling) { + std::cout << "\n=== Zero Size Handling ===" << std::endl; + + const size_t kSize = 4096; + void* buf_a = ctx_.AllocateGPUBuffer(kSize); + void* buf_b = ctx_.AllocateGPUBuffer(kSize); + ASSERT_NE(nullptr, buf_a); + ASSERT_NE(nullptr, buf_b); + + // As in TC_SWAP_007: a valid signal, so the rejection under test is the + // size check and not the completion-signal check that precedes it. + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + std::cout << " Testing size = 0..." << std::endl; + hsa_status_t status = + hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, 0, 0, nullptr, signal); + EXPECT_EQ(HSA_STATUS_ERROR_INVALID_ARGUMENT, status); + + std::cout << " [PASS] Zero size validation passed" << std::endl; + + hsa_signal_destroy(signal); + ctx_.Free(buf_a); + ctx_.Free(buf_b); +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_009_SameBuffer) { + std::cout << "\n=== Same Buffer (src == dst) ===" << std::endl; + + const size_t kSize = 4096; + void* buf = ctx_.AllocateGPUBuffer(kSize); + ASSERT_NE(nullptr, buf); + + // Fill with pattern + SwapCopyTestUtils::FillPatternA(buf, kSize); + auto original = SwapCopyTestUtils::CreatePatternA(kSize); + + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + // Swap buffer with itself - should be a no-op or error + std::cout << " Swapping buffer with itself..." << std::endl; + hsa_status_t status = + hsa_amd_memory_swap_copy(buf, ctx_.gpu_agent, buf, ctx_.gpu_agent, kSize, 0, nullptr, signal); + + // Depending on implementation, this may succeed (no-op) or fail + if (status == HSA_STATUS_SUCCESS) { + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + // Verify data unchanged + EXPECT_TRUE(SwapCopyTestUtils::VerifyPattern(buf, original, "Buffer (should be unchanged)")); + } else { + std::cout << " Implementation rejects same-buffer swap (acceptable)" << std::endl; + } + + hsa_signal_destroy(signal); + ctx_.Free(buf); +} + +// ============================================================================= +// +// ============================================================================= +TEST_F(SwapCopyTest, TC_SWAP_010_ChainedSwaps) { + std::cout << "\n=== Chained Swaps (A<->B, B<->C, C<->A) ===" << std::endl; + + const size_t kSize = 4096; + + void* buf_a = ctx_.AllocateGPUBuffer(kSize); + void* buf_b = ctx_.AllocateGPUBuffer(kSize); + void* buf_c = ctx_.AllocateGPUBuffer(kSize); + ASSERT_NE(nullptr, buf_a); + ASSERT_NE(nullptr, buf_b); + ASSERT_NE(nullptr, buf_c); + + // Initialize: A=0xAA, B=0xBB, C=0xCC + memset(buf_a, 0xAA, kSize); + memset(buf_b, 0xBB, kSize); + memset(buf_c, 0xCC, kSize); + + hsa_signal_t signal; + ASSERT_EQ(HSA_STATUS_SUCCESS, hsa_signal_create(1, 0, nullptr, &signal)); + + // Swap A <-> B: A=0xBB, B=0xAA, C=0xCC + std::cout << " Step 1: A <-> B" << std::endl; + ASSERT_EQ(HSA_STATUS_SUCCESS, + hsa_amd_memory_swap_copy(buf_a, ctx_.gpu_agent, buf_b, ctx_.gpu_agent, kSize, 0, + nullptr, signal)); + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + // Swap B <-> C: A=0xBB, B=0xCC, C=0xAA + std::cout << " Step 2: B <-> C" << std::endl; + hsa_signal_store_relaxed(signal, 1); + ASSERT_EQ(HSA_STATUS_SUCCESS, + hsa_amd_memory_swap_copy(buf_b, ctx_.gpu_agent, buf_c, ctx_.gpu_agent, kSize, 0, + nullptr, signal)); + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + // Swap C <-> A: A=0xAA, B=0xCC, C=0xBB + std::cout << " Step 3: C <-> A" << std::endl; + hsa_signal_store_relaxed(signal, 1); + ASSERT_EQ(HSA_STATUS_SUCCESS, + hsa_amd_memory_swap_copy(buf_c, ctx_.gpu_agent, buf_a, ctx_.gpu_agent, kSize, 0, + nullptr, signal)); + hsa_signal_wait_relaxed(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + + // Final state: A=0xAA, B=0xCC, C=0xBB (cyclic rotation) + const uint8_t* ptr_a = static_cast(buf_a); + const uint8_t* ptr_b = static_cast(buf_b); + const uint8_t* ptr_c = static_cast(buf_c); + + bool correct = true; + for (size_t i = 0; i < kSize && correct; ++i) { + if (ptr_a[i] != 0xAA || ptr_b[i] != 0xCC || ptr_c[i] != 0xBB) { + correct = false; + } + } + + EXPECT_TRUE(correct) << "Chained swap result incorrect"; + if (correct) { + std::cout << " [PASS] Chained swaps verified: A=0xAA, B=0xCC, C=0xBB" << std::endl; + } + + hsa_signal_destroy(signal); + ctx_.Free(buf_a); + ctx_.Free(buf_b); + ctx_.Free(buf_c); +} + +// No main() here. These register with the rocrtst64 binary's own gtest main, +// as every other file in this directory does. diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/test_categories.yaml b/projects/rocr-runtime/rocrtst/suites/test_common/test_categories.yaml index 43c62b8d5e2d..d6017be4e49f 100644 --- a/projects/rocr-runtime/rocrtst/suites/test_common/test_categories.yaml +++ b/projects/rocr-runtime/rocrtst/suites/test_common/test_categories.yaml @@ -90,6 +90,8 @@ test_categories: - "rocrtstFunc.GpuDiscoveryDeprecatedDoorbellTest" - "DoorbellTypeValidation.*" - "BroadcastCapabilityQuery.*" + - "SwapCopyTest.*" + - "SwapCopyMultiEntry.*" - "rocrtstFunc.GpuCoreDump_DefaultPattern" - "rocrtstFunc.GpuCoreDump_CustomPattern" - "rocrtstFunc.GpuCoreDump_DisableFlag" @@ -194,6 +196,8 @@ test_categories: - "rocrtstFunc.GpuDiscoveryDeprecatedDoorbellTest" - "DoorbellTypeValidation.*" - "BroadcastCapabilityQuery.*" + - "SwapCopyTest.*" + - "SwapCopyMultiEntry.*" - "rocrtstFunc.GpuCoreDump_DefaultPattern" - "rocrtstFunc.GpuCoreDump_CustomPattern" - "rocrtstFunc.GpuCoreDump_DisableFlag" @@ -310,6 +314,8 @@ test_categories: - "rocrtstFunc.GpuDiscoveryDeprecatedDoorbellTest" - "DoorbellTypeValidation.*" - "BroadcastCapabilityQuery.*" + - "SwapCopyTest.*" + - "SwapCopyMultiEntry.*" - "rocrtstFunc.GpuCoreDump_DefaultPattern" - "rocrtstFunc.GpuCoreDump_CustomPattern" - "rocrtstFunc.GpuCoreDump_DisableFlag" From 32d0492c9e372cf3765849c0fd1c1d4b4ddf0da1 Mon Sep 17 00:00:00 2001 From: Shweta Khatri <155576586+shwetagkhatri@users.noreply.github.com> Date: Fri, 23 Oct 2026 17:49:27 -0400 Subject: [PATCH 4/4] rocr: keep gfx125+ on the native swap packet The swap capability flag is claimed by gfx94x+ only, so routing on it alone sends gfx125+ to the shader blit even though its SDMA engine reaches the swap packet through a separate condition. Exclude gfx125+ from the fallback so it keeps the engine path, and restore the matching term in the fan-out validation so the op is not refused one level up. Parts that claim neither are unaffected and still reach the shader blit. Co-Authored-By: Claude Opus 5 --- .../hsa-runtime/core/runtime/amd_gpu_agent.cpp | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index cc5e09735231..96bacbf3c274 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -1710,7 +1710,7 @@ hsa_status_t GpuAgent::DmaCopyFanOutOp( } if (op == HSA_AMD_MEMORY_COPY_OP_LINEAR_SWAP && - !coordinator->SwapSupported()) + !coordinator->SwapSupported() && !coordinator->IsGfx125Plus()) return HSA_STATUS_ERROR_INVALID_ARGUMENT; // Swap ops alignment validation @@ -2152,9 +2152,18 @@ hsa_status_t GpuAgent::DmaCopySwap( // form is the one a batch-copy caller builds, so restricting the fallback to // the scalar form would leave a batched swap with no route at all on an // agent whose SDMA engine lacks the swap sub-op. + // + // gfx125+ reaches the swap packet through its own condition rather than + // through swap_supported_, so it is native-capable here even though it does + // not claim the flag. Diverting it to the shader blit would move it off a + // working engine path, so it is excluded from the fallback. This mirrors the + // condition the fan-out validation and the SDMA submit guard already apply. lazy_ptr& sdma_blit = GetBlitObject(BlitHostToDev); - const bool sdma_swap = sdma_blit->isSDMA() && - static_cast((*sdma_blit).get())->SwapSupported(); + BlitSdmaBase* const sdma_base = sdma_blit->isSDMA() + ? static_cast((*sdma_blit).get()) + : nullptr; + const bool sdma_swap = sdma_base != nullptr && + (sdma_base->SwapSupported() || sdma_base->IsGfx125Plus()); if (!sdma_swap) return DmaCopyBatchFallback(op, dep_signals);