[OpenMP][libomptarget][AMDGPU] Enable active HSA wait state

[jplehr]

Summary

[OpenMP][libomptarget][AMDGPU] Enable active HSA wait state

Adds HSA timeout hint of 2 seconds to the AMDGPU nextgen-plugin to improve
performance of small kernels.
The HSA runtime may stay in HSA_WAIT_STATE_ACTIVE for up to the timeout
value before switching to HSA_WAIT_STATE_BLOCKED. This can improve
latency from which small kernels can benefit.
The value was determined via experimentation w/ different benchmarks.

The timeout value can be overriden using the environment variable
LIBOMPTARGET_AMDGPU_STREAM_BUSYWAIT with a value in microseconds.

[ye-luo]

Why separate KERNEL_BUSYWAIT and DATA_BUSYWAIT? I don't feel the need of two separate controls.
What are the units? Documentation?

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

521

When reading the implementation of wait(). I'm still wondering how error got handled when kernels finished (signal equals 0) with an error.

1245

Should this be private?

1575

Should they be private?

[kevinsala]

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

1245

Can we place this class data member just after the existing ones?

1575

Are these fields really needed? We can access the envars directly (which have the values cached in a variable), right?

[gregrodgers]

In D148808#4283842, @ye-luo wrote:

Why separate KERNEL_BUSYWAIT and DATA_BUSYWAIT? I don't feel the need of two separate controls.
What are the units? Documentation?

It is a matter of priorities. If you are doing async copies, you may not want them to every busy-wait because you have more important things to do on the CPU. However, in that same situation you may have high priority short kernels that you prefer to get the performance of busywait.

For long-term tuning, we may actually want more control over busy-wait timeouts. For now, we just separate the timeouts on data transfers and kernel completions.

[ye-luo]

In D148808#4286873, @gregrodgers wrote:

In D148808#4283842, @ye-luo wrote:

Why separate KERNEL_BUSYWAIT and DATA_BUSYWAIT? I don't feel the need of two separate controls.
What are the units? Documentation?

It is a matter of priorities. If you are doing async copies, you may not want them to every busy-wait because you have more important things to do on the CPU.

In the OpenMP context, when an OpenMP target without nowait reaches a signal wait. The user intention is busy/active waiting. There are no more important things to do on the core that OpenMP thread resides. If you were thinking of CPU cores being oversubscribed by threads, I would say one may gain or lose all depending on the application and there is no need to put such cases in high priority of discussion.
If nowait is added, events are used instead of signal wait. I think active waiting is always preferred.

[jplehr]

Thanks for all the comments, I think there is a bit of confusion due to my lack of documentation of the code.

1. Units of the variables: microseconds.
2. Wait state: This indicates whether the HSA runtime should actively wait. As far as I understand it, this means that it will likely not perform a context switch while waiting. This can improve the responsiveness to the signal value change. I did some basic profiles with babelstream using standard waiting and the timeout values we used in the old plugin. I can see improvements in the duration that the system stays within the hsa_signal_wait_scaqcuire. This suggests that this is indeed helping with latency, and therefore with the runtime "realizing" that a short running kernel has finished. Babelstream results do reproducibly improve.
3. As a result of 2., I think, we used different values for kernels and data movements, as the idea is that the data transfers may just be a little slower.

I will go ahead and update the patch to address all code-related comments.

[jplehr]

• rebase the patch
• add documentation
• address reviewer comments

[ye-luo]

From the meeting discussion.

1. default to active waiting rather than 0.
2. Consider a single LIBOMPTARGET_AMDGPU_STREAM_BUSYWAIT.

[jplehr]

One environment variable, defaulting to 2 seconds timeout for active waiting, rename, const

[ye-luo]

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

900

Could you move this right below line 891 mutex? Prefer not to have variables stuck between functions.

[ye-luo]

Updating the patch description is also needed.

[jplehr]

Address inline comment

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

struct AMDGPUSignalTy {
 
   /// Deinitialize the signal.
   Error deinit() {
     hsa_status_t Status = hsa_signal_destroy(Signal);
     return Plugin::check(Status, "Error in hsa_signal_destroy: %s");
   }
 
   /// Wait until the signal gets a zero value.
-  Error wait() const {
-    // TODO: Is it better to use busy waiting or blocking the thread?
+  Error wait(const uint64_t ActiveTimeout = 0) const {
+    if (ActiveTimeout) {
+      hsa_signal_value_t Got = 1;
+      Got = hsa_signal_wait_scacquire(Signal, HSA_SIGNAL_CONDITION_EQ, 0,
+                                      ActiveTimeout, HSA_WAIT_STATE_ACTIVE);
+      if (Got == 0)
+        return Plugin::success();
+    }
     while (hsa_signal_wait_scacquire(Signal, HSA_SIGNAL_CONDITION_EQ, 0,
                                      UINT64_MAX, HSA_WAIT_STATE_BLOCKED) != 0)
       ;
     return Plugin::success();
   }
 
   /// Load the value on the signal.
   hsa_signal_value_t load() const { return hsa_signal_load_scacquire(Signal); }
private:
 
   /// Return the current number of asychronous operations on the stream.
   uint32_t size() const { return NextSlot; }
 
   /// Return the last valid slot on the stream.
   uint32_t last() const { return size() - 1; }
 
   /// Consume one slot from the stream. Since the stream uses signals on demand
   /// and releases them once the slot is no longer used, the function requires
   /// an idle signal for the new consumed slot.
   std::pair<uint32_t, AMDGPUSignalTy *> consume(AMDGPUSignalTy *OutputSignal) {
     // Double the stream size if needed. Since we use std::deque, this operation
     // does not invalidate the already added slots.
     if (Slots.size() == NextSlot)
       Slots.resize(Slots.size() * 2);
 
     // Update the next available slot and the stream size.
if (InputSignal && InputSignal->load()) {
                                          &InputSignalRaw, OutputSignal->get());
     } else
       Status = hsa_amd_memory_async_copy(Dst, Agent, Inter, Agent, CopySize, 0,
                                          nullptr, OutputSignal->get());
 
     return Plugin::check(Status, "Error in hsa_amd_memory_async_copy: %s");
   }
 
+  uint64_t KernelBusyWaitTics; // initialized from AMDGPUDeviceTy
+
   /// Synchronize with the stream. The current thread waits until all operations
   /// are finalized and it performs the pending post actions (i.e., releasing
   /// intermediate buffers).
   Error synchronize() {
     std::lock_guard<std::mutex> Lock(Mutex);
 
     // No need to synchronize anything.
     if (size() == 0)
       return Plugin::success();
 
     // Wait until all previous operations on the stream have completed.
-    if (auto Err = Slots[last()].Signal->wait())
+    if (auto Err = Slots[last()].Signal->wait(KernelBusyWaitTics))
       return Err;
 
     // Reset the stream and perform all pending post actions.
     return complete();
   }
 
   /// Query the stream and complete pending post actions if operations finished.
   /// Return whether all the operations completed. This operation does not block
AMDGPUDeviceTy(int32_t DeviceId, int32_t NumDevices,
       : GenericDeviceTy(DeviceId, NumDevices, {0}), AMDGenericDeviceTy(),
         OMPX_NumQueues("LIBOMPTARGET_AMDGPU_NUM_HSA_QUEUES", 4),
         OMPX_QueueSize("LIBOMPTARGET_AMDGPU_HSA_QUEUE_SIZE", 512),
         OMPX_DefaultTeamsPerCU("LIBOMPTARGET_AMDGPU_TEAMS_PER_CU", 4),
         OMPX_MaxAsyncCopyBytes("LIBOMPTARGET_AMDGPU_MAX_ASYNC_COPY_BYTES",
                                1 * 1024 * 1024), // 1MB
         OMPX_InitialNumSignals("LIBOMPTARGET_AMDGPU_NUM_INITIAL_HSA_SIGNALS",
                                64),
+        OMPX_KernelBusyWait("LIBOMPTARGET_AMDGPU_KERNEL_BUSYWAIT", 0),
+        OMPX_DataBusyWait("LIBOMPTARGET_AMDGPU_DATA_BUSYWAIT", 0),
         AMDGPUStreamManager(*this), AMDGPUEventManager(*this),
         AMDGPUSignalManager(*this), Agent(Agent), HostDevice(HostDevice),
         Queues() {}
 
   ~AMDGPUDeviceTy() {}
 
+  uint64_t KernelBusyWaitTics;
+  uint64_t DataBusyWaitTics;
+
   /// Initialize the device, its resources and get its properties.
   Error initImpl(GenericPluginTy &Plugin) override {
     // First setup all the memory pools.
     if (auto Err = initMemoryPools())
       return Err;
 
     char GPUName[64];
     if (auto Err = getDeviceAttr(HSA_AGENT_INFO_NAME, GPUName))
Error initImpl(GenericPluginTy &Plugin) override {
     uint32_t MaxQueues;
     if (auto Err = getDeviceAttr(HSA_AGENT_INFO_QUEUES_MAX, MaxQueues))
       return Err;
 
     // Compute the number of queues and their size.
     const uint32_t NumQueues = std::min(OMPX_NumQueues.get(), MaxQueues);
     const uint32_t QueueSize = std::min(OMPX_QueueSize.get(), MaxQueueSize);
 
+    KernelBusyWaitTics = OMPX_KernelBusyWait;
+    DataBusyWaitTics = OMPX_DataBusyWait;
+
     // Construct and initialize each device queue.
     Queues = std::vector<AMDGPUQueueTy>(NumQueues);
     for (AMDGPUQueueTy &Queue : Queues)
       if (auto Err = Queue.init(Agent, QueueSize))
         return Err;
 
     // Initialize stream pool.
     if (auto Err = AMDGPUStreamManager.init(OMPX_InitialNumStreams))
Error deinitImpl() override {
     }
 
     // Invalidate agent reference.
     Agent = {0};
 
     return Plugin::success();
   }
 
+  uint64_t getDataBusyWaitTics() const { return DataBusyWaitTics; }
+  uint64_t getKernelBusyWaitTics() const { return KernelBusyWaitTics; }
+
   Expected<std::unique_ptr<MemoryBuffer>>
   doJITPostProcessing(std::unique_ptr<MemoryBuffer> MB) const override {
 
     // TODO: We should try to avoid materialization but there seems to be no
     // good linker interface w/o file i/o.
     SmallString<128> LinkerOutputFilePath;
     std::error_code EC = sys::fs::createTemporaryFile(
         "amdgpu-pre-link-jit", ".out", LinkerOutputFilePath);
if (Size >= OMPX_MaxAsyncCopyBytes) {
         return Err;
 
       Status = hsa_amd_memory_async_copy(TgtPtr, Agent, PinnedHstPtr, Agent,
                                          Size, 0, nullptr, Signal.get());
       if (auto Err =
               Plugin::check(Status, "Error in hsa_amd_memory_async_copy: %s"))
         return Err;
 
-      if (auto Err = Signal.wait())
+      if (auto Err = Signal.wait(getDataBusyWaitTics()))
         return Err;
 
       if (auto Err = Signal.deinit())
         return Err;
 
       Status = hsa_amd_memory_unlock(const_cast<void *>(HstPtr));
       return Plugin::check(Status, "Error in hsa_amd_memory_unlock: %s\n");
     }
if (Size >= OMPX_MaxAsyncCopyBytes) {
         return Err;
 
       Status = hsa_amd_memory_async_copy(PinnedHstPtr, Agent, TgtPtr, Agent,
                                          Size, 0, nullptr, Signal.get());
       if (auto Err =
               Plugin::check(Status, "Error in hsa_amd_memory_async_copy: %s"))
         return Err;
 
-      if (auto Err = Signal.wait())
+      if (auto Err = Signal.wait(getDataBusyWaitTics()))
         return Err;
 
       if (auto Err = Signal.deinit())
         return Err;
 
       Status = hsa_amd_memory_unlock(const_cast<void *>(HstPtr));
       return Plugin::check(Status, "Error in hsa_amd_memory_unlock: %s\n");
     }
private:
   UInt32Envar OMPX_MaxAsyncCopyBytes;
 
   /// Envar controlling the initial number of HSA signals per device. There is
   /// one manager of signals per device managing several pre-allocated signals.
   /// These signals are mainly used by AMDGPU streams. If needed, more signals
   /// will be created.
   UInt32Envar OMPX_InitialNumSignals;
 
+  /// Environment variables to set the time to wait in active state before
+  /// switching to blocked state. The default 0 goes directly to blocked state.
+  UInt32Envar OMPX_KernelBusyWait;
+  UInt32Envar OMPX_DataBusyWait;
+
   /// Stream manager for AMDGPU streams.
   AMDGPUStreamManagerTy AMDGPUStreamManager;
 
   /// Event manager for AMDGPU events.
   AMDGPUEventManagerTy AMDGPUEventManager;
 
   /// Signal manager for AMDGPU signals.
   AMDGPUSignalManagerTy AMDGPUSignalManager;
Error AMDGPUResourceRef<ResourceTy>::create(GenericDeviceTy &Device) {
 
   return Resource->init();
 }
 
 AMDGPUStreamTy::AMDGPUStreamTy(AMDGPUDeviceTy &Device)
     : Agent(Device.getAgent()), Queue(Device.getNextQueue()),
       SignalManager(Device.getSignalManager()),
       // Initialize the std::deque with some empty positions.
-      Slots(32), NextSlot(0), SyncCycle(0) {}
+      Slots(32), NextSlot(0), SyncCycle(0),
+      KernelBusyWaitTics(Device.getKernelBusyWaitTics()) {}
 
 /// Class implementing the AMDGPU-specific functionalities of the global
 /// handler.
 struct AMDGPUGlobalHandlerTy final : public GenericGlobalHandlerTy {
   /// Get the metadata of a global from the device. The name and size of the
   /// global is read from DeviceGlobal and the address of the global is written
   /// to DeviceGlobal.
   Error getGlobalMetadataFromDevice(GenericDeviceTy &Device,