From d16c392338d2ba8da749a9cf71f6d159328d8ccb Mon Sep 17 00:00:00 2001
From: Tony Tye <Tony.Tye@amd.com>
Date: Sat, 14 Oct 2023 16:44:15 +0000
Subject: [PATCH] Make intercept queue submission obstruction free

The intercept queue submit needs to be obstruction free as it can be
invoked by the runtime async handler helper thread. The code had a busy
wait loop waiting for a free slot to be available to add the retry
barrier packet. Blocking that thread prevents it servicing other async
handlers which may need to execute in order to allow packets on the
hardware queue to be processed to free up a slot.

Change the code to always leave one free slot unless there is a retry
barrier packet already on the queue.

Change-Id: If901c865550258b790b995d58037b0f99f1968cc
---
 .../core/runtime/intercept_queue.cpp          | 35 ++++++++++++-------
 1 file changed, 22 insertions(+), 13 deletions(-)

diff --git a/runtime/hsa-runtime/core/runtime/intercept_queue.cpp b/runtime/hsa-runtime/core/runtime/intercept_queue.cpp
index 004517a126..7490ba960b 100644
--- a/runtime/hsa-runtime/core/runtime/intercept_queue.cpp
+++ b/runtime/hsa-runtime/core/runtime/intercept_queue.cpp
@@ -195,25 +195,34 @@ bool InterceptQueue::Submit(const AqlPacket* packets, uint64_t count) {
     uint64_t write = wrapped->LoadWriteIndexRelaxed();
     uint64_t read = wrapped->LoadReadIndexRelaxed();
     uint64_t free_slots = wrapped->amd_queue_.hsa_queue.size - (write - read);
+    bool pending_retry_point = IsPendingRetryPoint(read);
 
-    // If out of space defer packet insertion.
-    if (free_slots <= count) {
+    // If out of space defer packet insertion. Always make sure there is a free
+    // slot available for the retry barrier packet if there is not already one
+    // present.
+    if (free_slots < count + (pending_retry_point ? 0 : 1)) {
       // If there is not already a pending retry point add one.
-      if (!IsPendingRetryPoint(read)) {
-        // Reserve and wait for one slot.
-        write = wrapped->AddWriteIndexRelaxed(1);
-        read = write - wrapped->amd_queue_.hsa_queue.size + 1;
-        while (wrapped->LoadReadIndexRelaxed() < read) os::YieldThread();
+      if (!pending_retry_point) {
+        // Reserve one slot for the barrier packet. There will always be at
+        // least one free slot.
+        assert(free_slots >= 1 &&
+               "Packet intercept error: there is no free slot for a retry barrier packet.\n");
+        // Reserve a slot for the barrier packet.
+        uint64_t barrier = wrapped->AddWriteIndexRelaxed(1);
+        assert(barrier == write &&
+               "Packet intercept error: wrapped queue has been updated by another thread.\n");
+        ++write;
 
-        // Submit barrer which will wake async queue processing.
-        ring[write & mask].barrier_and = kBarrierPacket;
-        ring[write & mask].barrier_and.completion_signal = Signal::Convert(async_doorbell_);
-        atomic::Store(&ring[write & mask].barrier_and.header, kBarrierHeader,
+        // Submit barrier which will wake async queue processing.
+        ring[barrier & mask].barrier_and = kBarrierPacket;
+        ring[barrier & mask].barrier_and.completion_signal = Signal::Convert(async_doorbell_);
+        atomic::Store(&ring[barrier & mask].barrier_and.header, kBarrierHeader,
                       std::memory_order_release);
-        HSA::hsa_signal_store_screlease(wrapped->amd_queue_.hsa_queue.doorbell_signal, write);
+        // Update the wrapped queue's doorbell so it knows there is a new packet in the queue.
+        HSA::hsa_signal_store_screlease(wrapped->amd_queue_.hsa_queue.doorbell_signal, barrier);
 
         // Record the retry point
-        retry_index_ = write;
+        retry_index_ = barrier;
       }
       return false;
     }