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: 395ad3b77b]

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;

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.");

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;