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 [ROCm/ROCR-Runtime commit: f312a7386e]
2017-10-20 07:12:04 -05:00
parent 2406218416
commit e2efba0676
13 changed files with 181 additions and 162 deletions
@@ -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;
 }

@@ -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) {
@@ -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));
 }
@@ -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);

@@ -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;
 }
@@ -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) {