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" 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..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 @@ -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. @@ -158,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 @@ -170,6 +210,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 +260,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..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 @@ -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,54 @@ 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; + + // 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_); + 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(); @@ -869,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) { @@ -1001,5 +1060,235 @@ 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); + + if (!ReservationFits(num_barrier_packet)) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + + 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); + }); + } + } + + // 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 + // 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..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 @@ -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 = @@ -2105,6 +2144,30 @@ 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. + // + // 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); + 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); + if (op.num_entries == 0) { const void* src_arr[1] = { op.src }; void* dst_arr[1] = { op.dst }; @@ -2187,16 +2250,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