P4 to Git Change 2009907 by ssahasra@ssahasra-hip-vdi on 2019/10/08 11:19:25

SWDEV-204782 - store extra information per HSA queue The new struct QueueInfo is used to store metadata about each HSA queue. For hostcall, this structure will eventually contain a pointer to the hostcall buffer allocated to each HSA queue. Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#135 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.hpp#41 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#88 edit [ROCm/clr commit: 1820fe21cf]
2019-10-08 10:27:38 -05:00
parent 2fb7d0e5cc
commit 6f4f575ac7
3 changed files with 78 additions and 39 deletions
@@ -1821,5 +1821,69 @@ bool Device::SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeI
  bool result = true;
  return result;
 }
+
+hsa_queue_t *Device::acquireQueue(uint32_t queue_size_hint) {
+  assert(queuePool_.size() <= GPU_MAX_HW_QUEUES);
+  LogPrintfInfo("number of allocated hardware queues: %d, maximum: %d",
+                queuePool_.size(), GPU_MAX_HW_QUEUES);
+
+  // If we have reached the max number of queues, reuse an existing queue,
+  // choosing the one with the least number of users.
+  if (queuePool_.size() == GPU_MAX_HW_QUEUES) {
+    typedef decltype(queuePool_)::const_reference PoolRef;
+    auto lowest = std::min_element(queuePool_.begin(), queuePool_.end(),
+                                   [] (PoolRef A, PoolRef B) {
+                                     return A.second.refCount < B.second.refCount;
+                                   });
+    LogPrintfInfo("selected queue with least refCount: %p (%d)",
+                  lowest->first, lowest->second.refCount);
+    lowest->second.refCount++;
+    return lowest->first;
+  }
+
+  // Else create a new queue. This also includes the initial state where there
+  // is no queue.
+  uint32_t queue_max_packets = 0;
+  if (HSA_STATUS_SUCCESS !=
+      hsa_agent_get_info(_bkendDevice, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_max_packets)) {
+    return nullptr;
+  }
+  auto queue_size = (queue_max_packets < queue_size_hint) ? queue_max_packets : queue_size_hint;
+
+  hsa_queue_t *queue;
+  while (hsa_queue_create(_bkendDevice, queue_size, HSA_QUEUE_TYPE_MULTI, nullptr, nullptr,
+                          std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
+                          &queue) != HSA_STATUS_SUCCESS) {
+    queue_size >>= 1;
+    if (queue_size < 64) {
+      return nullptr;
+    }
+  }
+  LogPrintfInfo("created hardware queue %p with size %d",
+                queue, queue_size);
+  hsa_amd_profiling_set_profiler_enabled(queue, 1);
+  auto result = queuePool_.emplace(std::make_pair(queue, QueueInfo()));
+  assert(result.second && "QueueInfo already exists");
+  auto &qInfo = result.first->second;
+  qInfo.refCount = 1;
+  return queue;
 }
+
+void Device::releaseQueue(hsa_queue_t* queue) {
+  auto qIter = queuePool_.find(queue);
+  assert(qIter != queuePool_.end());
+
+  auto &qInfo = qIter->second;
+  assert(qInfo.refCount > 0);
+  qInfo.refCount--;
+  if (qInfo.refCount != 0) {
+      return;
+  }
+  LogPrintfInfo("deleting hardware queue %p with refCount 0", queue);
+
+  hsa_queue_destroy(queue);
+  queuePool_.erase(qIter);
+}
+
+} // namespace roc
 #endif  // WITHOUT_HSA_BACKEND
@@ -406,8 +406,6 @@ class Device : public NullDevice {

  VirtualGPU* xferQueue() const;

-  std::map<hsa_queue_t*, int>& QueuePool() { return queue_pool_; }
-
  hsa_amd_memory_pool_t SystemSegment() const { return system_segment_; }

  hsa_amd_memory_pool_t SystemCoarseSegment() const { return system_coarse_segment_; }
@@ -442,10 +440,17 @@ class Device : public NullDevice {
  std::atomic<size_t> freeMem_;   //!< Total of free memory available
  mutable amd::Monitor vgpusAccess_;     //!< Lock to serialise virtual gpu list access
  bool hsa_exclusive_gpu_access_;  //!< TRUE if current device was moved into exclusive GPU access mode
-  std::map<hsa_queue_t*, int> queue_pool_;  //!< Pool of HSA queues for recycling
+
+  struct QueueInfo {
+    int refCount;
+  };
+  std::map<hsa_queue_t*, QueueInfo> queuePool_;  //!< Pool of HSA queues for recycling

 public:
  amd::Atomic<uint> numOfVgpus_;  //!< Virtual gpu unique index
+
+  hsa_queue_t *acquireQueue(uint32_t queue_size_hint);
+  void releaseQueue(hsa_queue_t*);
 };                                // class roc::Device
 }  // namespace roc

@@ -555,6 +555,7 @@ bool VirtualGPU::releaseGpuMemoryFence() {

 VirtualGPU::VirtualGPU(Device& device)
    : device::VirtualDevice(device),
+      gpu_queue_(nullptr),
      roc_device_(device),
      virtualQueue_(nullptr),
      deviceQueueSize_(0),
@@ -631,12 +632,9 @@ VirtualGPU::~VirtualGPU() {
  for (uint idx = index(); idx < roc_device_.vgpus().size(); ++idx) {
    roc_device_.vgpus()[idx]->index_--;
  }
-  // Decrement the counter
-  roc_device_.QueuePool()[gpu_queue_]--;
-  // Release the queue if the counter is 0
-  if (roc_device_.QueuePool()[gpu_queue_] == 0) {
-    hsa_status_t err = hsa_queue_destroy(gpu_queue_);
-    roc_device_.QueuePool().erase(gpu_queue_);
+
+  if (gpu_queue_) {
+    roc_device_.releaseQueue(gpu_queue_);
  }
 }

@@ -646,38 +644,10 @@ bool VirtualGPU::create(bool profilingEna) {
    return false;
  }

-  uint32_t queue_max_packets = 0;
-  if (HSA_STATUS_SUCCESS !=
-      hsa_agent_get_info(gpu_device_, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_max_packets)) {
-    return false;
-  }
-
  // Pick a reasonable queue size
  uint32_t queue_size = 1024;
-  queue_size = (queue_max_packets < queue_size) ? queue_max_packets : queue_size;
-  if (roc_device_.QueuePool().size() < GPU_MAX_HW_QUEUES) {
-    while (hsa_queue_create(gpu_device_, queue_size, HSA_QUEUE_TYPE_MULTI, nullptr, nullptr,
-                          std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
-                          &gpu_queue_) != HSA_STATUS_SUCCESS) {
-      queue_size >>= 1;
-      if (queue_size < 64) {
-        return false;
-      }
-    }
-    hsa_amd_profiling_set_profiler_enabled(gpu_queue(), 1);
-    roc_device_.QueuePool().insert({gpu_queue_, 1});
-  } else {
-    int usage = std::numeric_limits<int>::max();
-    // Loop through all allocated queues and find the lowest usage
-    for (const auto it : roc_device_.QueuePool()) {
-      if (it.second < usage) {
-        gpu_queue_ = it.first;
-        usage = it.second;
-      }
-    }
-    // Increment the usage of the current queue
-    roc_device_.QueuePool()[gpu_queue_]++;
-  }
+  gpu_queue_ = roc_device_.acquireQueue(queue_size);
+  if (!gpu_queue_) return false;

  if (!initPool(dev().settings().kernargPoolSize_, (profilingEna) ? queue_size : 0)) {
    LogError("Couldn't allocate arguments/signals for the queue");