Update ROCr implementation of Queue ID

Change-Id: Iec48b1978e4d01563e71cfb58aed8f1bbc446443 [ROCm/ROCR-Runtime commit: f7350c6020]
2020-06-23 15:31:17 -05:00
parent 950205ee0f
commit 3955da129e
5 changed files with 16 additions and 8 deletions
@@ -330,7 +330,14 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue {

  hsa_queue_t* public_handle_;

+  /// Next available queue id.
+  uint64_t GetQueueId() { return hsa_queue_counter_++; }
+
 private:
+
+  // HSA Queue ID - used to bind a unique ID
+  static std::atomic<uint64_t> hsa_queue_counter_;
+
  DISALLOW_COPY_AND_ASSIGN(Queue);
 };
 }   //  namespace core
@@ -251,11 +251,6 @@ class Runtime {
  /// set.
  hsa_status_t GetSystemInfo(hsa_system_info_t attribute, void* value);

-  /// @brief Query next available queue id.
-  ///
-  /// @retval Next available queue id.
-  uint32_t GetQueueId();
-
  /// @brief Register a callback function @p handler that is associated with
  /// @p signal to asynchronous event monitor thread.
  ///
@@ -225,9 +225,12 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
  // Complete populating the doorbell signal structure.
  signal_.legacy_hardware_doorbell_ptr =
      (volatile uint32_t*)queue_rsrc.Queue_DoorBell;
-  // Complete populating the amd_queue_ structure.
-  amd_queue_.hsa_queue.id = queue_id_ = queue_rsrc.QueueId;

+  // Bind Id of Queue such that is unique i.e. it is not re-used by another
+  // queue (AQL, HOST) in the same process during its lifetime.
+  amd_queue_.hsa_queue.id = this->GetQueueId();
+
+  queue_id_ = queue_rsrc.QueueId;
  MAKE_NAMED_SCOPE_GUARD(QueueGuard, [&]() { hsaKmtDestroyQueue(queue_id_); });

  MAKE_NAMED_SCOPE_GUARD(EventGuard, [&]() {
@@ -77,7 +77,7 @@ HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t
  amd_queue_.hsa_queue.base_address = ring_;
  amd_queue_.hsa_queue.size = size_;
  amd_queue_.hsa_queue.doorbell_signal = doorbell_signal;
-  amd_queue_.hsa_queue.id = queue_count_++;
+  amd_queue_.hsa_queue.id = this->GetQueueId();
  amd_queue_.hsa_queue.type = type;
  amd_queue_.hsa_queue.features = features;
 #ifdef HSA_LARGE_MODEL
@@ -46,6 +46,9 @@
 namespace rocr {
 namespace core {

+// HSA Queue ID - used to bind a unique ID
+std::atomic<uint64_t> Queue::hsa_queue_counter_(0);
+
 void Queue::DefaultErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data) {
  if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) {
    const char* msg = "UNKNOWN ERROR";