Initial GWS queue support.

Queues should transition to ref counting for all queues eventually. That cleanup will be part of shared queue pooling support. Change-Id: I217ff5d573156678b9559da6fb81baa8cd31c617 [ROCm/ROCR-Runtime commit: 0a43a107b1]
2019-10-11 20:21:56 -05:00
parent a161850fcf
commit 3bcde37b58
8 changed files with 101 additions and 9 deletions
@@ -74,6 +74,9 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
  /// @brief Change the scheduling priority of the queue
  hsa_status_t SetPriority(HSA_QUEUE_PRIORITY priority) override;

+  /// @brief Destroy ref counted queue
+  void Destroy() override;
+
  /// @brief Atomically reads the Read index of with Acquire semantics
  ///
  /// @return uint64_t Value of read index
@@ -256,6 +256,9 @@ class GpuAgent : public GpuAgentInt {
                           uint32_t group_segment_size,
                           core::Queue** queue) override;

+  // @brief Decrement GWS ref count.
+  void GWSRelease();
+
  // @brief Override from amd::GpuAgentInt.
  void AcquireQueueScratch(ScratchInfo& scratch) override;

@@ -488,6 +491,9 @@ class GpuAgent : public GpuAgentInt {
  // @brief Create internal queues and blits.
  void InitDma();

+  // @brief Setup GWS accessing queue.
+  void InitGWS();
+

  // Bind index of peer device that is connected via xGMI links
  lazy_ptr<core::Blit>& GetXgmiBlit(const core::Agent& peer_agent);
@@ -504,6 +510,13 @@ class GpuAgent : public GpuAgentInt {
  // @brief Alternative aperture size. Only on KV.
  size_t ape1_size_;

+  // @brief Queue with GWS access.
+  struct {
+    lazy_ptr<core::Queue> queue_;
+    int ref_ct_;
+    KernelMutex lock_;
+  } gws_queue_;
+
  DISALLOW_COPY_AND_ASSIGN(GpuAgent);
 };

@@ -150,6 +150,8 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue {

  virtual ~Queue() {}

+  virtual void Destroy() { delete this; }
+
  /// @brief Returns the handle of Queue's public data type
  ///
  /// @param queue Pointer to an instance of Queue implementation object
@@ -306,6 +306,14 @@ AqlQueue::~AqlQueue() {
  core::Runtime::runtime_singleton_->system_deallocator()(pm4_ib_buf_);
 }

+void AqlQueue::Destroy() {
+  if (amd_queue_.hsa_queue.type & HSA_QUEUE_TYPE_COOPERATIVE) {
+    agent_->GWSRelease();
+    return;
+  }
+  delete this;
+}
+
 uint64_t AqlQueue::LoadReadIndexAcquire() {
  return atomic::Load(&amd_queue_.read_dispatch_id, std::memory_order_acquire);
 }
@@ -580,6 +580,35 @@ void GpuAgent::InitDma() {
  for (uint32_t idx = DefaultBlitCount; idx < blit_cnt_; idx++) {
    blits_[idx].reset([blit_lambda, this]() { return blit_lambda(true, queues_[QueueUtility]); });
  }
+
+  // GWS queues.
+  InitGWS();
+}
+
+void GpuAgent::InitGWS() {
+  gws_queue_.queue_.reset([this]() {
+    if (properties_.NumGws == 0) return (core::Queue*)nullptr;
+    std::unique_ptr<core::Queue> queue(CreateInterceptibleQueue());
+    if (queue == nullptr)
+      throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
+                               "Internal queue creation failed.");
+
+    uint32_t discard;
+    auto status = hsaKmtAllocQueueGWS(queue->amd_queue_.hsa_queue.id, 1, &discard);
+    if (status != HSAKMT_STATUS_SUCCESS)
+      throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "GWS allocation failed.");
+
+    queue->amd_queue_.hsa_queue.type = HSA_QUEUE_TYPE_COOPERATIVE | HSA_QUEUE_TYPE_MULTI;
+    gws_queue_.ref_ct_ = 0;
+    return queue.release();
+  });
+}
+
+void GpuAgent::GWSRelease() {
+  ScopedAcquire<KernelMutex> lock(&gws_queue_.lock_);
+  gws_queue_.ref_ct_--;
+  if (gws_queue_.ref_ct_ != 0) return;
+  InitGWS();
 }

 void GpuAgent::PreloadBlits() {
@@ -869,6 +898,9 @@ hsa_status_t GpuAgent::GetInfo(hsa_agent_info_t attribute, void* value) const {
    case HSA_AMD_AGENT_INFO_DOMAIN:
      *((uint32_t*)value) = static_cast<uint32_t>(properties_.Domain);
      break;
+    case HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES:
+      *((bool*)value) = properties_.NumGws != 0;
+      break;
    default:
      return HSA_STATUS_ERROR_INVALID_ARGUMENT;
      break;
@@ -881,6 +913,18 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type,
                                   void* data, uint32_t private_segment_size,
                                   uint32_t group_segment_size,
                                   core::Queue** queue) {
+  // Handle GWS queues.
+  if (queue_type & HSA_QUEUE_TYPE_COOPERATIVE) {
+    ScopedAcquire<KernelMutex> lock(&gws_queue_.lock_);
+    auto ret = (*gws_queue_.queue_).get();
+    if (ret != nullptr) {
+      gws_queue_.ref_ct_++;
+      *queue = ret;
+      return HSA_STATUS_SUCCESS;
+    }
+    return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION;
+  }
+
  // AQL queues must be a power of two in length.
  if (!IsPowerOfTwo(size)) {
    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
@@ -687,8 +687,8 @@ hsa_status_t hsa_queue_create(
  TRY;
  IS_OPEN();

-  if ((queue == nullptr) || (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_COOPERATIVE)) {
    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
  }

@@ -701,7 +701,12 @@ hsa_status_t hsa_queue_create(
  assert(HSA_STATUS_SUCCESS == status);

  if (agent_queue_type == HSA_QUEUE_TYPE_SINGLE &&
-      type != HSA_QUEUE_TYPE_SINGLE) {
+      ((type & HSA_QUEUE_TYPE_SINGLE) != HSA_QUEUE_TYPE_SINGLE)) {
+    return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION;
+  }
+
+  if ((type & HSA_QUEUE_TYPE_COOPERATIVE) &&
+      ((type & HSA_QUEUE_TYPE_SINGLE) != HSA_QUEUE_TYPE_MULTI)) {
    return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION;
  }

@@ -758,7 +763,7 @@ hsa_status_t hsa_queue_destroy(hsa_queue_t* queue) {
  IS_BAD_PTR(queue);
  core::Queue* cmd_queue = core::Queue::Convert(queue);
  IS_VALID(cmd_queue);
-  delete cmd_queue;
+  cmd_queue->Destroy();
  return HSA_STATUS_SUCCESS;
  CATCH;
 }
@@ -2193,7 +2193,19 @@ typedef enum {
   * that support a single producer may be more efficient than queues supporting
   * multiple producers.
   */
-  HSA_QUEUE_TYPE_SINGLE = 1
+  HSA_QUEUE_TYPE_SINGLE = 1,
+  /**
+   * Queue supports cooperative dispatches able to use GWS synchronization.
+   * Queues of this type must also be of type HSA_QUEUE_TYPE_MULTI and
+   * may be limited in number.  The runtime may return the same queue to serve
+   * multiple hsa_queue_create calls when this type is given.  Callers must
+   * inspect the returned queue to discover queue size.  Queues of this type
+   * are reference counted and require a matching number of hsa_queue_destroy
+   * calls to release.  Use of multiproducer queue mechanics is required.  See
+   * ::HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES to query agent support for this
+   * type.
+   */
+  HSA_QUEUE_TYPE_COOPERATIVE = 2
 } hsa_queue_type_t;

 /**
@@ -2300,9 +2312,9 @@ typedef struct hsa_queue_s {
 * created queue is the maximum of @p size and the value of
 * ::HSA_AGENT_INFO_QUEUE_MIN_SIZE in @p agent.
 *
- * @param[in] type Type of the queue. If the value of
- * ::HSA_AGENT_INFO_QUEUE_TYPE in @p agent is ::HSA_QUEUE_TYPE_SINGLE, then @p
- * type must also be ::HSA_QUEUE_TYPE_SINGLE.
+ * @param[in] type Type of the queue, a bitwise OR of hsa_queue_type_t values.
+ * If the value of ::HSA_AGENT_INFO_QUEUE_TYPE in @p agent is ::HSA_QUEUE_TYPE_SINGLE,
+ * then @p type must also be ::HSA_QUEUE_TYPE_SINGLE.
 *
 * @param[in] callback Callback invoked by the HSA runtime for every
 * asynchronous event related to the newly created queue. May be NULL. The HSA
@@ -158,7 +158,12 @@ typedef enum hsa_amd_agent_info_s {
   * to give the full physical location of the Agent.
   * The type of this attribute is uint32_t.
   */
-  HSA_AMD_AGENT_INFO_DOMAIN = 0xA00F
+  HSA_AMD_AGENT_INFO_DOMAIN = 0xA00F,
+  /**
+   * Queries for support of cooperative queues.  See ::HSA_QUEUE_TYPE_COOPERATIVE.
+   * The type of this attribute is bool.
+   */
+  HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES = 0xA010
 } hsa_amd_agent_info_t;

 typedef struct hsa_amd_hdp_flush_s {