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

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;
     }