From 9cdf39a706dd70f39215ec2b825d99dc6631b854 Mon Sep 17 00:00:00 2001 From: Tony Tye Date: Sun, 15 Oct 2023 00:13:24 +0000 Subject: [PATCH] AQL packet header may need to be loaded atomically An AQL packet header field is stored using an atomic release, and needs to be read using atomic acquire if it may be written by another thread. Change-Id: I1d75587fd93f9c6216deebffc9a627b404a7e749 [ROCm/ROCR-Runtime commit: 395ad3b77b98f737729d47ee137251f219111025] --- .../runtime/hsa-runtime/core/inc/queue.h | 26 +++++++++++-------- .../core/runtime/amd_aql_queue.cpp | 8 ++++-- .../core/runtime/intercept_queue.cpp | 14 ++++++++-- 3 files changed, 33 insertions(+), 15 deletions(-) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h index 4cf175cec0..7f304d76a7 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h @@ -77,25 +77,29 @@ struct AqlPacket { hsa_agent_dispatch_packet_t agent; }; - uint8_t type() const { - return ((dispatch.header >> HSA_PACKET_HEADER_TYPE) & - ((1 << HSA_PACKET_HEADER_WIDTH_TYPE) - 1)); + // Access the type field from a packet header. The caller is responsible for + // loading the header using an atomic or ordinary load as appropriate. + static uint8_t type(uint16_t header) { + return ((header >> HSA_PACKET_HEADER_TYPE) & ((1 << HSA_PACKET_HEADER_WIDTH_TYPE) - 1)); } - bool IsValid() const { - return ((type() <= HSA_PACKET_TYPE_BARRIER_OR) && (type() != HSA_PACKET_TYPE_INVALID)); + // Determine if a packet is valid. The caller is responsible for loading the + // header using an atomic or ordinary load as appropriate. + static bool IsValid(uint16_t header) { + return ((type(header) <= HSA_PACKET_TYPE_BARRIER_OR) && + (type(header) != HSA_PACKET_TYPE_INVALID)); } std::string string() const { std::stringstream string; - uint8_t type = this->type(); + uint8_t t = type(packet.header); - const char* type_names[] = { + static const char* type_names[] = { "HSA_PACKET_TYPE_VENDOR_SPECIFIC", "HSA_PACKET_TYPE_INVALID", "HSA_PACKET_TYPE_KERNEL_DISPATCH", "HSA_PACKET_TYPE_BARRIER_AND", "HSA_PACKET_TYPE_AGENT_DISPATCH", "HSA_PACKET_TYPE_BARRIER_OR"}; - string << "type: " << type_names[type] + string << "type: " << type_names[t] << "\nbarrier: " << ((dispatch.header >> HSA_PACKET_HEADER_BARRIER) & ((1 << HSA_PACKET_HEADER_WIDTH_BARRIER) - 1)) << "\nacquire: " << ((dispatch.header >> HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) & @@ -103,7 +107,7 @@ struct AqlPacket { << "\nrelease: " << ((dispatch.header >> HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE) & ((1 << HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE) - 1)); - if (type == HSA_PACKET_TYPE_KERNEL_DISPATCH) { + if (t == HSA_PACKET_TYPE_KERNEL_DISPATCH) { string << "\nDim: " << dispatch.setup << "\nworkgroup_size: " << dispatch.workgroup_size_x << ", " << dispatch.workgroup_size_y << ", " << dispatch.workgroup_size_z @@ -116,8 +120,8 @@ struct AqlPacket { << "\nsignal: " << dispatch.completion_signal.handle; } - if ((type == HSA_PACKET_TYPE_BARRIER_AND) || - (type == HSA_PACKET_TYPE_BARRIER_OR)) { + if ((t == HSA_PACKET_TYPE_BARRIER_AND) || + (t == HSA_PACKET_TYPE_BARRIER_OR)) { for (int i = 0; i < 5; i++) string << "\ndep[" << i << "]: " << barrier_and.dep_signal[i].handle; string << "\nsignal: " << barrier_and.completion_signal.handle; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index f732017a28..01b18552f1 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -827,8 +827,12 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { core::AqlPacket& pkt = ((core::AqlPacket*)queue->amd_queue_.hsa_queue.base_address)[pkt_slot_idx]; - assert(pkt.IsValid() && "Invalid packet in dynamic scratch handler."); - assert(pkt.type() == HSA_PACKET_TYPE_KERNEL_DISPATCH && + // Load the packet header as atomic acquire as it it written by another + // thread as atomic release. This ensures the rest of the packet fields + // are visible. + uint16_t pkt_header = atomic::Load(&pkt.packet.header, std::memory_order_acquire); + assert(core::AqlPacket::IsValid(pkt_header) && "Invalid packet in dynamic scratch handler."); + assert(core::AqlPacket::type(pkt_header) == HSA_PACKET_TYPE_KERNEL_DISPATCH && "Invalid packet in dynamic scratch handler."); assert((pkt.dispatch.workgroup_size_x != 0) && (pkt.dispatch.workgroup_size_y != 0) && (pkt.dispatch.workgroup_size_z != 0) && "Invalid dispatch dimension."); diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/intercept_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/intercept_queue.cpp index 5b210bc8bf..7f82965ee2 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/intercept_queue.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/intercept_queue.cpp @@ -49,8 +49,13 @@ namespace core { namespace { +// Determine if a packet is the AMD_AQL_FORMAT_INTERCEPT_MARKER packet. Loads +// the packet header non-atomically. That is permissable if the calling thread +// has previously loaded the header atomically to determine if it is not an +// INVALID packet. Once a packet is no longer INVALID its ownership belongs to +// the packer processor. bool inline IsInterceptMarkerPacket(const AqlPacket* packet) { - return (packet->type() == HSA_PACKET_TYPE_VENDOR_SPECIFIC) && + return (AqlPacket::type(packet->packet.header) == HSA_PACKET_TYPE_VENDOR_SPECIFIC) && (packet->amd_vendor.format == AMD_AQL_FORMAT_INTERCEPT_MARKER); } @@ -348,7 +353,12 @@ void InterceptQueue::StoreRelaxed(hsa_signal_value_t value) { uint64_t i = next_packet_; while (i < end) { - if (!ring[i & mask].IsValid()) break; + // Load the packet header as atomic acquire as it may have been written by + // another thread as atomic release. This ensures the rest of the packet + // fields are visible. Once loaded and proven not to be INVALID, further + // loads by this thread can be non-atomic. + uint16_t header = atomic::Load(&ring[i & mask].packet.header, std::memory_order_acquire); + if (!AqlPacket::IsValid(header)) break; // Process callbacks. Cursor.interceptor_index = interceptors.size() - 1;