Exception support for Queue.

Remove "zombie" queue state and report queue creation failure via
exceptions.  Make Shared object a final container and support array
objects with Shared.  Add message printing to hsa_exception in
debug builds.

Change-Id: I459f38c80846018acbf45538874e95f91dd6b195

runtime/hsa-runtime/core/common/shared.h

@@ -52,7 +52,8 @@
 
 namespace core {
 /// @brief Base class encapsulating the allocator and deallocator for
-/// shared shared object.
+/// shared shared object.  As used this will allocate GPU visible host
+/// memory mapped to all GPUs.
 class BaseShared {
  public:
   static void SetAllocateAndFree(
@@ -67,10 +68,9 @@ class BaseShared {
   static std::function<void(void*)> free_;
 };
 
-/// @brief Base class for classes that encapsulates object shared between
-/// host and agents.  Alignment defaults to __alignof(T) but may be increased.
-template <typename T, size_t Align=0>
-class Shared : public BaseShared {
+/// @brief Container for object located in GPU visible host memory.
+/// Alignment defaults to __alignof(T) but may be increased.
+template <typename T, size_t Align = 0> class Shared final : private BaseShared {
  public:
   Shared() {
     assert(allocate_ != nullptr && free_ != nullptr &&
@@ -80,29 +80,105 @@ class Shared : public BaseShared {
 
     shared_object_ =
         reinterpret_cast<T*>(allocate_(sizeof(T), Max(__alignof(T), Align), 0));
+    if (shared_object_ == nullptr) throw std::bad_alloc();
 
-    if (shared_object_ != NULL) new (shared_object_) T;
+    MAKE_NAMED_SCOPE_GUARD(throwGuard, [&]() { free_(shared_object_); });
+
+    new (shared_object_) T;
+
+    throwGuard.Dismiss();
   }
 
-  virtual ~Shared() {
+  ~Shared() {
     assert(allocate_ != nullptr && free_ != nullptr &&
            "Shared object allocator is not set");
 
-    if (IsSharedObjectAllocationValid()) {
+    if (shared_object_ != nullptr) {
       shared_object_->~T();
       free_(shared_object_);
     }
   }
 
-  T* shared_object() const { return shared_object_; }
-
-  bool IsSharedObjectAllocationValid() const {
-    return (shared_object_ != NULL);
+  Shared(Shared&& rhs) {
+    this->~Shared();
+    shared_object_ = rhs.shared_object_;
+    rhs.shared_object_ = nullptr;
   }
+  Shared& operator=(Shared&& rhs) {
+    this->~Shared();
+    shared_object_ = rhs.shared_object_;
+    rhs.shared_object_ = nullptr;
+    return *this;
+  }
+
+  T* shared_object() const { return shared_object_; }
 
  private:
   T* shared_object_;
 };
 
+/// @brief Container for array located in GPU visible host memory.
+/// Alignment defaults to __alignof(T) but may be increased.
+template <typename T, size_t Align> class SharedArray final : private BaseShared {
+ public:
+  SharedArray() : shared_object_(nullptr) {}
+
+  explicit SharedArray(size_t length) : shared_object_(nullptr), len(length) {
+    assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set");
+    static_assert((__alignof(T) <= Align) || (Align == 0), "Align is less than alignof(T)");
+
+    shared_object_ =
+        reinterpret_cast<T*>(allocate_(sizeof(T) * length, Max(__alignof(T), Align), 0));
+    if (shared_object_ == nullptr) throw std::bad_alloc();
+
+    size_t i = 0;
+
+    MAKE_NAMED_SCOPE_GUARD(loopGuard, [&]() {
+      for (size_t t = 0; t < i - 1; t++) shared_object_[t].~T();
+      free_(shared_object_);
+    });
+
+    for (; i < length; i++) new (&shared_object_[i]) T;
+
+    loopGuard.Dismiss();
+  }
+
+  ~SharedArray() {
+    assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set");
+
+    if (shared_object_ != nullptr) {
+      for (size_t i = 0; i < len; i++) shared_object_[i].~T();
+      free_(shared_object_);
+    }
+  }
+
+  SharedArray(SharedArray&& rhs) {
+    this->~SharedArray();
+    shared_object_ = rhs.shared_object_;
+    rhs.shared_object_ = nullptr;
+    len = rhs.len;
+  }
+  SharedArray& operator=(SharedArray&& rhs) {
+    this->~SharedArray();
+    shared_object_ = rhs.shared_object_;
+    rhs.shared_object_ = nullptr;
+    len = rhs.len;
+    return *this;
+  }
+
+  T& operator[](size_t index) {
+    assert(index < len && "Index out of bounds.");
+    return shared_object_[index];
+  }
+  const T& operator[](size_t index) const {
+    assert(index < len && "Index out of bounds.");
+    return shared_object_[index];
+  }
+
+ private:
+  T* shared_object_;
+  size_t len;
+};
+
 }  // namespace core
 #endif  // header guard

runtime/hsa-runtime/core/inc/amd_aql_queue.h

@@ -68,9 +68,6 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig
 
   ~AqlQueue();
 
-  /// @brief Indicates if queue is valid or not
-  bool IsValid() const { return valid_; }
-
   /// @brief Queue interfaces
   hsa_status_t Inactivate() override;
 
@@ -368,11 +365,8 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig
   // Id of the Queue used in communication with thunk
   HSA_QUEUEID queue_id_;
 
-  // Indicates is queue is valid
-  bool valid_;
-
-  // Indicates if queue is inactive
-  int32_t active_;
+  // Indicates if queue is active
+  std::atomic<bool> active_;
 
   // Cached value of HsaNodeProperties.HSA_CAPABILITY.DoorbellType
   int doorbell_type_;
@@ -401,7 +395,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig
   static HsaEvent* queue_event_;
 
   // Queue count - used to ref count queue_event_
-  static volatile uint32_t queue_count_;
+  static std::atomic<uint32_t> queue_count_;
 
   // Mutex for queue_event_ manipulation
   static KernelMutex queue_lock_;
@@ -411,5 +405,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig
   // Forbid copying and moving of this object
   DISALLOW_COPY_AND_ASSIGN(AqlQueue);
 };
+
 }  // namespace amd
+
 #endif  // header guard

runtime/hsa-runtime/core/inc/host_queue.h

@@ -148,8 +148,6 @@ class HostQueue : public Queue {
     assert(false && "HostQueue::ExecutePM4 is unimplemented");
   }
 
-  bool active() const { return active_; }
-
   void* operator new(size_t size) {
     return _aligned_malloc(size, HSA_QUEUE_ALIGN_BYTES);
   }
@@ -167,7 +165,6 @@ class HostQueue : public Queue {
   static int rtti_id_;
   static const size_t kRingAlignment = 256;
   const uint32_t size_;
-  bool active_;
   void* ring_;
 
   // Host queue id counter, starting from 0x80000000 to avoid overlaping

runtime/hsa-runtime/core/inc/intercept_queue.h

@@ -63,10 +63,6 @@ class QueueWrapper : public Queue {
   std::unique_ptr<Queue> wrapped;
 
   explicit QueueWrapper(std::unique_ptr<Queue> queue) : Queue(), wrapped(std::move(queue)) {
-    if (Queue::Shared::shared_object() == NULL) {
-      return;
-    }
-
     memcpy(&amd_queue_, &wrapped->amd_queue_, sizeof(amd_queue_t));
     wrapped->set_public_handle(wrapped.get(), public_handle_);
   }
@@ -226,6 +222,9 @@ class InterceptQueue : public QueueProxy, private LocalSignal, public Signal {
   // Indicates queue active/inactive state.
   std::atomic<bool> active_;
 
+  // Proxy packet buffer
+  SharedArray<AqlPacket, 4096> buffer_;
+
   static const hsa_signal_value_t DOORBELL_MAX = 0xFFFFFFFFFFFFFFFFull;
 
   static bool HandleAsyncDoorbell(hsa_signal_value_t value, void* arg);

runtime/hsa-runtime/core/inc/queue.h

@@ -124,6 +124,14 @@ struct SharedQueue {
   Queue* core_queue;
 };
 
+class LocalQueue {
+ public:
+  SharedQueue* queue() const { return local_queue_.shared_object(); }
+
+ private:
+  Shared<SharedQueue, AMD_QUEUE_ALIGN_BYTES> local_queue_;
+};
+
 /// @brief Class Queue which encapsulate user mode queues and
 /// provides Api to access its Read, Write indices using Acquire,
 /// Release and Relaxed semantics.
@@ -132,16 +140,10 @@ Queue is intended to be an pure interface class and may be wrapped or replaced
 by tools.
 All funtions other than Convert and public_handle must be virtual.
 */
-class Queue : public Checked<0xFA3906A679F9DB49>,
-              public Shared<SharedQueue, AMD_QUEUE_ALIGN_BYTES> {
+class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue {
  public:
-  Queue() : Shared(), amd_queue_(shared_object()->amd_queue) {
-    if (!Shared::IsSharedObjectAllocationValid()) {
-      return;
-    }
-
-    shared_object()->core_queue = this;
-
+  Queue() : LocalQueue(), amd_queue_(queue()->amd_queue) {
+    queue()->core_queue = this;
     public_handle_ = Convert(this);
   }
 
@@ -153,9 +155,7 @@ class Queue : public Checked<0xFA3906A679F9DB49>,
   ///
   /// @return hsa_queue_t * Pointer to the public data type of a queue
   static __forceinline hsa_queue_t* Convert(Queue* queue) {
-    return ((queue != NULL) && (queue->IsSharedObjectAllocationValid()))
-               ? &queue->amd_queue_.hsa_queue
-               : NULL;
+    return (queue != nullptr) ? &queue->amd_queue_.hsa_queue : nullptr;
   }
 
   /// @brief Transform the public data type of a Queue's data type into an
@@ -165,14 +165,10 @@ class Queue : public Checked<0xFA3906A679F9DB49>,
   ///
   /// @return Queue * Pointer to the Queue's implementation object
   static __forceinline Queue* Convert(const hsa_queue_t* queue) {
-    return (queue != NULL)
-               ? reinterpret_cast<const SharedQueue*>(
-                     reinterpret_cast<uintptr_t>(queue) -
-                     (reinterpret_cast<uintptr_t>(
-                          &reinterpret_cast<SharedQueue*>(1234)
-                               ->amd_queue.hsa_queue) -
-                      uintptr_t(1234)))->core_queue
-               : NULL;
+    return (queue != nullptr)
+        ? reinterpret_cast<SharedQueue*>(reinterpret_cast<uintptr_t>(queue) -
+                                         offsetof(SharedQueue, amd_queue.hsa_queue))->core_queue
+        : nullptr;
   }
 
   /// @brief Inactivate the queue object. Once inactivate a

runtime/hsa-runtime/core/inc/signal.h

@@ -101,7 +101,6 @@ static_assert(std::is_trivially_destructible<SharedSignal>::value,
 class LocalSignal {
  public:
   explicit LocalSignal(hsa_signal_value_t initial_value) {
-    if (!local_signal_.IsSharedObjectAllocationValid()) throw std::bad_alloc();
     local_signal_.shared_object()->amd_signal.value = initial_value;
   }
 

runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp

@@ -73,30 +73,26 @@ namespace amd {
 const uint32_t kAmdQueueAlignBytes = 0x40;
 
 HsaEvent* AqlQueue::queue_event_ = NULL;
-volatile uint32_t AqlQueue::queue_count_ = 0;
+std::atomic<uint32_t> AqlQueue::queue_count_(0);
 KernelMutex AqlQueue::queue_lock_;
-int AqlQueue::rtti_id_;
+int AqlQueue::rtti_id_ = 0;
 
 AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, ScratchInfo& scratch,
                    core::HsaEventCallback callback, void* err_data, bool is_kv)
     : Queue(),
       LocalSignal(0),
       Signal(signal()),
-      ring_buf_(NULL),
+      ring_buf_(nullptr),
       ring_buf_alloc_bytes_(0),
       queue_id_(HSA_QUEUEID(-1)),
-      valid_(false),
+      active_(false),
       agent_(agent),
       queue_scratch_(scratch),
       errors_callback_(callback),
       errors_data_(err_data),
       is_kv_queue_(is_kv),
-      pm4_ib_buf_(NULL),
+      pm4_ib_buf_(nullptr),
       pm4_ib_size_b_(0x1000) {
-  if (!Queue::Shared::IsSharedObjectAllocationValid()) {
-    return;
-  }
-
   // When queue_full_workaround_ is set to 1, the ring buffer is internally
   // doubled in size. Virtual addresses in the upper half of the ring allocation
   // are mapped to the same set of pages backing the lower half.
@@ -122,14 +118,16 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
   queue_size_pkts = Max(queue_size_pkts, min_pkts);
 
   uint32_t queue_size_bytes = queue_size_pkts * sizeof(core::AqlPacket);
-  if ((queue_size_bytes & (queue_size_bytes - 1)) != 0) return;
+  if ((queue_size_bytes & (queue_size_bytes - 1)) != 0)
+    throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_QUEUE_CREATION,
+                             "Requested queue with non-power of two packet capacity.\n");
 
   // Allocate the AQL packet ring buffer.
   AllocRegisteredRingBuffer(queue_size_pkts);
-  if (ring_buf_ == NULL) return;
+  if (ring_buf_ == nullptr) throw std::bad_alloc();
   MAKE_NAMED_SCOPE_GUARD(RingGuard, [&]() { FreeRegisteredRingBuffer(); });
 
-  // Fill the ring buffer with ALWAYS_RESERVED packet headers.
+  // Fill the ring buffer with invalid packet headers.
   // Leave packet content uninitialized to help track errors.
   for (uint32_t pkt_id = 0; pkt_id < queue_size_pkts; ++pkt_id) {
     ((uint32_t*)ring_buf_)[16 * pkt_id] = HSA_PACKET_TYPE_INVALID;
@@ -154,7 +152,9 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
   kmt_status = hsaKmtCreateQueue(node_id, HSA_QUEUE_COMPUTE_AQL, 100,
                                  HSA_QUEUE_PRIORITY_NORMAL, ring_buf_,
                                  ring_buf_alloc_bytes_, NULL, &queue_rsrc);
-  if (kmt_status != HSAKMT_STATUS_SUCCESS) return;
+  if (kmt_status != HSAKMT_STATUS_SUCCESS)
+    throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
+                             "Queue create failed at hsaKmtCreateQueue\n");
   queue_id_ = queue_rsrc.QueueId;
   MAKE_NAMED_SCOPE_GUARD(QueueGuard, [&]() { hsaKmtDestroyQueue(queue_id_); });
 
@@ -228,7 +228,7 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
     queue_count_--;
     if (queue_count_ == 0) {
       core::InterruptSignal::DestroyEvent(queue_event_);
-      queue_event_ = NULL;
+      queue_event_ = nullptr;
     }
   });
 
@@ -239,35 +239,38 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
   if (core::g_use_interrupt_wait) {
     ScopedAcquire<KernelMutex> _lock(&queue_lock_);
     queue_count_++;
-    if (queue_event_ == NULL) {
+    if (queue_event_ == nullptr) {
       assert(queue_count_ == 1 && "Inconsistency in queue event reference counting found.\n");
 
       queue_event_ = core::InterruptSignal::CreateEvent(HSA_EVENTTYPE_SIGNAL, false);
-      if (queue_event_ == NULL) return;
+      if (queue_event_ == nullptr)
+        throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
+                                 "Queue event creation failed.\n");
     }
     auto Signal = new core::InterruptSignal(0, queue_event_);
-    if (Signal == nullptr) return;
+    assert(Signal != nullptr && "Should have thrown!\n");
     amd_queue_.queue_inactive_signal = core::InterruptSignal::Convert(Signal);
   } else {
     EventGuard.Dismiss();
     auto Signal = new core::DefaultSignal(0);
-    if (Signal == nullptr) return;
+    assert(Signal != nullptr && "Should have thrown!\n");
     amd_queue_.queue_inactive_signal = core::DefaultSignal::Convert(Signal);
   }
   if (AMD::hsa_amd_signal_async_handler(amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE,
                                         0, DynamicScratchHandler, this) != HSA_STATUS_SUCCESS)
-    return;
+    throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
+                             "Queue event handler failed registration.\n");
 
   pm4_ib_buf_ = core::Runtime::runtime_singleton_->system_allocator()(
       pm4_ib_size_b_, 0x1000, core::MemoryRegion::AllocateExecutable);
-  if (pm4_ib_buf_ == NULL) return;
+  if (pm4_ib_buf_ == nullptr)
+    throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "PM4 IB allocation failed.\n");
 
   MAKE_NAMED_SCOPE_GUARD(PM4IBGuard, [&]() {
     core::Runtime::runtime_singleton_->system_deallocator()(pm4_ib_buf_);
   });
 
-  valid_ = true;
-  active_ = 1;
+  active_ = true;
 
   PM4IBGuard.Dismiss();
   RingGuard.Dismiss();
@@ -277,11 +280,7 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
 }
 
 AqlQueue::~AqlQueue() {
-  if (!IsValid()) {
-    return;
-  }
-
-  if (active_ == 1) hsaKmtDestroyQueue(queue_id_);
+  Inactivate();
 
   FreeRegisteredRingBuffer();
   agent_->ReleaseQueueScratch(queue_scratch_.queue_base);
@@ -654,8 +653,11 @@ void AqlQueue::FreeRegisteredRingBuffer() {
 }
 
 hsa_status_t AqlQueue::Inactivate() {
-  int32_t active = atomic::Exchange((volatile int32_t*)&active_, 0);
-  if (active == 1) hsaKmtDestroyQueue(this->queue_id_);
+  bool active = active_.exchange(false, std::memory_order_relaxed);
+  if (active) {
+    auto err = hsaKmtDestroyQueue(this->queue_id_);
+    assert(err == HSAKMT_STATUS_SUCCESS && "hsaKmtDestroyQueue failed.");
+  }
   return HSA_STATUS_SUCCESS;
 }
 

runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp

@@ -910,26 +910,21 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type,
 
   const uint32_t num_cu = properties_.NumFComputeCores / properties_.NumSIMDPerCU;
   scratch.size = scratch.size_per_thread * 32 * 64 * num_cu;
-  scratch.queue_base = NULL;
+  scratch.queue_base = nullptr;
+
+  MAKE_NAMED_SCOPE_GUARD(scratchGuard, [&]() { ReleaseQueueScratch(scratch.queue_base); });
+
   if (scratch.size != 0) {
     AcquireQueueScratch(scratch);
-    if (scratch.queue_base == NULL) {
+    if (scratch.queue_base == nullptr) {
       return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
     }
   }
 
   // Create an HW AQL queue
-  AqlQueue* hw_queue = new AqlQueue(this, size, node_id(), scratch,
-                                    event_callback, data, is_kv_device_);
-  if (hw_queue && hw_queue->IsValid()) {
-    // return queue
-    *queue = hw_queue;
-    return HSA_STATUS_SUCCESS;
-  }
-  // If reached here its always an ERROR.
-  delete hw_queue;
-  ReleaseQueueScratch(scratch.queue_base);
-  return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
+  *queue = new AqlQueue(this, size, node_id(), scratch, event_callback, data, is_kv_device_);
+  scratchGuard.Dismiss();
+  return HSA_STATUS_SUCCESS;
 }
 
 void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) {

runtime/hsa-runtime/core/runtime/host_queue.cpp

@@ -50,23 +50,19 @@ namespace core {
 int HostQueue::rtti_id_ = 0;
 std::atomic<uint32_t> HostQueue::queue_count_(0x80000000);
 
-HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size,
-                     hsa_queue_type32_t type, uint32_t features,
-                     hsa_signal_t doorbell_signal)
-    : Queue(),
-      size_(ring_size),
-      active_(false) {
-  if (!Shared::IsSharedObjectAllocationValid()) {
-    return;
-  }
-
+HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t type,
+                     uint32_t features, hsa_signal_t doorbell_signal)
+    : Queue(), size_(ring_size) {
   HSA::hsa_memory_register(this, sizeof(HostQueue));
+  MAKE_NAMED_SCOPE_GUARD(registerGuard,
+                         [&]() { HSA::hsa_memory_deregister(this, sizeof(HostQueue)); });
 
   const size_t queue_buffer_size = size_ * sizeof(AqlPacket);
   if (HSA_STATUS_SUCCESS !=
       HSA::hsa_memory_allocate(region, queue_buffer_size, &ring_)) {
-    return;
+    throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "Host queue buffer alloc failed\n");
   }
+  MAKE_NAMED_SCOPE_GUARD(bufferGuard, [&]() { HSA::hsa_memory_free(&ring_); });
 
   assert(IsMultipleOf(ring_, kRingAlignment));
   assert(ring_ != NULL);
@@ -88,14 +84,11 @@ HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size,
   AMD_HSA_BITS_SET(
       amd_queue_.queue_properties, AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, 0);
 
-  active_ = true;
+  bufferGuard.Dismiss();
+  registerGuard.Dismiss();
 }
 
 HostQueue::~HostQueue() {
-  if (!Shared::IsSharedObjectAllocationValid()) {
-    return;
-  }
-
   HSA::hsa_memory_free(ring_);
   HSA::hsa_memory_deregister(this, sizeof(HostQueue));
 }

runtime/hsa-runtime/core/runtime/hsa.cpp

@@ -644,8 +644,8 @@ hsa_status_t hsa_queue_create(
   TRY;
   IS_OPEN();
 
-  if ((queue == NULL) || (size == 0) || (!IsPowerOfTwo(size)) ||
-      (type < HSA_QUEUE_TYPE_MULTI) || (type > HSA_QUEUE_TYPE_SINGLE)) {
+  if ((queue == nullptr) || (size == 0) || (!IsPowerOfTwo(size)) || (type < HSA_QUEUE_TYPE_MULTI) ||
+      (type > HSA_QUEUE_TYPE_SINGLE)) {
     return HSA_STATUS_ERROR_INVALID_ARGUMENT;
   }
 
@@ -662,22 +662,17 @@ hsa_status_t hsa_queue_create(
     return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION;
   }
 
-  if (callback == NULL) callback = core::Queue::DefaultErrorHandler;
+  if (callback == nullptr) callback = core::Queue::DefaultErrorHandler;
 
-  core::Queue* cmd_queue = NULL;
+  core::Queue* cmd_queue = nullptr;
   status = agent->QueueCreate(size, type, callback, data, private_segment_size,
                               group_segment_size, &cmd_queue);
-  if (cmd_queue != NULL) {
-    *queue = core::Queue::Convert(cmd_queue);
-    if (*queue == NULL) {
-      delete cmd_queue;
-      return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
-    }
-  } else {
-    *queue = NULL;
-  }
+  if (status != HSA_STATUS_SUCCESS) return status;
 
+  assert(cmd_queue != nullptr && "Queue not returned but status was success.\n");
+  *queue = core::Queue::Convert(cmd_queue);
   return status;
+
   CATCH;
 }
 
@@ -701,13 +696,7 @@ hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size,
   const core::Signal* signal = core::Signal::Convert(doorbell_signal);
   IS_VALID(signal);
 
-  core::HostQueue* host_queue =
-      new core::HostQueue(region, size, type, features, doorbell_signal);
-
-  if (!host_queue->active()) {
-    delete host_queue;
-    return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
-  }
+  core::HostQueue* host_queue = new core::HostQueue(region, size, type, features, doorbell_signal);
 
   *queue = core::Queue::Convert(host_queue);
 

runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp

@@ -130,9 +130,9 @@ hsa_status_t handleException() {
   } catch (const std::bad_alloc& e) {
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
   } catch (const hsa_exception& e) {
+    debug_print("HSA exception: %s\n", e.what());
     return e.error_code();
-    // } catch (std::nested_exception& e) {
-    // Rethrow exceptions caught from callbacks.
+    // } catch (std::nested_exception& e) {  // Rethrow exceptions from callbacks
     //  e.rethrow_nested();
     //  return HSA_STATUS_ERROR;
   } catch (const std::exception& e) {
@@ -803,6 +803,8 @@ hsa_status_t hsa_amd_queue_intercept_create(
     void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data,
     uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue) {
   TRY;
+  IS_OPEN();
+  IS_BAD_PTR(queue);
   hsa_queue_t* lower_queue;
   hsa_status_t err = HSA::hsa_queue_create(agent_handle, size, type, callback, data,
                                            private_segment_size, group_segment_size, &lower_queue);
@@ -811,11 +813,7 @@ hsa_status_t hsa_amd_queue_intercept_create(
 
   std::unique_ptr<core::InterceptQueue> upperQueue(new core::InterceptQueue(std::move(lowerQueue)));
 
-  if (!upperQueue->IsSharedObjectAllocationValid()) return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
-
-  *queue = core::Queue::Convert(upperQueue.get());
-  upperQueue.release();
-
+  *queue = core::Queue::Convert(upperQueue.release());
   return HSA_STATUS_SUCCESS;
   CATCH;
 }

runtime/hsa-runtime/core/runtime/intercept_queue.cpp

@@ -75,20 +75,8 @@ InterceptQueue::InterceptQueue(std::unique_ptr<Queue> queue)
       retry_index_(0),
       quit_(false),
       active_(true) {
-  if (Queue::Shared::shared_object() == NULL) {
-    return;
-    // TODO skeely: cleanup queue constructor failure.
-    // throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "Failed to allocate Queue ABI
-    // block.\n");
-  }
-
-  void* buffer = Runtime::runtime_singleton_->system_allocator()(
-      wrapped->amd_queue_.hsa_queue.size * sizeof(hsa_agent_dispatch_packet_t), 4096,
-      MemoryRegion::AllocateNoFlags);
-  if (buffer == nullptr) throw std::bad_alloc();
-  MAKE_NAMED_SCOPE_GUARD(buffGuard,
-                         [&]() { Runtime::runtime_singleton_->system_deallocator()(buffer); });
-  amd_queue_.hsa_queue.base_address = buffer;
+  buffer_ = SharedArray<AqlPacket, 4096>(wrapped->amd_queue_.hsa_queue.size);
+  amd_queue_.hsa_queue.base_address = reinterpret_cast<void*>(&buffer_[0]);
 
   // Match the queue's signal ABI block to async_doorbell_'s
   // This allows us to use the queue's signal ABI block from devices to trigger async_doorbell while
@@ -109,14 +97,9 @@ InterceptQueue::InterceptQueue(std::unique_ptr<Queue> queue)
   AddInterceptor(Submit, this);
 
   sigGuard.Dismiss();
-  buffGuard.Dismiss();
 }
 
 InterceptQueue::~InterceptQueue() {
-  if (Queue::Shared::shared_object() == NULL) {
-    return;
-  }
-
   active_ = false;
 
   // Kill the async doorbell handler
@@ -128,9 +111,6 @@ InterceptQueue::~InterceptQueue() {
   if (val != 0)
     async_doorbell_->WaitRelaxed(HSA_SIGNAL_CONDITION_EQ, 0, -1, HSA_WAIT_STATE_BLOCKED);
   async_doorbell_->DestroySignal();
-
-  // Release buffer resources
-  Runtime::runtime_singleton_->system_deallocator()(amd_queue_.hsa_queue.base_address);
 }
 
 bool InterceptQueue::HandleAsyncDoorbell(hsa_signal_value_t value, void* arg) {

runtime/hsa-runtime/core/util/utils.h

@@ -120,12 +120,11 @@ static __forceinline unsigned long long int strtoull(const char* str,
 #endif
 
 // A macro to disallow the copy and move constructor and operator= functions
-// This should be used in the private: declarations for a class
-#define DISALLOW_COPY_AND_ASSIGN(TypeName) \
-  TypeName(const TypeName&);               \
-  TypeName(TypeName&&);                    \
-  void operator=(const TypeName&);         \
-  void operator=(TypeName&&);
+#define DISALLOW_COPY_AND_ASSIGN(TypeName)                                                         \
+  TypeName(const TypeName&) = delete;                                                              \
+  TypeName(TypeName&&) = delete;                                                                   \
+  void operator=(const TypeName&) = delete;                                                        \
+  void operator=(TypeName&&) = delete;
 
 template <typename lambda>
 class ScopeGuard {