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;