Separate blit compute interface from queue creation

The runtime needs a queue on which to submit cache management commands.
Device-to-device blit copy already creates a queue unconditionally.
We can share this queue for both purposes.

This change restructures the BlitKernel interface to accept, rather than
create, a queue. GpuAgent creates queues as needed for both cache
management and blit compute.

Fix queue full detection in AcquireWriteIndex (<= vs <).

Change-Id: I61d0c6b9d04f2dba74872f0676ad791435778ba4


[ROCm/ROCR-Runtime commit: f7ab361347]

projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_blit_kernel.h

@@ -51,7 +51,7 @@
 namespace amd {
 class BlitKernel : public core::Blit {
  public:
-  explicit BlitKernel();
+  explicit BlitKernel(core::Queue* queue);
   virtual ~BlitKernel() override;
 
   /// @brief Initialize a blit kernel object.
@@ -175,12 +175,9 @@ class BlitKernel : public core::Blit {
   std::map<KernelType, KernelCode> kernels_;
 
   /// AQL queue for submitting the vector copy kernel.
-  hsa_queue_t* queue_;
+  core::Queue* queue_;
   uint32_t queue_bitmask_;
 
-  /// Index to track concurrent kernel launch.
-  volatile uint64_t cached_index_;
-
   /// Pointer to the kernel argument buffer.
   KernelArgs* kernarg_async_;
   uint32_t kernarg_async_mask_;

projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_agent.h

@@ -79,10 +79,11 @@ class GpuAgentInt : public core::Agent {
   // @retval HSA_STATUS_SUCCESS DMA queue initialization is successful.
   virtual void InitDma() = 0;
 
-  // @brief Initialize blit kernel object based on AQL queue.
+  // @brief Initialization hook invoked after tools library has loaded,
+  // to allow tools interception of interface functions.
   //
-  // @retval HSA_STATUS_SUCCESS blit kernel object initialization is successful.
-  virtual hsa_status_t InitBlitKernel() = 0;
+  // @retval HSA_STATUS_SUCCESS if initialization is successful.
+  virtual hsa_status_t PostToolsInit() = 0;
 
   // @brief Invoke the user provided callback for each region accessible by
   // this agent.
@@ -184,7 +185,7 @@ class GpuAgent : public GpuAgentInt {
   void InitDma() override;
 
   // @brief Override from core::Agent.
-  hsa_status_t InitBlitKernel() override;
+  hsa_status_t PostToolsInit() override;
 
   uint16_t GetMicrocodeVersion() const;
 
@@ -309,15 +310,18 @@ class GpuAgent : public GpuAgentInt {
   static const uint32_t minAqlSize_ = 0x1000;   // 4KB min
   static const uint32_t maxAqlSize_ = 0x20000;  // 8MB max
 
+  // @brief Create a queue through HSA API to allow tools to intercept.
+  core::Queue* CreateInterceptibleQueue();
+
   // @brief Create SDMA blit object.
   //
   // @retval NULL if SDMA blit creation and initialization failed.
   core::Blit* CreateBlitSdma();
 
-  // @brief Create Kernel blit object.
+  // @brief Create Kernel blit object using provided compute queue.
   //
   // @retval NULL if Kernel blit creation and initialization failed.
-  core::Blit* CreateBlitKernel();
+  core::Blit* CreateBlitKernel(core::Queue* queue);
 
   // @brief Invoke the user provided callback for every region in @p regions.
   //
@@ -359,15 +363,19 @@ class GpuAgent : public GpuAgentInt {
   // @brief Default scratch size per work item.
   size_t scratch_per_thread_;
 
-  // @brief Blit object to handle memory copy from system to device memory.
-  core::Blit* blit_h2d_;
+  // @brief Blit interfaces for each data path.
+  enum BlitEnum { BlitHostToDev, BlitDevToHost, BlitDevToDev, BlitCount };
 
-  // @brief Blit object to handle memory copy from device to system memory.
-  core::Blit* blit_d2h_;
+  core::Blit* blits_[BlitCount];
 
-  // @brief Blit object to handle memory copy from device to device memory, and
-  // memory fill.
-  core::Blit* blit_d2d_;
+  // @brief AQL queues for cache management and blit compute usage.
+  enum QueueEnum {
+    QueueUtility,   // Cache management and device to {host,device} blit compute
+    QueueBlitOnly,  // Host to device blit
+    QueueCount
+  };
+
+  core::Queue* queues_[QueueCount];
 
   // @brief Mutex to protect the update to coherency type.
   KernelMutex coherency_lock_;

projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_blit_kernel.cpp

@@ -515,10 +515,9 @@ static int kCopyMisalignedUnroll = GetKernelSourceParam("kCopyMisalignedUnroll")
 static int kFillVecWidth = GetKernelSourceParam("kFillVecWidth");
 static int kFillUnroll = GetKernelSourceParam("kFillUnroll");
 
-BlitKernel::BlitKernel()
+BlitKernel::BlitKernel(core::Queue* queue)
     : core::Blit(),
-      queue_(NULL),
-      cached_index_(0),
+      queue_(queue),
       kernarg_async_(NULL),
       kernarg_async_mask_(0),
       kernarg_async_counter_(0),
@@ -529,40 +528,18 @@ BlitKernel::BlitKernel()
 BlitKernel::~BlitKernel() {}
 
 hsa_status_t BlitKernel::Initialize(const core::Agent& agent) {
-  hsa_agent_t agent_handle = agent.public_handle();
+  queue_bitmask_ = queue_->public_handle()->size - 1;
 
-  uint32_t features = 0;
-  hsa_status_t status =
-      HSA::hsa_agent_get_info(agent_handle, HSA_AGENT_INFO_FEATURE, &features);
-  if (status != HSA_STATUS_SUCCESS) {
-    return status;
-  }
-
-  if ((features & HSA_AGENT_FEATURE_KERNEL_DISPATCH) == 0) {
-    return HSA_STATUS_ERROR;
-  }
-
-  status = HSA::hsa_queue_create(agent_handle, 1024, HSA_QUEUE_TYPE_MULTI, NULL,
-                                 NULL, 0, 0, &queue_);
-
-  if (HSA_STATUS_SUCCESS != status) {
-    return status;
-  }
-
-  queue_bitmask_ = queue_->size - 1;
-
-  cached_index_ = 0;
-
-  status = HSA::hsa_signal_create(1, 0, NULL, &completion_signal_);
+  hsa_status_t status = HSA::hsa_signal_create(1, 0, NULL, &completion_signal_);
   if (HSA_STATUS_SUCCESS != status) {
     return status;
   }
 
   kernarg_async_ = reinterpret_cast<KernelArgs*>(
       core::Runtime::runtime_singleton_->system_allocator()(
-          queue_->size * AlignUp(sizeof(KernelArgs), 16), 16));
+          queue_->public_handle()->size * AlignUp(sizeof(KernelArgs), 16), 16));
 
-  kernarg_async_mask_ = queue_->size - 1;
+  kernarg_async_mask_ = queue_->public_handle()->size - 1;
 
   // Obtain the number of compute units in the underlying agent.
   const GpuAgent& gpuAgent = static_cast<const GpuAgent&>(agent);
@@ -598,10 +575,6 @@ hsa_status_t BlitKernel::Destroy(const core::Agent& agent) {
                            kernel_pair.second.code_buf_size_);
   }
 
-  if (queue_ != NULL) {
-    HSA::hsa_queue_destroy(queue_);
-  }
-
   if (kernarg_async_ != NULL) {
     core::Runtime::runtime_singleton_->system_deallocator()(kernarg_async_);
   }
@@ -661,7 +634,8 @@ hsa_status_t BlitKernel::SubmitLinearCopyCommand(
   barrier_packet.header = HSA_PACKET_TYPE_INVALID;
 
   hsa_barrier_and_packet_t* queue_buffer =
-      reinterpret_cast<hsa_barrier_and_packet_t*>(queue_->base_address);
+      reinterpret_cast<hsa_barrier_and_packet_t*>(
+          queue_->public_handle()->base_address);
 
   const size_t dep_signal_count = dep_signals.size();
   for (size_t i = 0; i < dep_signal_count; ++i) {
@@ -803,26 +777,20 @@ hsa_status_t BlitKernel::SubmitLinearFillCommand(void* ptr, uint32_t value,
 }
 
 hsa_status_t BlitKernel::EnableProfiling(bool enable) {
-  core::Queue* cmd_queue = core::Queue::Convert(queue_);
-  if (cmd_queue != NULL) {
-    AMD_HSA_BITS_SET(cmd_queue->amd_queue_.queue_properties,
-                     AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, enable);
-    return HSA_STATUS_SUCCESS;
-  }
-
-  return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
+  AMD_HSA_BITS_SET(queue_->amd_queue_.queue_properties,
+                   AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, enable);
+  return HSA_STATUS_SUCCESS;
 }
 
 uint64_t BlitKernel::AcquireWriteIndex(uint32_t num_packet) {
-  assert(queue_->size >= num_packet);
+  assert(queue_->public_handle()->size >= num_packet);
 
-  uint64_t write_index =
-      HSA::hsa_queue_add_write_index_acq_rel(queue_, num_packet);
+  uint64_t write_index = queue_->AddWriteIndexAcqRel(num_packet);
 
   while (true) {
     // Wait until we have room in the queue;
-    const uint64_t read_index = HSA::hsa_queue_load_read_index_relaxed(queue_);
-    if ((write_index - read_index) < queue_->size) {
+    const uint64_t read_index = queue_->LoadReadIndexRelaxed();
+    if ((write_index - read_index) <= queue_->public_handle()->size) {
       break;
     }
   }
@@ -831,19 +799,10 @@ uint64_t BlitKernel::AcquireWriteIndex(uint32_t num_packet) {
 }
 
 void BlitKernel::ReleaseWriteIndex(uint64_t write_index, uint32_t num_packet) {
-  // Launch packet.
-  while (true) {
-    // Make sure that the address before ::current_offset is already released.
-    // Otherwise the packet processor may read invalid packets.
-    uint64_t expected_offset = write_index;
-    if (atomic::Cas(&cached_index_, write_index + num_packet, expected_offset,
-                    std::memory_order_release) == expected_offset) {
-      // Update doorbel register with last packet id.
-      HSA::hsa_signal_store_release(queue_->doorbell_signal,
-                                    write_index + num_packet - 1);
-      break;
-    }
-  }
+  // Update doorbel register with last packet id.
+  core::Signal* doorbell =
+      core::Signal::Convert(queue_->public_handle()->doorbell_signal);
+  doorbell->StoreRelease(write_index + num_packet - 1);
 }
 
 hsa_status_t BlitKernel::FenceRelease(uint64_t write_index,
@@ -871,7 +830,8 @@ hsa_status_t BlitKernel::FenceRelease(uint64_t write_index,
 
   // Populate queue buffer with AQL packet.
   hsa_barrier_and_packet_t* queue_buffer =
-      reinterpret_cast<hsa_barrier_and_packet_t*>(queue_->base_address);
+      reinterpret_cast<hsa_barrier_and_packet_t*>(
+          queue_->public_handle()->base_address);
   std::atomic_thread_fence(std::memory_order_acquire);
   queue_buffer[(write_index + num_copy_packet) & queue_bitmask_] = packet;
   std::atomic_thread_fence(std::memory_order_release);
@@ -921,7 +881,8 @@ void BlitKernel::PopulateQueue(uint64_t index, uint64_t code_handle, void* args,
 
   // Populate queue buffer with AQL packet.
   hsa_kernel_dispatch_packet_t* queue_buffer =
-      reinterpret_cast<hsa_kernel_dispatch_packet_t*>(queue_->base_address);
+      reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
+          queue_->public_handle()->base_address);
   std::atomic_thread_fence(std::memory_order_acquire);
   queue_buffer[index & queue_bitmask_] = packet;
   std::atomic_thread_fence(std::memory_order_release);

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

@@ -71,9 +71,8 @@ GpuAgent::GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props)
     : GpuAgentInt(node),
       properties_(node_props),
       current_coherency_type_(HSA_AMD_COHERENCY_TYPE_COHERENT),
-      blit_h2d_(NULL),
-      blit_d2h_(NULL),
-      blit_d2d_(NULL),
+      blits_(),
+      queues_(),
       local_region_(NULL),
       is_kv_device_(false),
       trap_code_buf_(NULL),
@@ -135,28 +134,16 @@ GpuAgent::GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props)
 }
 
 GpuAgent::~GpuAgent() {
-  if (blit_h2d_ != NULL) {
-    hsa_status_t status = blit_h2d_->Destroy(*this);
-    assert(status == HSA_STATUS_SUCCESS);
-
-    delete blit_h2d_;
-    blit_h2d_ = NULL;
+  for (int i = 0; i < BlitCount; ++i) {
+    if (blits_[i] != NULL) {
+      hsa_status_t status = blits_[i]->Destroy(*this);
+      assert(status == HSA_STATUS_SUCCESS);
+      delete blits_[i];
+    }
   }
 
-  if (blit_d2h_ != NULL && blit_d2h_ != blit_d2d_) {
-    hsa_status_t status = blit_d2h_->Destroy(*this);
-    assert(status == HSA_STATUS_SUCCESS);
-
-    delete blit_d2h_;
-    blit_d2h_ = NULL;
-  }
-
-  if (blit_d2d_ != NULL) {
-    hsa_status_t status = blit_d2d_->Destroy(*this);
-    assert(status == HSA_STATUS_SUCCESS);
-
-    delete blit_d2d_;
-    blit_d2d_ = NULL;
+  for (int i = 0; i < QueueCount; ++i) {
+    delete queues_[i];
   }
 
   if (end_ts_base_addr_ != NULL) {
@@ -501,6 +488,21 @@ hsa_status_t GpuAgent::VisitRegion(
   return HSA_STATUS_SUCCESS;
 }
 
+core::Queue* GpuAgent::CreateInterceptibleQueue() {
+  // Until tools runtime is merged in we need to use HSA API
+  // rather than GpuAgent::QueueCreate to allow interception.
+  hsa_queue_t* queue_handle;
+  hsa_status_t status =
+      HSA::hsa_queue_create(public_handle(), minAqlSize_, HSA_QUEUE_TYPE_MULTI,
+                            NULL, NULL, 0, 0, &queue_handle);
+
+  if (status != HSA_STATUS_SUCCESS) {
+    return NULL;
+  }
+
+  return core::Queue::Convert(queue_handle);
+}
+
 core::Blit* GpuAgent::CreateBlitSdma() {
   BlitSdma* sdma = new BlitSdma();
 
@@ -513,8 +515,8 @@ core::Blit* GpuAgent::CreateBlitSdma() {
   return sdma;
 }
 
-core::Blit* GpuAgent::CreateBlitKernel() {
-  BlitKernel* kernl = new BlitKernel();
+core::Blit* GpuAgent::CreateBlitKernel(core::Queue* queue) {
+  BlitKernel* kernl = new BlitKernel(queue);
 
   if (kernl->Initialize(*this) != HSA_STATUS_SUCCESS) {
     kernl->Destroy(*this);
@@ -536,25 +538,29 @@ void GpuAgent::InitDma() {
       // Try create SDMA blit first.
       if (core::Runtime::runtime_singleton_->flag().enable_sdma() &&
           (HSA_PROFILE_BASE == profile_)) {
-        blit_h2d_ = CreateBlitSdma();
-        blit_d2h_ = CreateBlitSdma();
+        blits_[BlitHostToDev] = CreateBlitSdma();
+        blits_[BlitDevToHost] = CreateBlitSdma();
 
-        if (blit_h2d_ != NULL && blit_d2h_ != NULL) {
+        if (blits_[BlitHostToDev] != NULL && blits_[BlitDevToHost] != NULL) {
           blit_initialized_.store(true, std::memory_order_release);
           return;
         }
       }
 
       // Fall back to blit kernel if SDMA is unavailable.
-      assert(blit_h2d_ == NULL || blit_d2h_ == NULL);
+      if (blits_[BlitHostToDev] == NULL) {
+        // Create a dedicated compute queue for host-to-device blits.
+        queues_[QueueBlitOnly] = CreateInterceptibleQueue();
+        assert(queues_[QueueBlitOnly] != NULL && "Queue creation failed");
 
-      if (blit_h2d_ == NULL) {
-        blit_h2d_ = CreateBlitKernel();
+        blits_[BlitHostToDev] = CreateBlitKernel(queues_[QueueBlitOnly]);
+        assert(blits_[BlitHostToDev] != NULL && "Blit creation failed");
       }
 
-      if (blit_d2h_ == NULL) {
-        // Share device-to-host queue with device-to-device.
-        blit_d2h_ = blit_d2d_;
+      if (blits_[BlitDevToHost] == NULL) {
+        // Share utility queue with device-to-host blits.
+        blits_[BlitDevToHost] = CreateBlitKernel(queues_[QueueUtility]);
+        assert(blits_[BlitDevToHost] != NULL && "Blit creation failed");
       }
 
       blit_initialized_.store(true, std::memory_order_release);
@@ -562,23 +568,26 @@ void GpuAgent::InitDma() {
   }
 }
 
-hsa_status_t GpuAgent::InitBlitKernel() {
-  // Unlike InitDma, this function is not designed for lazy initialization.
-  // So checking the state without  double checked locking is fine.
-  if (blit_d2d_ == NULL) {
-    blit_d2d_ = CreateBlitKernel();
-  }
+hsa_status_t GpuAgent::PostToolsInit() {
+  // Defer utility queue creation to allow tools to intercept.
+  queues_[QueueUtility] = CreateInterceptibleQueue();
 
-  return (blit_d2d_ != NULL) ? HSA_STATUS_SUCCESS
-                             : HSA_STATUS_ERROR_OUT_OF_RESOURCES;
-}
-
-hsa_status_t GpuAgent::DmaCopy(void* dst, const void* src, size_t size) {
-  if (blit_d2d_ == NULL) {
+  if (queues_[QueueUtility] == NULL) {
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
   }
 
-  return blit_d2d_->SubmitLinearCopyCommand(dst, src, size);
+  // Share utility queue with device-to-device blits.
+  blits_[BlitDevToDev] = CreateBlitKernel(queues_[QueueUtility]);
+
+  if (blits_[BlitDevToDev] == NULL) {
+    return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
+hsa_status_t GpuAgent::DmaCopy(void* dst, const void* src, size_t size) {
+  return blits_[BlitDevToDev]->SubmitLinearCopyCommand(dst, src, size);
 }
 
 hsa_status_t GpuAgent::DmaCopy(void* dst, core::Agent& dst_agent,
@@ -589,11 +598,11 @@ hsa_status_t GpuAgent::DmaCopy(void* dst, core::Agent& dst_agent,
   core::Blit* blit =
       (src_agent.device_type() == core::Agent::kAmdCpuDevice &&
        dst_agent.device_type() == core::Agent::kAmdGpuDevice)
-          ? blit_h2d_
+          ? blits_[BlitHostToDev]
           : (src_agent.device_type() == core::Agent::kAmdGpuDevice &&
              dst_agent.device_type() == core::Agent::kAmdCpuDevice)
-                ? blit_d2h_
-                : blit_d2d_;
+                ? blits_[BlitDevToHost]
+                : blits_[BlitDevToDev];
 
   if (blit == NULL) {
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
@@ -612,11 +621,7 @@ hsa_status_t GpuAgent::DmaCopy(void* dst, core::Agent& dst_agent,
 }
 
 hsa_status_t GpuAgent::DmaFill(void* ptr, uint32_t value, size_t count) {
-  if (blit_d2d_ == NULL) {
-    return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
-  }
-
-  return blit_d2d_->SubmitLinearFillCommand(ptr, value, count);
+  return blits_[BlitDevToDev]->SubmitLinearFillCommand(ptr, value, count);
 }
 
 hsa_status_t GpuAgent::EnableDmaProfiling(bool enable) {
@@ -624,10 +629,9 @@ hsa_status_t GpuAgent::EnableDmaProfiling(bool enable) {
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
   }
 
-  core::Blit* blit[3] = {blit_h2d_, blit_d2h_, blit_d2d_};
-  for (int i = 0; i < 3; ++i) {
-    if (blit[i] != NULL) {
-      const hsa_status_t stat = blit[i]->EnableProfiling(enable);
+  for (int i = 0; i < BlitCount; ++i) {
+    if (blits_[i] != NULL) {
+      const hsa_status_t stat = blits_[i]->EnableProfiling(enable);
       if (stat != HSA_STATUS_SUCCESS) {
         return stat;
       }

projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp

@@ -821,11 +821,9 @@ void Runtime::Load() {
   // Load tools libraries
   LoadTools();
 
-  // Initialize blit kernel object after tools is initialized to allow tools
-  // to overload blit kernel.
   for (core::Agent* agent : gpu_agents_) {
     const hsa_status_t stat =
-        reinterpret_cast<amd::GpuAgentInt*>(agent)->InitBlitKernel();
+        reinterpret_cast<amd::GpuAgentInt*>(agent)->PostToolsInit();
     assert(HSA_STATUS_SUCCESS == stat);
   }
 }