brintos / llvm-project-archived public Read only

1//===----RTLs/amdgpu/src/rtl.cpp - Target RTLs Implementation ----- C++ -*-===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// RTL NextGen for AMDGPU machine10//11//===----------------------------------------------------------------------===//12 13#include <atomic>14#include <cassert>15#include <cstddef>16#include <cstdint>17#include <deque>18#include <functional>19#include <mutex>20#include <string>21#include <system_error>22#include <unistd.h>23#include <unordered_map>24 25#include "ErrorReporting.h"26#include "Shared/APITypes.h"27#include "Shared/Debug.h"28#include "Shared/Environment.h"29#include "Shared/RefCnt.h"30#include "Shared/Utils.h"31#include "Utils/ELF.h"32 33#include "GlobalHandler.h"34#include "OpenMP/OMPT/Callback.h"35#include "PluginInterface.h"36#include "UtilitiesRTL.h"37#include "omptarget.h"38 39#include "llvm/ADT/SmallString.h"40#include "llvm/ADT/SmallVector.h"41#include "llvm/ADT/StringRef.h"42#include "llvm/BinaryFormat/ELF.h"43#include "llvm/Frontend/OpenMP/OMPConstants.h"44#include "llvm/Frontend/OpenMP/OMPGridValues.h"45#include "llvm/Support/Error.h"46#include "llvm/Support/FileOutputBuffer.h"47#include "llvm/Support/FileSystem.h"48#include "llvm/Support/MemoryBuffer.h"49#include "llvm/Support/Program.h"50#include "llvm/Support/Signals.h"51#include "llvm/Support/raw_ostream.h"52 53#if !defined(__BYTE_ORDER__) || !defined(__ORDER_LITTLE_ENDIAN__) ||           \54    !defined(__ORDER_BIG_ENDIAN__)55#error "Missing preprocessor definitions for endianness detection."56#endif57 58// The HSA headers require these definitions.59#if defined(__BYTE_ORDER__) && (__BYTE_ORDER__ == __ORDER_LITTLE_ENDIAN__)60#define LITTLEENDIAN_CPU61#elif defined(__BYTE_ORDER__) && (__BYTE_ORDER__ == __ORDER_BIG_ENDIAN__)62#define BIGENDIAN_CPU63#endif64 65#if defined(__has_include)66#if __has_include("hsa.h")67#include "hsa.h"68#include "hsa_ext_amd.h"69#elif __has_include("hsa/hsa.h")70#include "hsa/hsa.h"71#include "hsa/hsa_ext_amd.h"72#endif73#else74#include "hsa/hsa.h"75#include "hsa/hsa_ext_amd.h"76#endif77 78using namespace error;79 80namespace llvm {81namespace omp {82namespace target {83namespace plugin {84 85/// Forward declarations for all specialized data structures.86struct AMDGPUKernelTy;87struct AMDGPUDeviceTy;88struct AMDGPUPluginTy;89struct AMDGPUStreamTy;90struct AMDGPUEventTy;91struct AMDGPUStreamManagerTy;92struct AMDGPUEventManagerTy;93struct AMDGPUDeviceImageTy;94struct AMDGPUMemoryManagerTy;95struct AMDGPUMemoryPoolTy;96 97namespace hsa_utils {98 99/// Iterate elements using an HSA iterate function. Do not use this function100/// directly but the specialized ones below instead.101template <typename ElemTy, typename IterFuncTy, typename CallbackTy>102static hsa_status_t iterate(IterFuncTy Func, CallbackTy Cb) {103  auto L = [](ElemTy Elem, void *Data) -> hsa_status_t {104    CallbackTy *Unwrapped = static_cast<CallbackTy *>(Data);105    return (*Unwrapped)(Elem);106  };107  return Func(L, static_cast<void *>(&Cb));108}109 110/// Iterate elements using an HSA iterate function passing a parameter. Do not111/// use this function directly but the specialized ones below instead.112template <typename ElemTy, typename IterFuncTy, typename IterFuncArgTy,113          typename CallbackTy>114static hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg,115                            CallbackTy Cb) {116  auto L = [](ElemTy Elem, void *Data) -> hsa_status_t {117    CallbackTy *Unwrapped = static_cast<CallbackTy *>(Data);118    return (*Unwrapped)(Elem);119  };120  return Func(FuncArg, L, static_cast<void *>(&Cb));121}122 123/// Iterate elements using an HSA iterate function passing a parameter. Do not124/// use this function directly but the specialized ones below instead.125template <typename Elem1Ty, typename Elem2Ty, typename IterFuncTy,126          typename IterFuncArgTy, typename CallbackTy>127static hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg,128                            CallbackTy Cb) {129  auto L = [](Elem1Ty Elem1, Elem2Ty Elem2, void *Data) -> hsa_status_t {130    CallbackTy *Unwrapped = static_cast<CallbackTy *>(Data);131    return (*Unwrapped)(Elem1, Elem2);132  };133  return Func(FuncArg, L, static_cast<void *>(&Cb));134}135 136/// Iterate agents.137template <typename CallbackTy> static Error iterateAgents(CallbackTy Callback) {138  hsa_status_t Status = iterate<hsa_agent_t>(hsa_iterate_agents, Callback);139  return Plugin::check(Status, "error in hsa_iterate_agents: %s");140}141 142/// Iterate ISAs of an agent.143template <typename CallbackTy>144static Error iterateAgentISAs(hsa_agent_t Agent, CallbackTy Cb) {145  hsa_status_t Status = iterate<hsa_isa_t>(hsa_agent_iterate_isas, Agent, Cb);146  return Plugin::check(Status, "error in hsa_agent_iterate_isas: %s");147}148 149/// Iterate memory pools of an agent.150template <typename CallbackTy>151static Error iterateAgentMemoryPools(hsa_agent_t Agent, CallbackTy Cb) {152  hsa_status_t Status = iterate<hsa_amd_memory_pool_t>(153      hsa_amd_agent_iterate_memory_pools, Agent, Cb);154  return Plugin::check(Status,155                       "error in hsa_amd_agent_iterate_memory_pools: %s");156}157 158/// Dispatches an asynchronous memory copy.159/// Enables different SDMA engines for the dispatch in a round-robin fashion.160static Error asyncMemCopy(bool UseMultipleSdmaEngines, void *Dst,161                          hsa_agent_t DstAgent, const void *Src,162                          hsa_agent_t SrcAgent, size_t Size,163                          uint32_t NumDepSignals,164                          const hsa_signal_t *DepSignals,165                          hsa_signal_t CompletionSignal) {166  if (!UseMultipleSdmaEngines) {167    hsa_status_t S =168        hsa_amd_memory_async_copy(Dst, DstAgent, Src, SrcAgent, Size,169                                  NumDepSignals, DepSignals, CompletionSignal);170    return Plugin::check(S, "error in hsa_amd_memory_async_copy: %s");171  }172 173// This solution is probably not the best174#if !(HSA_AMD_INTERFACE_VERSION_MAJOR >= 1 &&                                  \175      HSA_AMD_INTERFACE_VERSION_MINOR >= 2)176  return Plugin::error(ErrorCode::UNSUPPORTED,177                       "async copy on selected SDMA requires ROCm 5.7");178#else179  static std::atomic<int> SdmaEngine{1};180 181  // This atomics solution is probably not the best, but should be sufficient182  // for now.183  // In a worst case scenario, in which threads read the same value, they will184  // dispatch to the same SDMA engine. This may result in sub-optimal185  // performance. However, I think the possibility to be fairly low.186  int LocalSdmaEngine = SdmaEngine.load(std::memory_order_acquire);187  // This call is only avail in ROCm >= 5.7188  hsa_status_t S = hsa_amd_memory_async_copy_on_engine(189      Dst, DstAgent, Src, SrcAgent, Size, NumDepSignals, DepSignals,190      CompletionSignal, (hsa_amd_sdma_engine_id_t)LocalSdmaEngine,191      /*force_copy_on_sdma=*/true);192  // Increment to use one of two SDMA engines: 0x1, 0x2193  LocalSdmaEngine = (LocalSdmaEngine << 1) % 3;194  SdmaEngine.store(LocalSdmaEngine, std::memory_order_relaxed);195 196  return Plugin::check(S, "error in hsa_amd_memory_async_copy_on_engine: %s");197#endif198}199 200static Error getTargetTripleAndFeatures(hsa_agent_t Agent,201                                        SmallVector<SmallString<32>> &Targets) {202  auto Err = hsa_utils::iterateAgentISAs(Agent, [&](hsa_isa_t ISA) {203    uint32_t Length;204    hsa_status_t Status;205    Status = hsa_isa_get_info_alt(ISA, HSA_ISA_INFO_NAME_LENGTH, &Length);206    if (Status != HSA_STATUS_SUCCESS)207      return Status;208 209    llvm::SmallVector<char> ISAName(Length);210    Status = hsa_isa_get_info_alt(ISA, HSA_ISA_INFO_NAME, ISAName.begin());211    if (Status != HSA_STATUS_SUCCESS)212      return Status;213 214    llvm::StringRef TripleTarget(ISAName.begin(), Length);215    if (TripleTarget.consume_front("amdgcn-amd-amdhsa")) {216      auto Target = TripleTarget.ltrim('-').rtrim('\0');217      Targets.push_back(Target);218    }219    return HSA_STATUS_SUCCESS;220  });221  return Err;222}223} // namespace hsa_utils224 225/// Utility class representing generic resource references to AMDGPU resources.226template <typename ResourceTy>227struct AMDGPUResourceRef : public GenericDeviceResourceRef {228  /// The underlying handle type for resources.229  using HandleTy = ResourceTy *;230 231  /// Create an empty reference to an invalid resource.232  AMDGPUResourceRef() : Resource(nullptr) {}233 234  /// Create a reference to an existing resource.235  AMDGPUResourceRef(HandleTy Resource) : Resource(Resource) {}236 237  virtual ~AMDGPUResourceRef() {}238 239  /// Create a new resource and save the reference. The reference must be empty240  /// before calling to this function.241  Error create(GenericDeviceTy &Device) override;242 243  /// Destroy the referenced resource and invalidate the reference. The244  /// reference must be to a valid resource before calling to this function.245  Error destroy(GenericDeviceTy &Device) override {246    if (!Resource)247      return Plugin::error(ErrorCode::INVALID_ARGUMENT,248                           "destroying an invalid resource");249 250    if (auto Err = Resource->deinit())251      return Err;252 253    delete Resource;254 255    Resource = nullptr;256    return Plugin::success();257  }258 259  /// Get the underlying resource handle.260  operator HandleTy() const { return Resource; }261 262private:263  /// The handle to the actual resource.264  HandleTy Resource;265};266 267/// Class holding an HSA memory pool.268struct AMDGPUMemoryPoolTy {269  /// Create a memory pool from an HSA memory pool.270  AMDGPUMemoryPoolTy(hsa_amd_memory_pool_t MemoryPool)271      : MemoryPool(MemoryPool), GlobalFlags(0) {}272 273  /// Initialize the memory pool retrieving its properties.274  Error init() {275    if (auto Err = getAttr(HSA_AMD_MEMORY_POOL_INFO_SEGMENT, Segment))276      return Err;277 278    if (auto Err = getAttr(HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, GlobalFlags))279      return Err;280 281    return Plugin::success();282  }283 284  /// Getter of the HSA memory pool.285  hsa_amd_memory_pool_t get() const { return MemoryPool; }286 287  /// Indicate the segment which belongs to.288  bool isGlobal() const { return (Segment == HSA_AMD_SEGMENT_GLOBAL); }289  bool isReadOnly() const { return (Segment == HSA_AMD_SEGMENT_READONLY); }290  bool isPrivate() const { return (Segment == HSA_AMD_SEGMENT_PRIVATE); }291  bool isGroup() const { return (Segment == HSA_AMD_SEGMENT_GROUP); }292 293  /// Indicate if it is fine-grained memory. Valid only for global.294  bool isFineGrained() const {295    assert(isGlobal() && "Not global memory");296    return (GlobalFlags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED);297  }298 299  /// Indicate if it is coarse-grained memory. Valid only for global.300  bool isCoarseGrained() const {301    assert(isGlobal() && "Not global memory");302    return (GlobalFlags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED);303  }304 305  /// Indicate if it supports storing kernel arguments. Valid only for global.306  bool supportsKernelArgs() const {307    assert(isGlobal() && "Not global memory");308    return (GlobalFlags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT);309  }310 311  /// Allocate memory on the memory pool.312  Error allocate(size_t Size, void **PtrStorage) {313    hsa_status_t Status =314        hsa_amd_memory_pool_allocate(MemoryPool, Size, 0, PtrStorage);315    return Plugin::check(Status, "error in hsa_amd_memory_pool_allocate: %s");316  }317 318  /// Return memory to the memory pool.319  Error deallocate(void *Ptr) {320    hsa_status_t Status = hsa_amd_memory_pool_free(Ptr);321    return Plugin::check(Status, "error in hsa_amd_memory_pool_free: %s");322  }323 324  /// Returns if the \p Agent can access the memory pool.325  bool canAccess(hsa_agent_t Agent) {326    hsa_amd_memory_pool_access_t Access;327    if (hsa_amd_agent_memory_pool_get_info(328            Agent, MemoryPool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &Access))329      return false;330    return Access != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED;331  }332 333  /// Allow the device to access a specific allocation.334  Error enableAccess(void *Ptr, int64_t Size,335                     const llvm::SmallVector<hsa_agent_t> &Agents) const {336#ifdef OMPTARGET_DEBUG337    for (hsa_agent_t Agent : Agents) {338      hsa_amd_memory_pool_access_t Access;339      if (auto Err =340              getAttr(Agent, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, Access))341        return Err;342 343      // The agent is not allowed to access the memory pool in any case. Do not344      // continue because otherwise it result in undefined behavior.345      if (Access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED)346        return Plugin::error(ErrorCode::INVALID_VALUE,347                             "an agent is not allowed to access a memory pool");348    }349#endif350 351    // We can access but it is disabled by default. Enable the access then.352    hsa_status_t Status =353        hsa_amd_agents_allow_access(Agents.size(), Agents.data(), nullptr, Ptr);354    return Plugin::check(Status, "error in hsa_amd_agents_allow_access: %s");355  }356 357  /// Get attribute from the memory pool.358  template <typename Ty>359  Error getAttr(hsa_amd_memory_pool_info_t Kind, Ty &Value) const {360    hsa_status_t Status;361    Status = hsa_amd_memory_pool_get_info(MemoryPool, Kind, &Value);362    return Plugin::check(Status, "error in hsa_amd_memory_pool_get_info: %s");363  }364 365  template <typename Ty>366  hsa_status_t getAttrRaw(hsa_amd_memory_pool_info_t Kind, Ty &Value) const {367    return hsa_amd_memory_pool_get_info(MemoryPool, Kind, &Value);368  }369 370  /// Get attribute from the memory pool relating to an agent.371  template <typename Ty>372  Error getAttr(hsa_agent_t Agent, hsa_amd_agent_memory_pool_info_t Kind,373                Ty &Value) const {374    hsa_status_t Status;375    Status =376        hsa_amd_agent_memory_pool_get_info(Agent, MemoryPool, Kind, &Value);377    return Plugin::check(Status,378                         "error in hsa_amd_agent_memory_pool_get_info: %s");379  }380 381private:382  /// The HSA memory pool.383  hsa_amd_memory_pool_t MemoryPool;384 385  /// The segment where the memory pool belongs to.386  hsa_amd_segment_t Segment;387 388  /// The global flags of memory pool. Only valid if the memory pool belongs to389  /// the global segment.390  uint32_t GlobalFlags;391};392 393/// Class that implements a memory manager that gets memory from a specific394/// memory pool.395struct AMDGPUMemoryManagerTy : public DeviceAllocatorTy {396 397  /// Create an empty memory manager.398  AMDGPUMemoryManagerTy(AMDGPUPluginTy &Plugin)399      : Plugin(Plugin), MemoryPool(nullptr), MemoryManager(nullptr) {}400 401  /// Initialize the memory manager from a memory pool.402  Error init(AMDGPUMemoryPoolTy &MemoryPool) {403    const uint32_t Threshold = 1 << 30;404    this->MemoryManager = new MemoryManagerTy(*this, Threshold);405    this->MemoryPool = &MemoryPool;406    return Plugin::success();407  }408 409  /// Deinitialize the memory manager and free its allocations.410  Error deinit() {411    assert(MemoryManager && "Invalid memory manager");412 413    // Delete and invalidate the memory manager. At this point, the memory414    // manager will deallocate all its allocations.415    delete MemoryManager;416    MemoryManager = nullptr;417 418    return Plugin::success();419  }420 421  /// Reuse or allocate memory through the memory manager.422  Error allocate(size_t Size, void **PtrStorage) {423    assert(MemoryManager && "Invalid memory manager");424    assert(PtrStorage && "Invalid pointer storage");425 426    auto PtrStorageOrErr = MemoryManager->allocate(Size, nullptr);427    if (!PtrStorageOrErr)428      return PtrStorageOrErr.takeError();429 430    *PtrStorage = *PtrStorageOrErr;431    if (Size && *PtrStorage == nullptr)432      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,433                           "failure to allocate from AMDGPU memory manager");434 435    return Plugin::success();436  }437 438  /// Release an allocation to be reused.439  Error deallocate(void *Ptr) {440    if (MemoryManager->free(Ptr))441      return Plugin::error(ErrorCode::UNKNOWN,442                           "failure to deallocate from AMDGPU memory manager");443 444    return Plugin::success();445  }446 447private:448  /// Allocation callback that will be called once the memory manager does not449  /// have more previously allocated buffers.450  Expected<void *> allocate(size_t Size, void *HstPtr,451                            TargetAllocTy Kind) override;452 453  /// Deallocation callback that will be called by the memory manager.454  Error free(void *TgtPtr, TargetAllocTy Kind) override {455    return MemoryPool->deallocate(TgtPtr);456  }457 458  /// The underlying plugin that owns this memory manager.459  AMDGPUPluginTy &Plugin;460 461  /// The memory pool used to allocate memory.462  AMDGPUMemoryPoolTy *MemoryPool;463 464  /// Reference to the actual memory manager.465  MemoryManagerTy *MemoryManager;466};467 468/// Class implementing the AMDGPU device images' properties.469struct AMDGPUDeviceImageTy : public DeviceImageTy {470  /// Create the AMDGPU image with the id and the target image pointer.471  AMDGPUDeviceImageTy(int32_t ImageId, GenericDeviceTy &Device,472                      std::unique_ptr<MemoryBuffer> &&TgtImage)473      : DeviceImageTy(ImageId, Device, std::move(TgtImage)) {}474 475  /// Prepare and load the executable corresponding to the image.476  Error loadExecutable(const AMDGPUDeviceTy &Device);477 478  /// Unload the executable.479  Error unloadExecutable() {480    hsa_status_t Status = hsa_executable_destroy(Executable);481    return Plugin::check(Status, "error in hsa_executable_destroy: %s");482  }483 484  /// Get the executable.485  hsa_executable_t getExecutable() const { return Executable; }486 487  /// Get to Code Object Version of the ELF488  uint16_t getELFABIVersion() const { return ELFABIVersion; }489 490  /// Find an HSA device symbol by its name on the executable.491  Expected<hsa_executable_symbol_t>492  findDeviceSymbol(GenericDeviceTy &Device, StringRef SymbolName) const;493 494  /// Get additional info for kernel, e.g., register spill counts495  std::optional<offloading::amdgpu::AMDGPUKernelMetaData>496  getKernelInfo(StringRef Identifier) const {497    auto It = KernelInfoMap.find(Identifier);498 499    if (It == KernelInfoMap.end())500      return {};501 502    return It->second;503  }504 505private:506  /// The executable loaded on the agent.507  hsa_executable_t Executable;508  StringMap<offloading::amdgpu::AMDGPUKernelMetaData> KernelInfoMap;509  uint16_t ELFABIVersion;510};511 512/// Class implementing the AMDGPU kernel functionalities which derives from the513/// generic kernel class.514struct AMDGPUKernelTy : public GenericKernelTy {515  /// Create an AMDGPU kernel with a name and an execution mode.516  AMDGPUKernelTy(const char *Name) : GenericKernelTy(Name) {}517 518  /// Initialize the AMDGPU kernel.519  Error initImpl(GenericDeviceTy &Device, DeviceImageTy &Image) override {520    AMDGPUDeviceImageTy &AMDImage = static_cast<AMDGPUDeviceImageTy &>(Image);521 522    // Kernel symbols have a ".kd" suffix.523    std::string KernelName(getName());524    KernelName += ".kd";525 526    // Find the symbol on the device executable.527    auto SymbolOrErr = AMDImage.findDeviceSymbol(Device, KernelName);528    if (!SymbolOrErr)529      return SymbolOrErr.takeError();530 531    hsa_executable_symbol_t Symbol = *SymbolOrErr;532    hsa_symbol_kind_t SymbolType;533    hsa_status_t Status;534 535    // Retrieve different properties of the kernel symbol.536    std::pair<hsa_executable_symbol_info_t, void *> RequiredInfos[] = {537        {HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &SymbolType},538        {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &KernelObject},539        {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &ArgsSize},540        {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &GroupSize},541        {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK, &DynamicStack},542        {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &PrivateSize}};543 544    for (auto &Info : RequiredInfos) {545      Status = hsa_executable_symbol_get_info(Symbol, Info.first, Info.second);546      if (auto Err = Plugin::check(547              Status, "error in hsa_executable_symbol_get_info: %s"))548        return Err;549    }550 551    // Make sure it is a kernel symbol.552    if (SymbolType != HSA_SYMBOL_KIND_KERNEL)553      return Plugin::error(ErrorCode::INVALID_BINARY,554                           "symbol %s is not a kernel function");555 556    // TODO: Read the kernel descriptor for the max threads per block. May be557    // read from the image.558 559    ImplicitArgsSize =560        hsa_utils::getImplicitArgsSize(AMDImage.getELFABIVersion());561    DP("ELFABIVersion: %d\n", AMDImage.getELFABIVersion());562 563    // Get additional kernel info read from image564    KernelInfo = AMDImage.getKernelInfo(getName());565    if (!KernelInfo.has_value())566      INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device.getDeviceId(),567           "Could not read extra information for kernel %s.", getName());568 569    return Plugin::success();570  }571 572  /// Launch the AMDGPU kernel function.573  Error launchImpl(GenericDeviceTy &GenericDevice, uint32_t NumThreads[3],574                   uint32_t NumBlocks[3], KernelArgsTy &KernelArgs,575                   KernelLaunchParamsTy LaunchParams,576                   AsyncInfoWrapperTy &AsyncInfoWrapper) const override;577 578  /// Return maximum block size for maximum occupancy579  ///580  /// TODO: This needs to be implemented for amdgpu581  Expected<uint64_t> maxGroupSize(GenericDeviceTy &GenericDevice,582                                  uint64_t DynamicMemSize) const override {583    return Plugin::error(584        ErrorCode::UNSUPPORTED,585        "occupancy calculations for AMDGPU are not yet implemented");586  }587 588  /// Print more elaborate kernel launch info for AMDGPU589  Error printLaunchInfoDetails(GenericDeviceTy &GenericDevice,590                               KernelArgsTy &KernelArgs, uint32_t NumThreads[3],591                               uint32_t NumBlocks[3]) const override;592 593  /// Get group and private segment kernel size.594  uint32_t getGroupSize() const { return GroupSize; }595  uint32_t getPrivateSize() const { return PrivateSize; }596 597  /// Get the HSA kernel object representing the kernel function.598  uint64_t getKernelObject() const { return KernelObject; }599 600  /// Get the size of implicitargs based on the code object version.601  uint32_t getImplicitArgsSize() const { return ImplicitArgsSize; }602 603  /// Indicates whether or not we need to set up our own private segment size.604  bool usesDynamicStack() const { return DynamicStack; }605 606private:607  /// The kernel object to execute.608  uint64_t KernelObject;609 610  /// The args, group and private segments sizes required by a kernel instance.611  uint32_t ArgsSize;612  uint32_t GroupSize;613  uint32_t PrivateSize;614  bool DynamicStack;615 616  /// The size of implicit kernel arguments.617  uint32_t ImplicitArgsSize;618 619  /// Additional Info for the AMD GPU Kernel620  std::optional<offloading::amdgpu::AMDGPUKernelMetaData> KernelInfo;621};622 623/// Class representing an HSA signal. Signals are used to define dependencies624/// between asynchronous operations: kernel launches and memory transfers.625struct AMDGPUSignalTy {626  /// Create an empty signal.627  AMDGPUSignalTy() : HSASignal({0}), UseCount() {}628  AMDGPUSignalTy(AMDGPUDeviceTy &Device) : HSASignal({0}), UseCount() {}629 630  /// Initialize the signal with an initial value.631  Error init(uint32_t InitialValue = 1) {632    hsa_status_t Status =633        hsa_amd_signal_create(InitialValue, 0, nullptr, 0, &HSASignal);634    return Plugin::check(Status, "error in hsa_signal_create: %s");635  }636 637  /// Deinitialize the signal.638  Error deinit() {639    hsa_status_t Status = hsa_signal_destroy(HSASignal);640    return Plugin::check(Status, "error in hsa_signal_destroy: %s");641  }642 643  /// Wait until the signal gets a zero value.644  Error wait(const uint64_t ActiveTimeout = 0,645             GenericDeviceTy *Device = nullptr) const {646    if (ActiveTimeout) {647      hsa_signal_value_t Got = 1;648      Got = hsa_signal_wait_scacquire(HSASignal, HSA_SIGNAL_CONDITION_EQ, 0,649                                      ActiveTimeout, HSA_WAIT_STATE_ACTIVE);650      if (Got == 0)651        return Plugin::success();652    }653 654    // If there is an RPC device attached to this stream we run it as a server.655    uint64_t Timeout = UINT64_MAX;656    auto WaitState = HSA_WAIT_STATE_BLOCKED;657    while (hsa_signal_wait_scacquire(HSASignal, HSA_SIGNAL_CONDITION_EQ, 0,658                                     Timeout, WaitState) != 0)659      ;660    return Plugin::success();661  }662 663  /// Load the value on the signal.664  hsa_signal_value_t load() const {665    return hsa_signal_load_scacquire(HSASignal);666  }667 668  /// Signal decrementing by one.669  void signal() {670    assert(load() > 0 && "Invalid signal value");671    hsa_signal_subtract_screlease(HSASignal, 1);672  }673 674  /// Reset the signal value before reusing the signal. Do not call this675  /// function if the signal is being currently used by any watcher, such as a676  /// plugin thread or the HSA runtime.677  void reset() { hsa_signal_store_screlease(HSASignal, 1); }678 679  /// Increase the number of concurrent uses.680  void increaseUseCount() { UseCount.increase(); }681 682  /// Decrease the number of concurrent uses and return whether was the last.683  bool decreaseUseCount() { return UseCount.decrease(); }684 685  hsa_signal_t get() const { return HSASignal; }686 687private:688  /// The underlying HSA signal.689  hsa_signal_t HSASignal;690 691  /// Reference counter for tracking the concurrent use count. This is mainly692  /// used for knowing how many streams are using the signal.693  RefCountTy<> UseCount;694};695 696/// Classes for holding AMDGPU signals and managing signals.697using AMDGPUSignalRef = AMDGPUResourceRef<AMDGPUSignalTy>;698using AMDGPUSignalManagerTy = GenericDeviceResourceManagerTy<AMDGPUSignalRef>;699 700/// Class holding an HSA queue to submit kernel and barrier packets.701struct AMDGPUQueueTy {702  /// Create an empty queue.703  AMDGPUQueueTy() : Queue(nullptr), Mutex(), NumUsers(0) {}704 705  /// Lazily initialize a new queue belonging to a specific agent.706  Error init(GenericDeviceTy &Device, hsa_agent_t Agent, int32_t QueueSize) {707    if (Queue)708      return Plugin::success();709    hsa_status_t Status =710        hsa_queue_create(Agent, QueueSize, HSA_QUEUE_TYPE_MULTI, callbackError,711                         &Device, UINT32_MAX, UINT32_MAX, &Queue);712    return Plugin::check(Status, "error in hsa_queue_create: %s");713  }714 715  /// Deinitialize the queue and destroy its resources.716  Error deinit() {717    std::lock_guard<std::mutex> Lock(Mutex);718    if (!Queue)719      return Plugin::success();720    hsa_status_t Status = hsa_queue_destroy(Queue);721    return Plugin::check(Status, "error in hsa_queue_destroy: %s");722  }723 724  /// Returns the number of streams, this queue is currently assigned to.725  bool getUserCount() const { return NumUsers; }726 727  /// Returns if the underlying HSA queue is initialized.728  bool isInitialized() { return Queue != nullptr; }729 730  /// Decrement user count of the queue object.731  void removeUser() { --NumUsers; }732 733  /// Increase user count of the queue object.734  void addUser() { ++NumUsers; }735 736  /// Push a kernel launch to the queue. The kernel launch requires an output737  /// signal and can define an optional input signal (nullptr if none).738  Error pushKernelLaunch(const AMDGPUKernelTy &Kernel, void *KernelArgs,739                         uint32_t NumThreads[3], uint32_t NumBlocks[3],740                         uint32_t GroupSize, uint64_t StackSize,741                         AMDGPUSignalTy *OutputSignal,742                         AMDGPUSignalTy *InputSignal) {743    assert(OutputSignal && "Invalid kernel output signal");744 745    // Lock the queue during the packet publishing process. Notice this blocks746    // the addition of other packets to the queue. The following piece of code747    // should be lightweight; do not block the thread, allocate memory, etc.748    std::lock_guard<std::mutex> Lock(Mutex);749    assert(Queue && "Interacted with a non-initialized queue!");750 751    // Add a barrier packet before the kernel packet in case there is a pending752    // preceding operation. The barrier packet will delay the processing of753    // subsequent queue's packets until the barrier input signal are satisfied.754    // No need output signal needed because the dependency is already guaranteed755    // by the queue barrier itself.756    if (InputSignal && InputSignal->load())757      if (auto Err = pushBarrierImpl(nullptr, InputSignal))758        return Err;759 760    // Now prepare the kernel packet.761    uint64_t PacketId;762    hsa_kernel_dispatch_packet_t *Packet = acquirePacket(PacketId);763    assert(Packet && "Invalid packet");764 765    // The first 32 bits of the packet are written after the other fields766    uint16_t Dims = NumBlocks[2] * NumThreads[2] > 1767                        ? 3768                        : 1 + (NumBlocks[1] * NumThreads[1] != 1);769    uint16_t Setup = UINT16_C(Dims)770                     << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;771    Packet->workgroup_size_x = NumThreads[0];772    Packet->workgroup_size_y = NumThreads[1];773    Packet->workgroup_size_z = NumThreads[2];774    Packet->reserved0 = 0;775    Packet->grid_size_x = NumBlocks[0] * NumThreads[0];776    Packet->grid_size_y = NumBlocks[1] * NumThreads[1];777    Packet->grid_size_z = NumBlocks[2] * NumThreads[2];778    Packet->private_segment_size =779        Kernel.usesDynamicStack() ? StackSize : Kernel.getPrivateSize();780    Packet->group_segment_size = GroupSize;781    Packet->kernel_object = Kernel.getKernelObject();782    Packet->kernarg_address = KernelArgs;783    Packet->reserved2 = 0;784    Packet->completion_signal = OutputSignal->get();785 786    // Publish the packet. Do not modify the packet after this point.787    publishKernelPacket(PacketId, Setup, Packet);788 789    return Plugin::success();790  }791 792  /// Push a barrier packet that will wait up to two input signals. All signals793  /// are optional (nullptr if none).794  Error pushBarrier(AMDGPUSignalTy *OutputSignal,795                    const AMDGPUSignalTy *InputSignal1,796                    const AMDGPUSignalTy *InputSignal2) {797    // Lock the queue during the packet publishing process.798    std::lock_guard<std::mutex> Lock(Mutex);799    assert(Queue && "Interacted with a non-initialized queue!");800 801    // Push the barrier with the lock acquired.802    return pushBarrierImpl(OutputSignal, InputSignal1, InputSignal2);803  }804 805private:806  /// Push a barrier packet that will wait up to two input signals. Assumes the807  /// the queue lock is acquired.808  Error pushBarrierImpl(AMDGPUSignalTy *OutputSignal,809                        const AMDGPUSignalTy *InputSignal1,810                        const AMDGPUSignalTy *InputSignal2 = nullptr) {811    // Add a queue barrier waiting on both the other stream's operation and the812    // last operation on the current stream (if any).813    uint64_t PacketId;814    hsa_barrier_and_packet_t *Packet =815        (hsa_barrier_and_packet_t *)acquirePacket(PacketId);816    assert(Packet && "Invalid packet");817 818    Packet->reserved0 = 0;819    Packet->reserved1 = 0;820    Packet->dep_signal[0] = {0};821    Packet->dep_signal[1] = {0};822    Packet->dep_signal[2] = {0};823    Packet->dep_signal[3] = {0};824    Packet->dep_signal[4] = {0};825    Packet->reserved2 = 0;826    Packet->completion_signal = {0};827 828    // Set input and output dependencies if needed.829    if (OutputSignal)830      Packet->completion_signal = OutputSignal->get();831    if (InputSignal1)832      Packet->dep_signal[0] = InputSignal1->get();833    if (InputSignal2)834      Packet->dep_signal[1] = InputSignal2->get();835 836    // Publish the packet. Do not modify the packet after this point.837    publishBarrierPacket(PacketId, Packet);838 839    return Plugin::success();840  }841 842  /// Acquire a packet from the queue. This call may block the thread if there843  /// is no space in the underlying HSA queue. It may need to wait until the HSA844  /// runtime processes some packets. Assumes the queue lock is acquired.845  hsa_kernel_dispatch_packet_t *acquirePacket(uint64_t &PacketId) {846    // Increase the queue index with relaxed memory order. Notice this will need847    // another subsequent atomic operation with acquire order.848    PacketId = hsa_queue_add_write_index_relaxed(Queue, 1);849 850    // Wait for the package to be available. Notice the atomic operation uses851    // the acquire memory order.852    while (PacketId - hsa_queue_load_read_index_scacquire(Queue) >= Queue->size)853      ;854 855    // Return the packet reference.856    const uint32_t Mask = Queue->size - 1; // The size is a power of 2.857    return (hsa_kernel_dispatch_packet_t *)Queue->base_address +858           (PacketId & Mask);859  }860 861  /// Publish the kernel packet so that the HSA runtime can start processing862  /// the kernel launch. Do not modify the packet once this function is called.863  /// Assumes the queue lock is acquired.864  void publishKernelPacket(uint64_t PacketId, uint16_t Setup,865                           hsa_kernel_dispatch_packet_t *Packet) {866    uint32_t *PacketPtr = reinterpret_cast<uint32_t *>(Packet);867 868    uint16_t Header = HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;869    Header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;870    Header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;871 872    // Publish the packet. Do not modify the package after this point.873    uint32_t HeaderWord = Header | (Setup << 16u);874    __atomic_store_n(PacketPtr, HeaderWord, __ATOMIC_RELEASE);875 876    // Signal the doorbell about the published packet.877    hsa_signal_store_relaxed(Queue->doorbell_signal, PacketId);878  }879 880  /// Publish the barrier packet so that the HSA runtime can start processing881  /// the barrier. Next packets in the queue will not be processed until all882  /// barrier dependencies (signals) are satisfied. Assumes the queue is locked883  void publishBarrierPacket(uint64_t PacketId,884                            hsa_barrier_and_packet_t *Packet) {885    uint32_t *PacketPtr = reinterpret_cast<uint32_t *>(Packet);886    uint16_t Setup = 0;887    uint16_t Header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;888    Header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;889    Header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;890 891    // Publish the packet. Do not modify the package after this point.892    uint32_t HeaderWord = Header | (Setup << 16u);893    __atomic_store_n(PacketPtr, HeaderWord, __ATOMIC_RELEASE);894 895    // Signal the doorbell about the published packet.896    hsa_signal_store_relaxed(Queue->doorbell_signal, PacketId);897  }898 899  /// Callback that will be called when an error is detected on the HSA queue.900  static void callbackError(hsa_status_t Status, hsa_queue_t *Source,901                            void *Data);902 903  /// The HSA queue.904  hsa_queue_t *Queue;905 906  /// Mutex to protect the acquiring and publishing of packets. For the moment,907  /// we need this mutex to prevent publishing packets that are not ready to be908  /// published in a multi-thread scenario. Without a queue lock, a thread T1909  /// could acquire packet P and thread T2 acquire packet P+1. Thread T2 could910  /// publish its packet P+1 (signaling the queue's doorbell) before packet P911  /// from T1 is ready to be processed. That scenario should be invalid. Thus,912  /// we use the following mutex to make packet acquiring and publishing atomic.913  /// TODO: There are other more advanced approaches to avoid this mutex using914  /// atomic operations. We can further investigate it if this is a bottleneck.915  std::mutex Mutex;916 917  /// The number of streams, this queue is currently assigned to. A queue is918  /// considered idle when this is zero, otherwise: busy.919  uint32_t NumUsers;920};921 922/// Struct that implements a stream of asynchronous operations for AMDGPU923/// devices. This class relies on signals to implement streams and define the924/// dependencies between asynchronous operations.925struct AMDGPUStreamTy {926public:927  /// Function pointer type for `pushHostCallback`928  using HostFnType = void (*)(void *);929 930private:931  /// Utility struct holding arguments for async H2H memory copies.932  struct MemcpyArgsTy {933    void *Dst;934    const void *Src;935    size_t Size;936    size_t NumTimes;937  };938 939  /// Utility struct holding arguments for freeing buffers to memory managers.940  struct ReleaseBufferArgsTy {941    void *Buffer;942    AMDGPUMemoryManagerTy *MemoryManager;943  };944 945  /// Utility struct holding arguments for releasing signals to signal managers.946  struct ReleaseSignalArgsTy {947    AMDGPUSignalTy *Signal;948    AMDGPUSignalManagerTy *SignalManager;949  };950 951  using AMDGPUStreamCallbackTy = Error(void *Data);952 953  /// The stream is composed of N stream's slots. The struct below represents954  /// the fields of each slot. Each slot has a signal and an optional action955  /// function. When appending an HSA asynchronous operation to the stream, one956  /// slot is consumed and used to store the operation's information. The957  /// operation's output signal is set to the consumed slot's signal. If there958  /// is a previous asynchronous operation on the previous slot, the HSA async959  /// operation's input signal is set to the signal of the previous slot. This960  /// way, we obtain a chain of dependent async operations. The action is a961  /// function that will be executed eventually after the operation is962  /// completed, e.g., for releasing a buffer.963  struct StreamSlotTy {964    /// The output signal of the stream operation. May be used by the subsequent965    /// operation as input signal.966    AMDGPUSignalTy *Signal;967 968    /// The actions that must be performed after the operation's completion. Set969    /// to nullptr when there is no action to perform.970    llvm::SmallVector<AMDGPUStreamCallbackTy *> Callbacks;971 972    /// Space for the action's arguments. A pointer to these arguments is passed973    /// to the action function. Notice the space of arguments is limited.974    union ActionArgsTy {975      MemcpyArgsTy MemcpyArgs;976      ReleaseBufferArgsTy ReleaseBufferArgs;977      ReleaseSignalArgsTy ReleaseSignalArgs;978      void *CallbackArgs;979    };980 981    llvm::SmallVector<ActionArgsTy> ActionArgs;982 983    /// Create an empty slot.984    StreamSlotTy() : Signal(nullptr), Callbacks({}), ActionArgs({}) {}985 986    /// Schedule a host memory copy action on the slot.987    ///988    /// Num times will repeat the copy that many times, sequentually in the dest989    /// buffer.990    Error schedHostMemoryCopy(void *Dst, const void *Src, size_t Size,991                              size_t NumTimes = 1) {992      Callbacks.emplace_back(memcpyAction);993      ActionArgs.emplace_back().MemcpyArgs =994          MemcpyArgsTy{Dst, Src, Size, NumTimes};995      return Plugin::success();996    }997 998    /// Schedule a release buffer action on the slot.999    Error schedReleaseBuffer(void *Buffer, AMDGPUMemoryManagerTy &Manager) {1000      Callbacks.emplace_back(releaseBufferAction);1001      ActionArgs.emplace_back().ReleaseBufferArgs =1002          ReleaseBufferArgsTy{Buffer, &Manager};1003      return Plugin::success();1004    }1005 1006    /// Schedule a signal release action on the slot.1007    Error schedReleaseSignal(AMDGPUSignalTy *SignalToRelease,1008                             AMDGPUSignalManagerTy *SignalManager) {1009      Callbacks.emplace_back(releaseSignalAction);1010      ActionArgs.emplace_back().ReleaseSignalArgs =1011          ReleaseSignalArgsTy{SignalToRelease, SignalManager};1012      return Plugin::success();1013    }1014 1015    /// Register a callback to be called on compleition1016    Error schedCallback(AMDGPUStreamCallbackTy *Func, void *Data) {1017      Callbacks.emplace_back(Func);1018      ActionArgs.emplace_back().CallbackArgs = Data;1019 1020      return Plugin::success();1021    }1022 1023    // Perform the action if needed.1024    Error performAction() {1025      if (Callbacks.empty())1026        return Plugin::success();1027 1028      assert(Callbacks.size() == ActionArgs.size() && "Size mismatch");1029      for (auto [Callback, ActionArg] : llvm::zip(Callbacks, ActionArgs)) {1030        // Perform the action.1031        if (Callback == memcpyAction) {1032          if (auto Err = memcpyAction(&ActionArg))1033            return Err;1034        } else if (Callback == releaseBufferAction) {1035          if (auto Err = releaseBufferAction(&ActionArg))1036            return Err;1037        } else if (Callback == releaseSignalAction) {1038          if (auto Err = releaseSignalAction(&ActionArg))1039            return Err;1040        } else if (Callback) {1041          if (auto Err = Callback(ActionArg.CallbackArgs))1042            return Err;1043        }1044      }1045 1046      // Invalidate the action.1047      Callbacks.clear();1048      ActionArgs.clear();1049 1050      return Plugin::success();1051    }1052  };1053 1054  /// The device agent where the stream was created.1055  hsa_agent_t Agent;1056 1057  /// The queue that the stream uses to launch kernels.1058  AMDGPUQueueTy *Queue;1059 1060  /// The manager of signals to reuse signals.1061  AMDGPUSignalManagerTy &SignalManager;1062 1063  /// A reference to the associated device.1064  GenericDeviceTy &Device;1065 1066  /// Array of stream slots. Use std::deque because it can dynamically grow1067  /// without invalidating the already inserted elements. For instance, the1068  /// std::vector may invalidate the elements by reallocating the internal1069  /// array if there is not enough space on new insertions.1070  std::deque<StreamSlotTy> Slots;1071 1072  /// The next available slot on the queue. This is reset to zero each time the1073  /// stream is synchronized. It also indicates the current number of consumed1074  /// slots at a given time.1075  uint32_t NextSlot;1076 1077  /// The synchronization id. This number is increased each time the stream is1078  /// synchronized. It is useful to detect if an AMDGPUEventTy points to an1079  /// operation that was already finalized in a previous stream sycnhronize.1080  uint32_t SyncCycle;1081 1082  /// Mutex to protect stream's management.1083  mutable std::mutex Mutex;1084 1085  /// Timeout hint for HSA actively waiting for signal value to change1086  const uint64_t StreamBusyWaitMicroseconds;1087 1088  /// Indicate to spread data transfers across all available SDMAs1089  bool UseMultipleSdmaEngines;1090 1091  struct CallbackDataType {1092    HostFnType UserFn;1093    void *UserData;1094    AMDGPUSignalTy *OutputSignal;1095  };1096  /// Wrapper function for implementing host callbacks1097  static bool callbackWrapper([[maybe_unused]] hsa_signal_value_t Signal,1098                              void *UserData) {1099    auto CallbackData = reinterpret_cast<CallbackDataType *>(UserData);1100    CallbackData->UserFn(CallbackData->UserData);1101    CallbackData->OutputSignal->signal();1102    delete CallbackData;1103    return false;1104  }1105 1106  /// Return the current number of asynchronous operations on the stream.1107  uint32_t size() const { return NextSlot; }1108 1109  /// Return the last valid slot on the stream.1110  uint32_t last() const { return size() - 1; }1111 1112  /// Consume one slot from the stream. Since the stream uses signals on demand1113  /// and releases them once the slot is no longer used, the function requires1114  /// an idle signal for the new consumed slot.1115  std::pair<uint32_t, AMDGPUSignalTy *> consume(AMDGPUSignalTy *OutputSignal) {1116    // Double the stream size if needed. Since we use std::deque, this operation1117    // does not invalidate the already added slots.1118    if (Slots.size() == NextSlot)1119      Slots.resize(Slots.size() * 2);1120 1121    // Update the next available slot and the stream size.1122    uint32_t Curr = NextSlot++;1123 1124    // Retrieve the input signal, if any, of the current operation.1125    AMDGPUSignalTy *InputSignal = (Curr > 0) ? Slots[Curr - 1].Signal : nullptr;1126 1127    // Set the output signal of the current slot.1128    Slots[Curr].Signal = OutputSignal;1129 1130    return std::make_pair(Curr, InputSignal);1131  }1132 1133  /// Complete all pending post actions and reset the stream after synchronizing1134  /// or positively querying the stream.1135  Error complete() {1136    for (uint32_t Slot = 0; Slot < NextSlot; ++Slot) {1137      // Take the post action of the operation if any.1138      if (auto Err = Slots[Slot].performAction())1139        return Err;1140 1141      // Release the slot's signal if possible. Otherwise, another user will.1142      if (Slots[Slot].Signal->decreaseUseCount())1143        if (auto Err = SignalManager.returnResource(Slots[Slot].Signal))1144          return Err;1145 1146      Slots[Slot].Signal = nullptr;1147    }1148 1149    // Reset the stream slots to zero.1150    NextSlot = 0;1151 1152    // Increase the synchronization id since the stream completed a sync cycle.1153    SyncCycle += 1;1154 1155    return Plugin::success();1156  }1157 1158  /// Complete pending post actions until and including the event in target1159  /// slot.1160  Error completeUntil(uint32_t TargetSlot) {1161    for (uint32_t Slot = 0; Slot <= TargetSlot; ++Slot) {1162      // Take the post action of the operation if any.1163      if (auto Err = Slots[Slot].performAction())1164        return Err;1165    }1166 1167    return Plugin::success();1168  }1169 1170  /// Make the current stream wait on a specific operation of another stream.1171  /// The idea is to make the current stream waiting on two signals: 1) the last1172  /// signal of the current stream, and 2) the last signal of the other stream.1173  /// Use a barrier packet with two input signals.1174  Error waitOnStreamOperation(AMDGPUStreamTy &OtherStream, uint32_t Slot) {1175    if (Queue == nullptr)1176      return Plugin::error(ErrorCode::INVALID_NULL_POINTER,1177                           "target queue was nullptr");1178 1179    /// The signal that we must wait from the other stream.1180    AMDGPUSignalTy *OtherSignal = OtherStream.Slots[Slot].Signal;1181 1182    // Prevent the release of the other stream's signal.1183    OtherSignal->increaseUseCount();1184 1185    // Retrieve an available signal for the operation's output.1186    AMDGPUSignalTy *OutputSignal = nullptr;1187    if (auto Err = SignalManager.getResource(OutputSignal))1188      return Err;1189    OutputSignal->reset();1190    OutputSignal->increaseUseCount();1191 1192    // Consume stream slot and compute dependencies.1193    auto [Curr, InputSignal] = consume(OutputSignal);1194 1195    // Setup the post action to release the signal.1196    if (auto Err = Slots[Curr].schedReleaseSignal(OtherSignal, &SignalManager))1197      return Err;1198 1199    // Push a barrier into the queue with both input signals.1200    return Queue->pushBarrier(OutputSignal, InputSignal, OtherSignal);1201  }1202 1203  /// Callback for running a specific asynchronous operation. This callback is1204  /// used for hsa_amd_signal_async_handler. The argument is the operation that1205  /// should be executed. Notice we use the post action mechanism to codify the1206  /// asynchronous operation.1207  static bool asyncActionCallback(hsa_signal_value_t Value, void *Args) {1208    StreamSlotTy *Slot = reinterpret_cast<StreamSlotTy *>(Args);1209    assert(Slot && "Invalid slot");1210    assert(Slot->Signal && "Invalid signal");1211 1212    // This thread is outside the stream mutex. Make sure the thread sees the1213    // changes on the slot.1214    std::atomic_thread_fence(std::memory_order_acquire);1215 1216    // Perform the operation.1217    if (auto Err = Slot->performAction())1218      FATAL_MESSAGE(1, "Error performing post action: %s",1219                    toString(std::move(Err)).data());1220 1221    // Signal the output signal to notify the asynchronous operation finalized.1222    Slot->Signal->signal();1223 1224    // Unregister callback.1225    return false;1226  }1227 1228  // Callback for host-to-host memory copies. This is an asynchronous action.1229  static Error memcpyAction(void *Data) {1230    MemcpyArgsTy *Args = reinterpret_cast<MemcpyArgsTy *>(Data);1231    assert(Args && "Invalid arguments");1232    assert(Args->Dst && "Invalid destination buffer");1233    assert(Args->Src && "Invalid source buffer");1234 1235    auto *BasePtr = Args->Dst;1236    for (size_t I = 0; I < Args->NumTimes; I++) {1237      std::memcpy(BasePtr, Args->Src, Args->Size);1238      BasePtr = reinterpret_cast<uint8_t *>(BasePtr) + Args->Size;1239    }1240 1241    return Plugin::success();1242  }1243 1244  /// Releasing a memory buffer to a memory manager. This is a post completion1245  /// action. There are two kinds of memory buffers:1246  ///   1. For kernel arguments. This buffer can be freed after receiving the1247  ///   kernel completion signal.1248  ///   2. For H2D transfers that need pinned memory space for staging. This1249  ///   buffer can be freed after receiving the transfer completion signal.1250  ///   3. For D2H transfers that need pinned memory space for staging. This1251  ///   buffer cannot be freed after receiving the transfer completion signal1252  ///   because of the following asynchronous H2H callback.1253  ///      For this reason, This action can only be taken at1254  ///      AMDGPUStreamTy::complete()1255  /// Because of the case 3, all releaseBufferActions are taken at1256  /// AMDGPUStreamTy::complete() in the current implementation.1257  static Error releaseBufferAction(void *Data) {1258    ReleaseBufferArgsTy *Args = reinterpret_cast<ReleaseBufferArgsTy *>(Data);1259    assert(Args && "Invalid arguments");1260    assert(Args->MemoryManager && "Invalid memory manager");1261 1262    // Release the allocation to the memory manager.1263    return Args->MemoryManager->deallocate(Args->Buffer);1264  }1265 1266  /// Releasing a signal object back to SignalManager. This is a post completion1267  /// action. This action can only be taken at AMDGPUStreamTy::complete()1268  static Error releaseSignalAction(void *Data) {1269    ReleaseSignalArgsTy *Args = reinterpret_cast<ReleaseSignalArgsTy *>(Data);1270    assert(Args && "Invalid arguments");1271    assert(Args->Signal && "Invalid signal");1272    assert(Args->SignalManager && "Invalid signal manager");1273 1274    // Release the signal if needed.1275    if (Args->Signal->decreaseUseCount())1276      if (auto Err = Args->SignalManager->returnResource(Args->Signal))1277        return Err;1278 1279    return Plugin::success();1280  }1281 1282public:1283  /// Create an empty stream associated with a specific device.1284  AMDGPUStreamTy(AMDGPUDeviceTy &Device);1285 1286  /// Initialize the stream's signals.1287  Error init() { return Plugin::success(); }1288 1289  /// Deinitialize the stream's signals.1290  Error deinit() { return Plugin::success(); }1291 1292  /// Push a asynchronous kernel to the stream. The kernel arguments must be1293  /// placed in a special allocation for kernel args and must keep alive until1294  /// the kernel finalizes. Once the kernel is finished, the stream will release1295  /// the kernel args buffer to the specified memory manager.1296  Error pushKernelLaunch(const AMDGPUKernelTy &Kernel, void *KernelArgs,1297                         uint32_t NumThreads[3], uint32_t NumBlocks[3],1298                         uint32_t GroupSize, uint64_t StackSize,1299                         AMDGPUMemoryManagerTy &MemoryManager) {1300    if (Queue == nullptr)1301      return Plugin::error(ErrorCode::INVALID_NULL_POINTER,1302                           "target queue was nullptr");1303 1304    // Retrieve an available signal for the operation's output.1305    AMDGPUSignalTy *OutputSignal = nullptr;1306    if (auto Err = SignalManager.getResource(OutputSignal))1307      return Err;1308    OutputSignal->reset();1309    OutputSignal->increaseUseCount();1310 1311    std::lock_guard<std::mutex> StreamLock(Mutex);1312 1313    // Consume stream slot and compute dependencies.1314    auto [Curr, InputSignal] = consume(OutputSignal);1315 1316    // Setup the post action to release the kernel args buffer.1317    if (auto Err = Slots[Curr].schedReleaseBuffer(KernelArgs, MemoryManager))1318      return Err;1319 1320    // If we are running an RPC server we want to wake up the server thread1321    // whenever there is a kernel running and let it sleep otherwise.1322    if (Device.getRPCServer())1323      Device.Plugin.getRPCServer().Thread->notify();1324 1325    // Push the kernel with the output signal and an input signal (optional)1326    if (auto Err = Queue->pushKernelLaunch(Kernel, KernelArgs, NumThreads,1327                                           NumBlocks, GroupSize, StackSize,1328                                           OutputSignal, InputSignal))1329      return Err;1330 1331    // Register a callback to indicate when the kernel is complete.1332    if (Device.getRPCServer()) {1333      if (auto Err = Slots[Curr].schedCallback(1334              [](void *Data) -> llvm::Error {1335                GenericPluginTy &Plugin =1336                    *reinterpret_cast<GenericPluginTy *>(Data);1337                Plugin.getRPCServer().Thread->finish();1338                return Error::success();1339              },1340              &Device.Plugin))1341        return Err;1342    }1343    return Plugin::success();1344  }1345 1346  /// Push an asynchronous memory copy between pinned memory buffers.1347  Error pushPinnedMemoryCopyAsync(void *Dst, const void *Src,1348                                  uint64_t CopySize) {1349    // Retrieve an available signal for the operation's output.1350    AMDGPUSignalTy *OutputSignal = nullptr;1351    if (auto Err = SignalManager.getResource(OutputSignal))1352      return Err;1353    OutputSignal->reset();1354    OutputSignal->increaseUseCount();1355 1356    std::lock_guard<std::mutex> Lock(Mutex);1357 1358    // Consume stream slot and compute dependencies.1359    auto [Curr, InputSignal] = consume(OutputSignal);1360 1361    // Issue the async memory copy.1362    if (InputSignal && InputSignal->load()) {1363      hsa_signal_t InputSignalRaw = InputSignal->get();1364      return hsa_utils::asyncMemCopy(UseMultipleSdmaEngines, Dst, Agent, Src,1365                                     Agent, CopySize, 1, &InputSignalRaw,1366                                     OutputSignal->get());1367    }1368 1369    return hsa_utils::asyncMemCopy(UseMultipleSdmaEngines, Dst, Agent, Src,1370                                   Agent, CopySize, 0, nullptr,1371                                   OutputSignal->get());1372  }1373 1374  /// Push an asynchronous memory copy device-to-host involving an unpinned1375  /// memory buffer. The operation consists of a two-step copy from the1376  /// device buffer to an intermediate pinned host buffer, and then, to a1377  /// unpinned host buffer. Both operations are asynchronous and dependent.1378  /// The intermediate pinned buffer will be released to the specified memory1379  /// manager once the operation completes.1380  Error pushMemoryCopyD2HAsync(void *Dst, const void *Src, void *Inter,1381                               uint64_t CopySize,1382                               AMDGPUMemoryManagerTy &MemoryManager) {1383    // Retrieve available signals for the operation's outputs.1384    AMDGPUSignalTy *OutputSignals[2] = {};1385    if (auto Err = SignalManager.getResources(/*Num=*/2, OutputSignals))1386      return Err;1387    for (auto *Signal : OutputSignals) {1388      Signal->reset();1389      Signal->increaseUseCount();1390    }1391 1392    std::lock_guard<std::mutex> Lock(Mutex);1393 1394    // Consume stream slot and compute dependencies.1395    auto [Curr, InputSignal] = consume(OutputSignals[0]);1396 1397    // Setup the post action for releasing the intermediate buffer.1398    if (auto Err = Slots[Curr].schedReleaseBuffer(Inter, MemoryManager))1399      return Err;1400 1401    // Issue the first step: device to host transfer. Avoid defining the input1402    // dependency if already satisfied.1403    if (InputSignal && InputSignal->load()) {1404      hsa_signal_t InputSignalRaw = InputSignal->get();1405      if (auto Err = hsa_utils::asyncMemCopy(1406              UseMultipleSdmaEngines, Inter, Agent, Src, Agent, CopySize, 1,1407              &InputSignalRaw, OutputSignals[0]->get()))1408        return Err;1409    } else {1410      if (auto Err = hsa_utils::asyncMemCopy(UseMultipleSdmaEngines, Inter,1411                                             Agent, Src, Agent, CopySize, 0,1412                                             nullptr, OutputSignals[0]->get()))1413        return Err;1414    }1415 1416    // Consume another stream slot and compute dependencies.1417    std::tie(Curr, InputSignal) = consume(OutputSignals[1]);1418    assert(InputSignal && "Invalid input signal");1419 1420    // The std::memcpy is done asynchronously using an async handler. We store1421    // the function's information in the action but it's not actually an action.1422    if (auto Err = Slots[Curr].schedHostMemoryCopy(Dst, Inter, CopySize))1423      return Err;1424 1425    // Make changes on this slot visible to the async handler's thread.1426    std::atomic_thread_fence(std::memory_order_release);1427 1428    // Issue the second step: host to host transfer.1429    hsa_status_t Status = hsa_amd_signal_async_handler(1430        InputSignal->get(), HSA_SIGNAL_CONDITION_EQ, 0, asyncActionCallback,1431        (void *)&Slots[Curr]);1432 1433    return Plugin::check(Status, "error in hsa_amd_signal_async_handler: %s");1434  }1435 1436  /// Push an asynchronous memory copy host-to-device involving an unpinned1437  /// memory buffer. The operation consists of a two-step copy from the1438  /// unpinned host buffer to an intermediate pinned host buffer, and then, to1439  /// the pinned host buffer. Both operations are asynchronous and dependent.1440  /// The intermediate pinned buffer will be released to the specified memory1441  /// manager once the operation completes.1442  Error pushMemoryCopyH2DAsync(void *Dst, const void *Src, void *Inter,1443                               uint64_t CopySize,1444                               AMDGPUMemoryManagerTy &MemoryManager,1445                               size_t NumTimes = 1) {1446    // Retrieve available signals for the operation's outputs.1447    AMDGPUSignalTy *OutputSignals[2] = {};1448    if (auto Err = SignalManager.getResources(/*Num=*/2, OutputSignals))1449      return Err;1450    for (auto *Signal : OutputSignals) {1451      Signal->reset();1452      Signal->increaseUseCount();1453    }1454 1455    AMDGPUSignalTy *OutputSignal = OutputSignals[0];1456 1457    std::lock_guard<std::mutex> Lock(Mutex);1458 1459    // Consume stream slot and compute dependencies.1460    auto [Curr, InputSignal] = consume(OutputSignal);1461 1462    // Issue the first step: host to host transfer.1463    if (InputSignal && InputSignal->load()) {1464      // The std::memcpy is done asynchronously using an async handler. We store1465      // the function's information in the action but it is not actually a1466      // post action.1467      if (auto Err =1468              Slots[Curr].schedHostMemoryCopy(Inter, Src, CopySize, NumTimes))1469        return Err;1470 1471      // Make changes on this slot visible to the async handler's thread.1472      std::atomic_thread_fence(std::memory_order_release);1473 1474      hsa_status_t Status = hsa_amd_signal_async_handler(1475          InputSignal->get(), HSA_SIGNAL_CONDITION_EQ, 0, asyncActionCallback,1476          (void *)&Slots[Curr]);1477 1478      if (auto Err = Plugin::check(Status,1479                                   "error in hsa_amd_signal_async_handler: %s"))1480        return Err;1481 1482      // Let's use now the second output signal.1483      OutputSignal = OutputSignals[1];1484 1485      // Consume another stream slot and compute dependencies.1486      std::tie(Curr, InputSignal) = consume(OutputSignal);1487    } else {1488      // All preceding operations completed, copy the memory synchronously.1489      auto *InterPtr = Inter;1490      for (size_t I = 0; I < NumTimes; I++) {1491        std::memcpy(InterPtr, Src, CopySize);1492        InterPtr = reinterpret_cast<uint8_t *>(InterPtr) + CopySize;1493      }1494 1495      // Return the second signal because it will not be used.1496      OutputSignals[1]->decreaseUseCount();1497      if (auto Err = SignalManager.returnResource(OutputSignals[1]))1498        return Err;1499    }1500 1501    // Setup the post action to release the intermediate pinned buffer.1502    if (auto Err = Slots[Curr].schedReleaseBuffer(Inter, MemoryManager))1503      return Err;1504 1505    // Issue the second step: host to device transfer. Avoid defining the input1506    // dependency if already satisfied.1507    if (InputSignal && InputSignal->load()) {1508      hsa_signal_t InputSignalRaw = InputSignal->get();1509      return hsa_utils::asyncMemCopy(UseMultipleSdmaEngines, Dst, Agent, Inter,1510                                     Agent, CopySize * NumTimes, 1,1511                                     &InputSignalRaw, OutputSignal->get());1512    }1513    return hsa_utils::asyncMemCopy(UseMultipleSdmaEngines, Dst, Agent, Inter,1514                                   Agent, CopySize * NumTimes, 0, nullptr,1515                                   OutputSignal->get());1516  }1517 1518  // AMDGPUDeviceTy is incomplete here, passing the underlying agent instead1519  Error pushMemoryCopyD2DAsync(void *Dst, hsa_agent_t DstAgent, const void *Src,1520                               hsa_agent_t SrcAgent, uint64_t CopySize) {1521    AMDGPUSignalTy *OutputSignal;1522    if (auto Err = SignalManager.getResources(/*Num=*/1, &OutputSignal))1523      return Err;1524    OutputSignal->reset();1525    OutputSignal->increaseUseCount();1526 1527    std::lock_guard<std::mutex> Lock(Mutex);1528 1529    // Consume stream slot and compute dependencies.1530    auto [Curr, InputSignal] = consume(OutputSignal);1531 1532    // The agents need to have access to the corresponding memory1533    // This is presently only true if the pointers were originally1534    // allocated by this runtime or the caller made the appropriate1535    // access calls.1536 1537    if (InputSignal && InputSignal->load()) {1538      hsa_signal_t InputSignalRaw = InputSignal->get();1539      return hsa_utils::asyncMemCopy(UseMultipleSdmaEngines, Dst, DstAgent, Src,1540                                     SrcAgent, CopySize, 1, &InputSignalRaw,1541                                     OutputSignal->get());1542    }1543    return hsa_utils::asyncMemCopy(UseMultipleSdmaEngines, Dst, DstAgent, Src,1544                                   SrcAgent, CopySize, 0, nullptr,1545                                   OutputSignal->get());1546  }1547 1548  Error pushHostCallback(HostFnType Callback, void *UserData) {1549    // Retrieve an available signal for the operation's output.1550    AMDGPUSignalTy *OutputSignal = nullptr;1551    if (auto Err = SignalManager.getResource(OutputSignal))1552      return Err;1553    OutputSignal->reset();1554    OutputSignal->increaseUseCount();1555 1556    AMDGPUSignalTy *InputSignal;1557    {1558      std::lock_guard<std::mutex> Lock(Mutex);1559 1560      // Consume stream slot and compute dependencies.1561      InputSignal = consume(OutputSignal).second;1562    }1563 1564    auto *CallbackData = new CallbackDataType{Callback, UserData, OutputSignal};1565    if (InputSignal && InputSignal->load()) {1566      hsa_status_t Status = hsa_amd_signal_async_handler(1567          InputSignal->get(), HSA_SIGNAL_CONDITION_EQ, 0, callbackWrapper,1568          CallbackData);1569 1570      return Plugin::check(Status, "error in hsa_amd_signal_async_handler: %s");1571    }1572 1573    // No dependencies - schedule it now.1574    // Using a seperate thread because this function should run asynchronously1575    // and not block the main thread.1576    std::thread([](void *CallbackData) { callbackWrapper(0, CallbackData); },1577                CallbackData)1578        .detach();1579    return Plugin::success();1580  }1581 1582  /// Synchronize with the stream. The current thread waits until all operations1583  /// are finalized and it performs the pending post actions (i.e., releasing1584  /// intermediate buffers).1585  Error synchronize() {1586    std::lock_guard<std::mutex> Lock(Mutex);1587 1588    // No need to synchronize anything.1589    if (size() == 0)1590      return Plugin::success();1591 1592    // Wait until all previous operations on the stream have completed.1593    if (auto Err =1594            Slots[last()].Signal->wait(StreamBusyWaitMicroseconds, &Device))1595      return Err;1596 1597    // Reset the stream and perform all pending post actions.1598    return complete();1599  }1600 1601  /// Synchronize the stream until the given event. The current thread waits1602  /// until the provided event is finalized, and it performs the pending post1603  /// actions for that and prior events.1604  Error synchronizeOn(AMDGPUEventTy &Event);1605 1606  /// Return true if the event from this queue is complete1607  Expected<bool> isEventComplete(const AMDGPUEventTy &Event);1608 1609  /// Query the stream and complete pending post actions if operations finished.1610  /// Return whether all the operations completed. This operation does not block1611  /// the calling thread.1612  Expected<bool> query() {1613    std::lock_guard<std::mutex> Lock(Mutex);1614 1615    // No need to query anything.1616    if (size() == 0)1617      return true;1618 1619    // The last operation did not complete yet. Return directly.1620    if (Slots[last()].Signal->load())1621      return false;1622 1623    // Reset the stream and perform all pending post actions.1624    if (auto Err = complete())1625      return std::move(Err);1626 1627    return true;1628  }1629 1630  const AMDGPUQueueTy *getQueue() const { return Queue; }1631 1632  /// Record the state of the stream on an event.1633  Error recordEvent(AMDGPUEventTy &Event) const;1634 1635  /// Make the stream wait on an event.1636  Error waitEvent(const AMDGPUEventTy &Event);1637 1638  friend struct AMDGPUStreamManagerTy;1639};1640 1641/// Class representing an event on AMDGPU. The event basically stores some1642/// information regarding the state of the recorded stream.1643struct AMDGPUEventTy {1644  /// Create an empty event.1645  AMDGPUEventTy(AMDGPUDeviceTy &Device)1646      : RecordedStream(nullptr), RecordedSlot(-1), RecordedSyncCycle(-1) {}1647 1648  /// Initialize and deinitialize.1649  Error init() { return Plugin::success(); }1650  Error deinit() { return Plugin::success(); }1651 1652  /// Record the state of a stream on the event.1653  Error record(AMDGPUStreamTy &Stream) {1654    std::lock_guard<std::mutex> Lock(Mutex);1655 1656    // Ignore the last recorded stream.1657    RecordedStream = &Stream;1658 1659    return Stream.recordEvent(*this);1660  }1661 1662  /// Make a stream wait on the current event.1663  Error wait(AMDGPUStreamTy &Stream) {1664    std::lock_guard<std::mutex> Lock(Mutex);1665 1666    if (!RecordedStream)1667      return Plugin::error(ErrorCode::INVALID_ARGUMENT,1668                           "event does not have any recorded stream");1669 1670    // Synchronizing the same stream. Do nothing.1671    if (RecordedStream == &Stream)1672      return Plugin::success();1673 1674    // No need to wait anything, the recorded stream already finished the1675    // corresponding operation.1676    if (RecordedSlot < 0)1677      return Plugin::success();1678 1679    return Stream.waitEvent(*this);1680  }1681 1682  Error sync() {1683    std::lock_guard<std::mutex> Lock(Mutex);1684 1685    if (!RecordedStream)1686      return Plugin::error(ErrorCode::INVALID_ARGUMENT,1687                           "event does not have any recorded stream");1688 1689    // No need to wait on anything, the recorded stream already finished the1690    // corresponding operation.1691    if (RecordedSlot < 0)1692      return Plugin::success();1693 1694    return RecordedStream->synchronizeOn(*this);1695  }1696 1697protected:1698  /// The stream registered in this event.1699  AMDGPUStreamTy *RecordedStream;1700 1701  /// The recordered operation on the recorded stream.1702  int64_t RecordedSlot;1703 1704  /// The sync cycle when the stream was recorded. Used to detect stale events.1705  int64_t RecordedSyncCycle;1706 1707  /// Mutex to safely access event fields.1708  mutable std::mutex Mutex;1709 1710  friend struct AMDGPUStreamTy;1711};1712 1713Error AMDGPUStreamTy::recordEvent(AMDGPUEventTy &Event) const {1714  std::lock_guard<std::mutex> Lock(Mutex);1715 1716  if (size() > 0) {1717    // Record the synchronize identifier (to detect stale recordings) and1718    // the last valid stream's operation.1719    Event.RecordedSyncCycle = SyncCycle;1720    Event.RecordedSlot = last();1721 1722    assert(Event.RecordedSyncCycle >= 0 && "Invalid recorded sync cycle");1723    assert(Event.RecordedSlot >= 0 && "Invalid recorded slot");1724  } else {1725    // The stream is empty, everything already completed, record nothing.1726    Event.RecordedSyncCycle = -1;1727    Event.RecordedSlot = -1;1728  }1729  return Plugin::success();1730}1731 1732Error AMDGPUStreamTy::waitEvent(const AMDGPUEventTy &Event) {1733  // Retrieve the recorded stream on the event.1734  AMDGPUStreamTy &RecordedStream = *Event.RecordedStream;1735 1736  std::scoped_lock<std::mutex, std::mutex> Lock(Mutex, RecordedStream.Mutex);1737 1738  // The recorded stream already completed the operation because the synchronize1739  // identifier is already outdated.1740  if (RecordedStream.SyncCycle != (uint32_t)Event.RecordedSyncCycle)1741    return Plugin::success();1742 1743  // Again, the recorded stream already completed the operation, the last1744  // operation's output signal is satisfied.1745  if (!RecordedStream.Slots[Event.RecordedSlot].Signal->load())1746    return Plugin::success();1747 1748  // Otherwise, make the current stream wait on the other stream's operation.1749  return waitOnStreamOperation(RecordedStream, Event.RecordedSlot);1750}1751 1752Error AMDGPUStreamTy::synchronizeOn(AMDGPUEventTy &Event) {1753  std::lock_guard<std::mutex> Lock(Mutex);1754 1755  // If this event was for an older sync cycle, it has already been finalized1756  if (Event.RecordedSyncCycle < SyncCycle)1757    return Plugin::success();1758  assert(Event.RecordedSyncCycle == SyncCycle && "event is from the future?");1759 1760  // Wait until the requested slot has completed1761  if (auto Err = Slots[Event.RecordedSlot].Signal->wait(1762          StreamBusyWaitMicroseconds, &Device))1763    return Err;1764 1765  // If the event is the last one in the stream, just do a full finalize1766  if (Event.RecordedSlot == last())1767    return complete();1768 1769  // Otherwise, only finalize until the appropriate event1770  return completeUntil(Event.RecordedSlot);1771}1772 1773Expected<bool> AMDGPUStreamTy::isEventComplete(const AMDGPUEventTy &Event) {1774  std::lock_guard<std::mutex> Lock(Mutex);1775  assert(Event.RecordedStream == this && "event is for a different stream");1776 1777  if (Event.RecordedSyncCycle < SyncCycle) {1778    return true;1779  }1780  assert(Event.RecordedSyncCycle == SyncCycle && "event is from the future?");1781 1782  return !Slots[Event.RecordedSlot].Signal->load();1783}1784 1785struct AMDGPUStreamManagerTy final1786    : GenericDeviceResourceManagerTy<AMDGPUResourceRef<AMDGPUStreamTy>> {1787  using ResourceRef = AMDGPUResourceRef<AMDGPUStreamTy>;1788  using ResourcePoolTy = GenericDeviceResourceManagerTy<ResourceRef>;1789 1790  AMDGPUStreamManagerTy(GenericDeviceTy &Device, hsa_agent_t HSAAgent)1791      : GenericDeviceResourceManagerTy(Device), Device(Device),1792        OMPX_QueueTracking("LIBOMPTARGET_AMDGPU_HSA_QUEUE_BUSY_TRACKING", true),1793        NextQueue(0), Agent(HSAAgent) {}1794 1795  Error init(uint32_t InitialSize, int NumHSAQueues, int HSAQueueSize) {1796    Queues = std::vector<AMDGPUQueueTy>(NumHSAQueues);1797    QueueSize = HSAQueueSize;1798    MaxNumQueues = NumHSAQueues;1799    // Initialize one queue eagerly1800    if (auto Err = Queues.front().init(Device, Agent, QueueSize))1801      return Err;1802 1803    return GenericDeviceResourceManagerTy::init(InitialSize);1804  }1805 1806  /// Deinitialize the resource pool and delete all resources. This function1807  /// must be called before the destructor.1808  Error deinit() override {1809    // De-init all queues1810    for (AMDGPUQueueTy &Queue : Queues) {1811      if (auto Err = Queue.deinit())1812        return Err;1813    }1814 1815    return GenericDeviceResourceManagerTy::deinit();1816  }1817 1818  /// Get a single stream from the pool or create new resources.1819  virtual Error getResource(AMDGPUStreamTy *&StreamHandle) override {1820    return getResourcesImpl(1, &StreamHandle, [this](AMDGPUStreamTy *&Handle) {1821      return assignNextQueue(Handle);1822    });1823  }1824 1825  /// Return stream to the pool.1826  virtual Error returnResource(AMDGPUStreamTy *StreamHandle) override {1827    return returnResourceImpl(StreamHandle, [](AMDGPUStreamTy *Handle) {1828      Handle->Queue->removeUser();1829      return Plugin::success();1830    });1831  }1832 1833private:1834  /// Search for and assign an preferably idle queue to the given Stream. If1835  /// there is no queue without current users, choose the queue with the lowest1836  /// user count. If utilization is ignored: use round robin selection.1837  inline Error assignNextQueue(AMDGPUStreamTy *Stream) {1838    // Start from zero when tracking utilization, otherwise: round robin policy.1839    uint32_t Index = OMPX_QueueTracking ? 0 : NextQueue++ % MaxNumQueues;1840 1841    if (OMPX_QueueTracking) {1842      // Find the least used queue.1843      for (uint32_t I = 0; I < MaxNumQueues; ++I) {1844        // Early exit when an initialized queue is idle.1845        if (Queues[I].isInitialized() && Queues[I].getUserCount() == 0) {1846          Index = I;1847          break;1848        }1849 1850        // Update the least used queue.1851        if (Queues[Index].getUserCount() > Queues[I].getUserCount())1852          Index = I;1853      }1854    }1855 1856    // Make sure the queue is initialized, then add user & assign.1857    if (auto Err = Queues[Index].init(Device, Agent, QueueSize))1858      return Err;1859    Queues[Index].addUser();1860    Stream->Queue = &Queues[Index];1861 1862    return Plugin::success();1863  }1864 1865  /// The device associated with this stream.1866  GenericDeviceTy &Device;1867 1868  /// Envar for controlling the tracking of busy HSA queues.1869  BoolEnvar OMPX_QueueTracking;1870 1871  /// The next queue index to use for round robin selection.1872  uint32_t NextQueue;1873 1874  /// The queues which are assigned to requested streams.1875  std::vector<AMDGPUQueueTy> Queues;1876 1877  /// The corresponding device as HSA agent.1878  hsa_agent_t Agent;1879 1880  /// The maximum number of queues.1881  uint32_t MaxNumQueues;1882 1883  /// The size of created queues.1884  uint32_t QueueSize;1885};1886 1887/// Abstract class that holds the common members of the actual kernel devices1888/// and the host device. Both types should inherit from this class.1889struct AMDGenericDeviceTy {1890  AMDGenericDeviceTy() {}1891 1892  virtual ~AMDGenericDeviceTy() {}1893 1894  /// Create all memory pools which the device has access to and classify them.1895  Error initMemoryPools() {1896    // Retrieve all memory pools from the device agent(s).1897    Error Err = retrieveAllMemoryPools();1898    if (Err)1899      return Err;1900 1901    for (AMDGPUMemoryPoolTy *MemoryPool : AllMemoryPools) {1902      // Initialize the memory pool and retrieve some basic info.1903      Error Err = MemoryPool->init();1904      if (Err)1905        return Err;1906 1907      if (!MemoryPool->isGlobal())1908        continue;1909 1910      // Classify the memory pools depending on their properties.1911      if (MemoryPool->isFineGrained()) {1912        FineGrainedMemoryPools.push_back(MemoryPool);1913        if (MemoryPool->supportsKernelArgs())1914          ArgsMemoryPools.push_back(MemoryPool);1915      } else if (MemoryPool->isCoarseGrained()) {1916        CoarseGrainedMemoryPools.push_back(MemoryPool);1917      }1918    }1919    return Plugin::success();1920  }1921 1922  /// Destroy all memory pools.1923  Error deinitMemoryPools() {1924    for (AMDGPUMemoryPoolTy *Pool : AllMemoryPools)1925      delete Pool;1926 1927    AllMemoryPools.clear();1928    FineGrainedMemoryPools.clear();1929    CoarseGrainedMemoryPools.clear();1930    ArgsMemoryPools.clear();1931 1932    return Plugin::success();1933  }1934 1935  /// Retrieve and construct all memory pools from the device agent(s).1936  virtual Error retrieveAllMemoryPools() = 0;1937 1938  /// Get the device agent.1939  virtual hsa_agent_t getAgent() const = 0;1940 1941protected:1942  /// Array of all memory pools available to the host agents.1943  llvm::SmallVector<AMDGPUMemoryPoolTy *> AllMemoryPools;1944 1945  /// Array of fine-grained memory pools available to the host agents.1946  llvm::SmallVector<AMDGPUMemoryPoolTy *> FineGrainedMemoryPools;1947 1948  /// Array of coarse-grained memory pools available to the host agents.1949  llvm::SmallVector<AMDGPUMemoryPoolTy *> CoarseGrainedMemoryPools;1950 1951  /// Array of kernel args memory pools available to the host agents.1952  llvm::SmallVector<AMDGPUMemoryPoolTy *> ArgsMemoryPools;1953};1954 1955/// Class representing the host device. This host device may have more than one1956/// HSA host agent. We aggregate all its resources into the same instance.1957struct AMDHostDeviceTy : public AMDGenericDeviceTy {1958  /// Create a host device from an array of host agents.1959  AMDHostDeviceTy(AMDGPUPluginTy &Plugin,1960                  const llvm::SmallVector<hsa_agent_t> &HostAgents)1961      : AMDGenericDeviceTy(), Agents(HostAgents), ArgsMemoryManager(Plugin),1962        PinnedMemoryManager(Plugin) {1963    assert(HostAgents.size() && "No host agent found");1964  }1965 1966  /// Initialize the host device memory pools and the memory managers for1967  /// kernel args and host pinned memory allocations.1968  Error init() {1969    if (auto Err = initMemoryPools())1970      return Err;1971 1972    if (auto Err = ArgsMemoryManager.init(getArgsMemoryPool()))1973      return Err;1974 1975    if (auto Err = PinnedMemoryManager.init(getFineGrainedMemoryPool()))1976      return Err;1977 1978    return Plugin::success();1979  }1980 1981  /// Deinitialize memory pools and managers.1982  Error deinit() {1983    if (auto Err = deinitMemoryPools())1984      return Err;1985 1986    if (auto Err = ArgsMemoryManager.deinit())1987      return Err;1988 1989    if (auto Err = PinnedMemoryManager.deinit())1990      return Err;1991 1992    return Plugin::success();1993  }1994 1995  /// Retrieve and construct all memory pools from the host agents.1996  Error retrieveAllMemoryPools() override {1997    // Iterate through the available pools across the host agents.1998    for (hsa_agent_t Agent : Agents) {1999      Error Err = hsa_utils::iterateAgentMemoryPools(2000          Agent, [&](hsa_amd_memory_pool_t HSAMemoryPool) {2001            AMDGPUMemoryPoolTy *MemoryPool =2002                new AMDGPUMemoryPoolTy(HSAMemoryPool);2003            AllMemoryPools.push_back(MemoryPool);2004            return HSA_STATUS_SUCCESS;2005          });2006      if (Err)2007        return Err;2008    }2009    return Plugin::success();2010  }2011 2012  /// Get one of the host agents. Return always the first agent.2013  hsa_agent_t getAgent() const override { return Agents[0]; }2014 2015  /// Get a memory pool for fine-grained allocations.2016  AMDGPUMemoryPoolTy &getFineGrainedMemoryPool() {2017    assert(!FineGrainedMemoryPools.empty() && "No fine-grained mempool");2018    // Retrieve any memory pool.2019    return *FineGrainedMemoryPools[0];2020  }2021 2022  AMDGPUMemoryPoolTy &getCoarseGrainedMemoryPool() {2023    assert(!CoarseGrainedMemoryPools.empty() && "No coarse-grained mempool");2024    // Retrieve any memory pool.2025    return *CoarseGrainedMemoryPools[0];2026  }2027 2028  /// Get a memory pool for kernel args allocations.2029  AMDGPUMemoryPoolTy &getArgsMemoryPool() {2030    assert(!ArgsMemoryPools.empty() && "No kernelargs mempool");2031    // Retrieve any memory pool.2032    return *ArgsMemoryPools[0];2033  }2034 2035  /// Getters for kernel args and host pinned memory managers.2036  AMDGPUMemoryManagerTy &getArgsMemoryManager() { return ArgsMemoryManager; }2037  AMDGPUMemoryManagerTy &getPinnedMemoryManager() {2038    return PinnedMemoryManager;2039  }2040 2041private:2042  /// Array of agents on the host side.2043  const llvm::SmallVector<hsa_agent_t> Agents;2044 2045  // Memory manager for kernel arguments.2046  AMDGPUMemoryManagerTy ArgsMemoryManager;2047 2048  // Memory manager for pinned memory.2049  AMDGPUMemoryManagerTy PinnedMemoryManager;2050};2051 2052/// Class implementing the AMDGPU device functionalities which derives from the2053/// generic device class.2054struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {2055  // Create an AMDGPU device with a device id and default AMDGPU grid values.2056  AMDGPUDeviceTy(GenericPluginTy &Plugin, int32_t DeviceId, int32_t NumDevices,2057                 AMDHostDeviceTy &HostDevice, hsa_agent_t Agent)2058      : GenericDeviceTy(Plugin, DeviceId, NumDevices, {}), AMDGenericDeviceTy(),2059        OMPX_NumQueues("LIBOMPTARGET_AMDGPU_NUM_HSA_QUEUES", 4),2060        OMPX_QueueSize("LIBOMPTARGET_AMDGPU_HSA_QUEUE_SIZE", 512),2061        OMPX_DefaultTeamsPerCU("LIBOMPTARGET_AMDGPU_TEAMS_PER_CU", 4),2062        OMPX_MaxAsyncCopyBytes("LIBOMPTARGET_AMDGPU_MAX_ASYNC_COPY_BYTES",2063                               1 * 1024 * 1024), // 1MB2064        OMPX_InitialNumSignals("LIBOMPTARGET_AMDGPU_NUM_INITIAL_HSA_SIGNALS",2065                               64),2066        OMPX_StreamBusyWait("LIBOMPTARGET_AMDGPU_STREAM_BUSYWAIT", 2000000),2067        OMPX_UseMultipleSdmaEngines(2068            "LIBOMPTARGET_AMDGPU_USE_MULTIPLE_SDMA_ENGINES", false),2069        OMPX_ApuMaps("OMPX_APU_MAPS", false), AMDGPUStreamManager(*this, Agent),2070        AMDGPUEventManager(*this), AMDGPUSignalManager(*this), Agent(Agent),2071        HostDevice(HostDevice) {}2072 2073  ~AMDGPUDeviceTy() {}2074 2075  /// Initialize the device, its resources and get its properties.2076  Error initImpl(GenericPluginTy &Plugin) override {2077    // First setup all the memory pools.2078    if (auto Err = initMemoryPools())2079      return Err;2080 2081    char GPUName[64];2082    if (auto Err = getDeviceAttr(HSA_AGENT_INFO_NAME, GPUName))2083      return Err;2084    ComputeUnitKind = GPUName;2085 2086    // From the ROCm HSA documentation:2087    // Query the UUID of the agent. The value is an Ascii string with a maximum2088    // of 21 chars including NUL. The string value consists of two parts: header2089    // and body. The header identifies the device type (GPU, CPU, DSP) while the2090    // body encodes the UUID as a 16 digit hex string.2091    //2092    // Agents that do not support UUID will return the string "GPU-XX" or2093    // "CPU-XX" or "DSP-XX" depending on their device type.2094    char UUID[24] = {0};2095    if (auto Err = getDeviceAttr(HSA_AMD_AGENT_INFO_UUID, UUID))2096      return Err;2097    if (!StringRef(UUID).ends_with("-XX"))2098      setDeviceUidFromVendorUid(UUID);2099 2100    // Get the wavefront size.2101    uint32_t WavefrontSize = 0;2102    if (auto Err = getDeviceAttr(HSA_AGENT_INFO_WAVEFRONT_SIZE, WavefrontSize))2103      return Err;2104    GridValues.GV_Warp_Size = WavefrontSize;2105 2106    // Get the frequency of the steady clock. If the attribute is missing2107    // assume running on an older libhsa and default to 0, omp_get_wtime2108    // will be inaccurate but otherwise programs can still run.2109    if (getDeviceAttrRaw(HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY,2110                         ClockFrequency) != HSA_STATUS_SUCCESS)2111      ClockFrequency = 0;2112 2113    // Load the grid values depending on the wavefront.2114    if (WavefrontSize == 32)2115      GridValues = getAMDGPUGridValues<32>();2116    else if (WavefrontSize == 64)2117      GridValues = getAMDGPUGridValues<64>();2118    else2119      return Plugin::error(ErrorCode::UNSUPPORTED,2120                           "unexpected AMDGPU wavefront %d", WavefrontSize);2121 2122    // Get maximum number of workitems per workgroup.2123    uint16_t WorkgroupMaxDim[3];2124    if (auto Err =2125            getDeviceAttr(HSA_AGENT_INFO_WORKGROUP_MAX_DIM, WorkgroupMaxDim))2126      return Err;2127    GridValues.GV_Max_WG_Size = WorkgroupMaxDim[0];2128 2129    // Get maximum number of workgroups.2130    hsa_dim3_t GridMaxDim;2131    if (auto Err = getDeviceAttr(HSA_AGENT_INFO_GRID_MAX_DIM, GridMaxDim))2132      return Err;2133 2134    GridValues.GV_Max_Teams = GridMaxDim.x / GridValues.GV_Max_WG_Size;2135    if (GridValues.GV_Max_Teams == 0)2136      return Plugin::error(ErrorCode::INVALID_ARGUMENT,2137                           "maximum number of teams cannot be zero");2138 2139    // Compute the default number of teams.2140    uint32_t ComputeUnits = 0;2141    if (auto Err =2142            getDeviceAttr(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, ComputeUnits))2143      return Err;2144    GridValues.GV_Default_Num_Teams = ComputeUnits * OMPX_DefaultTeamsPerCU;2145 2146    uint32_t WavesPerCU = 0;2147    if (auto Err =2148            getDeviceAttr(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU, WavesPerCU))2149      return Err;2150    HardwareParallelism = ComputeUnits * WavesPerCU;2151 2152    // Get maximum size of any device queues and maximum number of queues.2153    uint32_t MaxQueueSize;2154    if (auto Err = getDeviceAttr(HSA_AGENT_INFO_QUEUE_MAX_SIZE, MaxQueueSize))2155      return Err;2156 2157    uint32_t MaxQueues;2158    if (auto Err = getDeviceAttr(HSA_AGENT_INFO_QUEUES_MAX, MaxQueues))2159      return Err;2160 2161    // Compute the number of queues and their size.2162    OMPX_NumQueues = std::max(1U, std::min(OMPX_NumQueues.get(), MaxQueues));2163    OMPX_QueueSize = std::min(OMPX_QueueSize.get(), MaxQueueSize);2164 2165    // Initialize stream pool.2166    if (auto Err = AMDGPUStreamManager.init(OMPX_InitialNumStreams,2167                                            OMPX_NumQueues, OMPX_QueueSize))2168      return Err;2169 2170    // Initialize event pool.2171    if (auto Err = AMDGPUEventManager.init(OMPX_InitialNumEvents))2172      return Err;2173 2174    // Initialize signal pool.2175    if (auto Err = AMDGPUSignalManager.init(OMPX_InitialNumSignals))2176      return Err;2177 2178    // Detect if XNACK is enabled2179    SmallVector<SmallString<32>> Targets;2180    if (auto Err = hsa_utils::getTargetTripleAndFeatures(Agent, Targets))2181      return Err;2182    if (!Targets.empty() && Targets[0].str().contains("xnack+"))2183      IsXnackEnabled = true;2184 2185    // detect if device is an APU.2186    if (auto Err = checkIfAPU())2187      return Err;2188 2189    // Retrieve the size of the group memory.2190    for (const auto *Pool : AllMemoryPools) {2191      if (Pool->isGroup()) {2192        if (auto Err = Pool->getAttr(HSA_AMD_MEMORY_POOL_INFO_SIZE,2193                                     MaxBlockSharedMemSize))2194          return Err;2195        break;2196      }2197    }2198 2199    return Plugin::success();2200  }2201 2202  Error unloadBinaryImpl(DeviceImageTy *Image) override {2203    AMDGPUDeviceImageTy &AMDImage = static_cast<AMDGPUDeviceImageTy &>(*Image);2204 2205    // Unload the executable of the image.2206    if (auto Err = AMDImage.unloadExecutable())2207      return Err;2208 2209    // Destroy the associated memory and invalidate the object.2210    Plugin.free(Image);2211    return Error::success();2212  }2213 2214  /// Deinitialize the device and release its resources.2215  Error deinitImpl() override {2216    // Deinitialize the stream and event pools.2217    if (auto Err = AMDGPUStreamManager.deinit())2218      return Err;2219 2220    if (auto Err = AMDGPUEventManager.deinit())2221      return Err;2222 2223    if (auto Err = AMDGPUSignalManager.deinit())2224      return Err;2225 2226    // Invalidate agent reference.2227    Agent = {0};2228 2229    return Plugin::success();2230  }2231 2232  virtual Error callGlobalConstructors(GenericPluginTy &Plugin,2233                                       DeviceImageTy &Image) override {2234    return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/true);2235  }2236 2237  virtual Error callGlobalDestructors(GenericPluginTy &Plugin,2238                                      DeviceImageTy &Image) override {2239    return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/false);2240  }2241 2242  uint64_t getStreamBusyWaitMicroseconds() const { return OMPX_StreamBusyWait; }2243 2244  Expected<std::unique_ptr<MemoryBuffer>>2245  doJITPostProcessing(std::unique_ptr<MemoryBuffer> MB) const override {2246 2247    // TODO: We should try to avoid materialization but there seems to be no2248    // good linker interface w/o file i/o.2249    SmallString<128> LinkerInputFilePath;2250    std::error_code EC = sys::fs::createTemporaryFile("amdgpu-pre-link-jit",2251                                                      "o", LinkerInputFilePath);2252    if (EC)2253      return Plugin::error(ErrorCode::HOST_IO,2254                           "failed to create temporary file for linker");2255 2256    // Write the file's contents to the output file.2257    Expected<std::unique_ptr<FileOutputBuffer>> OutputOrErr =2258        FileOutputBuffer::create(LinkerInputFilePath, MB->getBuffer().size());2259    if (!OutputOrErr)2260      return OutputOrErr.takeError();2261    std::unique_ptr<FileOutputBuffer> Output = std::move(*OutputOrErr);2262    llvm::copy(MB->getBuffer(), Output->getBufferStart());2263    if (Error E = Output->commit())2264      return std::move(E);2265 2266    SmallString<128> LinkerOutputFilePath;2267    EC = sys::fs::createTemporaryFile("amdgpu-pre-link-jit", "so",2268                                      LinkerOutputFilePath);2269    if (EC)2270      return Plugin::error(ErrorCode::HOST_IO,2271                           "failed to create temporary file for linker");2272 2273    const auto &ErrorOrPath = sys::findProgramByName("lld");2274    if (!ErrorOrPath)2275      return createStringError(ErrorCode::HOST_TOOL_NOT_FOUND,2276                               "failed to find `lld` on the PATH.");2277 2278    std::string LLDPath = ErrorOrPath.get();2279    INFO(OMP_INFOTYPE_PLUGIN_KERNEL, getDeviceId(),2280         "Using `%s` to link JITed amdgcn output.", LLDPath.c_str());2281 2282    std::string MCPU = "-plugin-opt=mcpu=" + getComputeUnitKind();2283    StringRef Args[] = {LLDPath,2284                        "-flavor",2285                        "gnu",2286                        "--no-undefined",2287                        "-shared",2288                        MCPU,2289                        "-o",2290                        LinkerOutputFilePath.data(),2291                        LinkerInputFilePath.data()};2292 2293    std::string Error;2294    int RC = sys::ExecuteAndWait(LLDPath, Args, std::nullopt, {}, 0, 0, &Error);2295    if (RC)2296      return Plugin::error(ErrorCode::LINK_FAILURE,2297                           "linking optimized bitcode failed: %s",2298                           Error.c_str());2299 2300    auto BufferOrErr = MemoryBuffer::getFileOrSTDIN(LinkerOutputFilePath);2301    if (!BufferOrErr)2302      return Plugin::error(ErrorCode::HOST_IO,2303                           "failed to open temporary file for lld");2304 2305    // Clean up the temporary files afterwards.2306    if (sys::fs::remove(LinkerOutputFilePath))2307      return Plugin::error(ErrorCode::HOST_IO,2308                           "failed to remove temporary output file for lld");2309    if (sys::fs::remove(LinkerInputFilePath))2310      return Plugin::error(ErrorCode::HOST_IO,2311                           "failed to remove temporary input file for lld");2312 2313    return std::move(*BufferOrErr);2314  }2315 2316  /// See GenericDeviceTy::getComputeUnitKind().2317  std::string getComputeUnitKind() const override { return ComputeUnitKind; }2318 2319  /// Returns the clock frequency for the given AMDGPU device.2320  uint64_t getClockFrequency() const override { return ClockFrequency; }2321 2322  /// Allocate and construct an AMDGPU kernel.2323  Expected<GenericKernelTy &> constructKernel(const char *Name) override {2324    // Allocate and construct the AMDGPU kernel.2325    AMDGPUKernelTy *AMDGPUKernel = Plugin.allocate<AMDGPUKernelTy>();2326    if (!AMDGPUKernel)2327      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,2328                           "failed to allocate memory for AMDGPU kernel");2329 2330    new (AMDGPUKernel) AMDGPUKernelTy(Name);2331 2332    return *AMDGPUKernel;2333  }2334 2335  /// Set the current context to this device's context. Do nothing since the2336  /// AMDGPU devices do not have the concept of contexts.2337  Error setContext() override { return Plugin::success(); }2338 2339  /// AMDGPU returns the product of the number of compute units and the waves2340  /// per compute unit.2341  uint64_t getHardwareParallelism() const override {2342    return HardwareParallelism;2343  }2344 2345  /// We want to set up the RPC server for host services to the GPU if it is2346  /// available.2347  bool shouldSetupRPCServer() const override { return true; }2348 2349  /// The RPC interface should have enough space for all available parallelism.2350  uint64_t requestedRPCPortCount() const override {2351    return getHardwareParallelism();2352  }2353 2354  /// Get the stream of the asynchronous info structure or get a new one.2355  Error getStream(AsyncInfoWrapperTy &AsyncInfoWrapper,2356                  AMDGPUStreamTy *&Stream) {2357    auto WrapperStream =2358        AsyncInfoWrapper.getOrInitQueue<AMDGPUStreamTy *>(AMDGPUStreamManager);2359    if (!WrapperStream)2360      return WrapperStream.takeError();2361    Stream = *WrapperStream;2362    return Plugin::success();2363  }2364 2365  /// Load the binary image into the device and allocate an image object.2366  Expected<DeviceImageTy *>2367  loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,2368                 int32_t ImageId) override {2369    // Allocate and initialize the image object.2370    AMDGPUDeviceImageTy *AMDImage = Plugin.allocate<AMDGPUDeviceImageTy>();2371    new (AMDImage) AMDGPUDeviceImageTy(ImageId, *this, std::move(TgtImage));2372 2373    // Load the HSA executable.2374    if (Error Err = AMDImage->loadExecutable(*this))2375      return std::move(Err);2376 2377    return AMDImage;2378  }2379 2380  /// Allocate memory on the device or related to the device.2381  Expected<void *> allocate(size_t Size, void *, TargetAllocTy Kind) override;2382 2383  /// Deallocate memory on the device or related to the device.2384  Error free(void *TgtPtr, TargetAllocTy Kind) override {2385    if (TgtPtr == nullptr)2386      return Plugin::success();2387 2388    AMDGPUMemoryPoolTy *MemoryPool = nullptr;2389    switch (Kind) {2390    case TARGET_ALLOC_DEFAULT:2391    case TARGET_ALLOC_DEVICE:2392      MemoryPool = CoarseGrainedMemoryPools[0];2393      break;2394    case TARGET_ALLOC_HOST:2395      MemoryPool = &HostDevice.getFineGrainedMemoryPool();2396      break;2397    case TARGET_ALLOC_SHARED:2398      MemoryPool = &HostDevice.getFineGrainedMemoryPool();2399      break;2400    }2401 2402    if (!MemoryPool)2403      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,2404                           "no memory pool for the specified allocation kind");2405 2406    if (auto Err = MemoryPool->deallocate(TgtPtr))2407      return Err;2408 2409    return Plugin::success();2410  }2411 2412  /// Synchronize current thread with the pending operations on the async info.2413  Error synchronizeImpl(__tgt_async_info &AsyncInfo,2414                        bool ReleaseQueue) override {2415    AMDGPUStreamTy *Stream =2416        reinterpret_cast<AMDGPUStreamTy *>(AsyncInfo.Queue);2417    assert(Stream && "Invalid stream");2418 2419    if (auto Err = Stream->synchronize())2420      return Err;2421 2422    // Once the stream is synchronized, return it to stream pool and reset2423    // AsyncInfo. This is to make sure the synchronization only works for its2424    // own tasks.2425    if (ReleaseQueue) {2426      AsyncInfo.Queue = nullptr;2427      return AMDGPUStreamManager.returnResource(Stream);2428    }2429    return Plugin::success();2430  }2431 2432  /// Query for the completion of the pending operations on the async info.2433  Error queryAsyncImpl(__tgt_async_info &AsyncInfo) override {2434    AMDGPUStreamTy *Stream =2435        reinterpret_cast<AMDGPUStreamTy *>(AsyncInfo.Queue);2436    assert(Stream && "Invalid stream");2437 2438    auto CompletedOrErr = Stream->query();2439    if (!CompletedOrErr)2440      return CompletedOrErr.takeError();2441 2442    // Return if it the stream did not complete yet.2443    if (!(*CompletedOrErr))2444      return Plugin::success();2445 2446    // Once the stream is completed, return it to stream pool and reset2447    // AsyncInfo. This is to make sure the synchronization only works for its2448    // own tasks.2449    AsyncInfo.Queue = nullptr;2450    return AMDGPUStreamManager.returnResource(Stream);2451  }2452 2453  /// Pin the host buffer and return the device pointer that should be used for2454  /// device transfers.2455  Expected<void *> dataLockImpl(void *HstPtr, int64_t Size) override {2456    void *PinnedPtr = nullptr;2457 2458    hsa_status_t Status =2459        hsa_amd_memory_lock(HstPtr, Size, nullptr, 0, &PinnedPtr);2460    if (auto Err = Plugin::check(Status, "error in hsa_amd_memory_lock: %s\n"))2461      return std::move(Err);2462 2463    return PinnedPtr;2464  }2465 2466  /// Unpin the host buffer.2467  Error dataUnlockImpl(void *HstPtr) override {2468    hsa_status_t Status = hsa_amd_memory_unlock(HstPtr);2469    return Plugin::check(Status, "error in hsa_amd_memory_unlock: %s\n");2470  }2471 2472  /// Check through the HSA runtime whether the \p HstPtr buffer is pinned.2473  Expected<bool> isPinnedPtrImpl(void *HstPtr, void *&BaseHstPtr,2474                                 void *&BaseDevAccessiblePtr,2475                                 size_t &BaseSize) const override {2476    hsa_amd_pointer_info_t Info;2477    Info.size = sizeof(hsa_amd_pointer_info_t);2478 2479    hsa_status_t Status = hsa_amd_pointer_info(2480        HstPtr, &Info, /*Allocator=*/nullptr, /*num_agents_accessible=*/nullptr,2481        /*accessible=*/nullptr);2482    if (auto Err = Plugin::check(Status, "error in hsa_amd_pointer_info: %s"))2483      return std::move(Err);2484 2485    // The buffer may be locked or allocated through HSA allocators. Assume that2486    // the buffer is host pinned if the runtime reports a HSA type.2487    if (Info.type != HSA_EXT_POINTER_TYPE_LOCKED &&2488        Info.type != HSA_EXT_POINTER_TYPE_HSA)2489      return false;2490 2491    assert(Info.hostBaseAddress && "Invalid host pinned address");2492    assert(Info.agentBaseAddress && "Invalid agent pinned address");2493    assert(Info.sizeInBytes > 0 && "Invalid pinned allocation size");2494 2495    // Save the allocation info in the output parameters.2496    BaseHstPtr = Info.hostBaseAddress;2497    BaseDevAccessiblePtr = Info.agentBaseAddress;2498    BaseSize = Info.sizeInBytes;2499 2500    return true;2501  }2502 2503  /// Submit data to the device (host to device transfer).2504  Error dataSubmitImpl(void *TgtPtr, const void *HstPtr, int64_t Size,2505                       AsyncInfoWrapperTy &AsyncInfoWrapper) override {2506    AMDGPUStreamTy *Stream = nullptr;2507    void *PinnedPtr = nullptr;2508 2509    // Use one-step asynchronous operation when host memory is already pinned.2510    if (void *PinnedPtr =2511            PinnedAllocs.getDeviceAccessiblePtrFromPinnedBuffer(HstPtr)) {2512      if (auto Err = getStream(AsyncInfoWrapper, Stream))2513        return Err;2514      return Stream->pushPinnedMemoryCopyAsync(TgtPtr, PinnedPtr, Size);2515    }2516 2517    // For large transfers use synchronous behavior.2518    if (Size >= OMPX_MaxAsyncCopyBytes) {2519      if (AsyncInfoWrapper.hasQueue())2520        if (auto Err = synchronize(AsyncInfoWrapper))2521          return Err;2522 2523      hsa_status_t Status;2524      Status = hsa_amd_memory_lock(const_cast<void *>(HstPtr), Size, nullptr, 0,2525                                   &PinnedPtr);2526      if (auto Err =2527              Plugin::check(Status, "error in hsa_amd_memory_lock: %s\n"))2528        return Err;2529 2530      AMDGPUSignalTy Signal;2531      if (auto Err = Signal.init())2532        return Err;2533 2534      if (auto Err = hsa_utils::asyncMemCopy(useMultipleSdmaEngines(), TgtPtr,2535                                             Agent, PinnedPtr, Agent, Size, 0,2536                                             nullptr, Signal.get()))2537        return Err;2538 2539      if (auto Err = Signal.wait(getStreamBusyWaitMicroseconds()))2540        return Err;2541 2542      if (auto Err = Signal.deinit())2543        return Err;2544 2545      Status = hsa_amd_memory_unlock(const_cast<void *>(HstPtr));2546      return Plugin::check(Status, "error in hsa_amd_memory_unlock: %s\n");2547    }2548 2549    // Otherwise, use two-step copy with an intermediate pinned host buffer.2550    AMDGPUMemoryManagerTy &PinnedMemoryManager =2551        HostDevice.getPinnedMemoryManager();2552    if (auto Err = PinnedMemoryManager.allocate(Size, &PinnedPtr))2553      return Err;2554 2555    if (auto Err = getStream(AsyncInfoWrapper, Stream))2556      return Err;2557 2558    return Stream->pushMemoryCopyH2DAsync(TgtPtr, HstPtr, PinnedPtr, Size,2559                                          PinnedMemoryManager);2560  }2561 2562  /// Retrieve data from the device (device to host transfer).2563  Error dataRetrieveImpl(void *HstPtr, const void *TgtPtr, int64_t Size,2564                         AsyncInfoWrapperTy &AsyncInfoWrapper) override {2565    AMDGPUStreamTy *Stream = nullptr;2566    void *PinnedPtr = nullptr;2567 2568    // Use one-step asynchronous operation when host memory is already pinned.2569    if (void *PinnedPtr =2570            PinnedAllocs.getDeviceAccessiblePtrFromPinnedBuffer(HstPtr)) {2571      if (auto Err = getStream(AsyncInfoWrapper, Stream))2572        return Err;2573 2574      return Stream->pushPinnedMemoryCopyAsync(PinnedPtr, TgtPtr, Size);2575    }2576 2577    // For large transfers use synchronous behavior.2578    if (Size >= OMPX_MaxAsyncCopyBytes) {2579      if (AsyncInfoWrapper.hasQueue())2580        if (auto Err = synchronize(AsyncInfoWrapper))2581          return Err;2582 2583      hsa_status_t Status;2584      Status = hsa_amd_memory_lock(const_cast<void *>(HstPtr), Size, nullptr, 0,2585                                   &PinnedPtr);2586      if (auto Err =2587              Plugin::check(Status, "error in hsa_amd_memory_lock: %s\n"))2588        return Err;2589 2590      AMDGPUSignalTy Signal;2591      if (auto Err = Signal.init())2592        return Err;2593 2594      if (auto Err = hsa_utils::asyncMemCopy(useMultipleSdmaEngines(),2595                                             PinnedPtr, Agent, TgtPtr, Agent,2596                                             Size, 0, nullptr, Signal.get()))2597        return Err;2598 2599      if (auto Err = Signal.wait(getStreamBusyWaitMicroseconds()))2600        return Err;2601 2602      if (auto Err = Signal.deinit())2603        return Err;2604 2605      Status = hsa_amd_memory_unlock(const_cast<void *>(HstPtr));2606      return Plugin::check(Status, "error in hsa_amd_memory_unlock: %s\n");2607    }2608 2609    // Otherwise, use two-step copy with an intermediate pinned host buffer.2610    AMDGPUMemoryManagerTy &PinnedMemoryManager =2611        HostDevice.getPinnedMemoryManager();2612    if (auto Err = PinnedMemoryManager.allocate(Size, &PinnedPtr))2613      return Err;2614 2615    if (auto Err = getStream(AsyncInfoWrapper, Stream))2616      return Err;2617 2618    return Stream->pushMemoryCopyD2HAsync(HstPtr, TgtPtr, PinnedPtr, Size,2619                                          PinnedMemoryManager);2620  }2621 2622  /// Exchange data between two devices within the plugin.2623  Error dataExchangeImpl(const void *SrcPtr, GenericDeviceTy &DstGenericDevice,2624                         void *DstPtr, int64_t Size,2625                         AsyncInfoWrapperTy &AsyncInfoWrapper) override {2626    AMDGPUDeviceTy &DstDevice = static_cast<AMDGPUDeviceTy &>(DstGenericDevice);2627 2628    // For large transfers use synchronous behavior.2629    if (Size >= OMPX_MaxAsyncCopyBytes) {2630      if (AsyncInfoWrapper.hasQueue())2631        if (auto Err = synchronize(AsyncInfoWrapper))2632          return Err;2633 2634      AMDGPUSignalTy Signal;2635      if (auto Err = Signal.init())2636        return Err;2637 2638      if (auto Err = hsa_utils::asyncMemCopy(2639              useMultipleSdmaEngines(), DstPtr, DstDevice.getAgent(), SrcPtr,2640              getAgent(), (uint64_t)Size, 0, nullptr, Signal.get()))2641        return Err;2642 2643      if (auto Err = Signal.wait(getStreamBusyWaitMicroseconds()))2644        return Err;2645 2646      return Signal.deinit();2647    }2648 2649    AMDGPUStreamTy *Stream = nullptr;2650    if (auto Err = getStream(AsyncInfoWrapper, Stream))2651      return Err;2652    if (Size <= 0)2653      return Plugin::success();2654 2655    return Stream->pushMemoryCopyD2DAsync(DstPtr, DstDevice.getAgent(), SrcPtr,2656                                          getAgent(), (uint64_t)Size);2657  }2658 2659  /// Insert a data fence between previous data operations and the following2660  /// operations. This is a no-op for AMDGPU devices as operations inserted into2661  /// a queue are in-order.2662  Error dataFence(__tgt_async_info *Async) override {2663    return Plugin::success();2664  }2665 2666  Error dataFillImpl(void *TgtPtr, const void *PatternPtr, int64_t PatternSize,2667                     int64_t Size,2668                     AsyncInfoWrapperTy &AsyncInfoWrapper) override {2669    // Fast case, where we can use the 4 byte hsa_amd_memory_fill2670    if (Size % 4 == 0 &&2671        (PatternSize == 4 || PatternSize == 2 || PatternSize == 1)) {2672      uint32_t Pattern;2673      if (PatternSize == 1) {2674        auto *Byte = reinterpret_cast<const uint8_t *>(PatternPtr);2675        Pattern = *Byte | *Byte << 8 | *Byte << 16 | *Byte << 24;2676      } else if (PatternSize == 2) {2677        auto *Word = reinterpret_cast<const uint16_t *>(PatternPtr);2678        Pattern = *Word | (*Word << 16);2679      } else if (PatternSize == 4) {2680        Pattern = *reinterpret_cast<const uint32_t *>(PatternPtr);2681      } else {2682        // Shouldn't be here if the pattern size is outwith those values2683        llvm_unreachable("Invalid pattern size");2684      }2685 2686      if (hasPendingWorkImpl(AsyncInfoWrapper)) {2687        AMDGPUStreamTy *Stream = nullptr;2688        if (auto Err = getStream(AsyncInfoWrapper, Stream))2689          return Err;2690 2691        struct MemFillArgsTy {2692          void *Dst;2693          uint32_t Pattern;2694          int64_t Size;2695        };2696        auto *Args = new MemFillArgsTy{TgtPtr, Pattern, Size / 4};2697        auto Fill = [](void *Data) {2698          MemFillArgsTy *Args = reinterpret_cast<MemFillArgsTy *>(Data);2699          assert(Args && "Invalid arguments");2700 2701          auto Status =2702              hsa_amd_memory_fill(Args->Dst, Args->Pattern, Args->Size);2703          delete Args;2704          auto Err =2705              Plugin::check(Status, "error in hsa_amd_memory_fill: %s\n");2706          if (Err) {2707            FATAL_MESSAGE(1, "error performing async fill: %s",2708                          toString(std::move(Err)).data());2709          }2710        };2711 2712        // hsa_amd_memory_fill doesn't signal completion using a signal, so use2713        // the existing host callback logic to handle that instead2714        return Stream->pushHostCallback(Fill, Args);2715      }2716      // If there is no pending work, do the fill synchronously2717      auto Status = hsa_amd_memory_fill(TgtPtr, Pattern, Size / 4);2718      return Plugin::check(Status, "error in hsa_amd_memory_fill: %s\n");2719    }2720 2721    // Slow case; allocate an appropriate memory size and enqueue copies2722    void *PinnedPtr = nullptr;2723    AMDGPUMemoryManagerTy &PinnedMemoryManager =2724        HostDevice.getPinnedMemoryManager();2725    if (auto Err = PinnedMemoryManager.allocate(Size, &PinnedPtr))2726      return Err;2727 2728    AMDGPUStreamTy *Stream = nullptr;2729    if (auto Err = getStream(AsyncInfoWrapper, Stream))2730      return Err;2731 2732    return Stream->pushMemoryCopyH2DAsync(TgtPtr, PatternPtr, PinnedPtr,2733                                          PatternSize, PinnedMemoryManager,2734                                          Size / PatternSize);2735  }2736 2737  /// Initialize the async info2738  Error initAsyncInfoImpl(AsyncInfoWrapperTy &AsyncInfoWrapper) override {2739    // TODO: Implement this function.2740    return Plugin::success();2741  }2742 2743  interop_spec_t selectInteropPreference(int32_t InteropType,2744                                         int32_t NumPrefers,2745                                         interop_spec_t *Prefers) override {2746    // TODO: update once targetsync is supported2747    if (InteropType == kmp_interop_type_target)2748      return interop_spec_t{tgt_fr_hsa, {false, 0}, 0};2749    return interop_spec_t{tgt_fr_none, {false, 0}, 0};2750  }2751 2752  Expected<omp_interop_val_t *>2753  createInterop(int32_t InteropType, interop_spec_t &InteropSpec) override {2754    auto *Ret = new omp_interop_val_t(2755        DeviceId, static_cast<kmp_interop_type_t>(InteropType));2756    Ret->fr_id = tgt_fr_hsa;2757    Ret->vendor_id = omp_vendor_amd;2758 2759    // TODO: implement targetsync support2760 2761    Ret->device_info.Platform = nullptr;2762    Ret->device_info.Device = reinterpret_cast<void *>(Agent.handle);2763    Ret->device_info.Context = nullptr;2764 2765    return Ret;2766  }2767 2768  Error releaseInterop(omp_interop_val_t *Interop) override {2769    if (Interop)2770      delete Interop;2771    return Plugin::success();2772  }2773 2774  Error enqueueHostCallImpl(AMDGPUStreamTy::HostFnType Callback, void *UserData,2775                            AsyncInfoWrapperTy &AsyncInfo) override {2776    AMDGPUStreamTy *Stream = nullptr;2777    if (auto Err = getStream(AsyncInfo, Stream))2778      return Err;2779 2780    return Stream->pushHostCallback(Callback, UserData);2781  };2782 2783  /// Create an event.2784  Error createEventImpl(void **EventPtrStorage) override {2785    AMDGPUEventTy **Event = reinterpret_cast<AMDGPUEventTy **>(EventPtrStorage);2786    return AMDGPUEventManager.getResource(*Event);2787  }2788 2789  /// Destroy a previously created event.2790  Error destroyEventImpl(void *EventPtr) override {2791    AMDGPUEventTy *Event = reinterpret_cast<AMDGPUEventTy *>(EventPtr);2792    return AMDGPUEventManager.returnResource(Event);2793  }2794 2795  /// Record the event.2796  Error recordEventImpl(void *EventPtr,2797                        AsyncInfoWrapperTy &AsyncInfoWrapper) override {2798    AMDGPUEventTy *Event = reinterpret_cast<AMDGPUEventTy *>(EventPtr);2799    assert(Event && "Invalid event");2800 2801    AMDGPUStreamTy *Stream = nullptr;2802    if (auto Err = getStream(AsyncInfoWrapper, Stream))2803      return Err;2804 2805    return Event->record(*Stream);2806  }2807 2808  /// Make the stream wait on the event.2809  Error waitEventImpl(void *EventPtr,2810                      AsyncInfoWrapperTy &AsyncInfoWrapper) override {2811    AMDGPUEventTy *Event = reinterpret_cast<AMDGPUEventTy *>(EventPtr);2812 2813    AMDGPUStreamTy *Stream = nullptr;2814    if (auto Err = getStream(AsyncInfoWrapper, Stream))2815      return Err;2816 2817    return Event->wait(*Stream);2818  }2819 2820  Expected<bool> hasPendingWorkImpl(AsyncInfoWrapperTy &AsyncInfo) override {2821    auto *Stream = AsyncInfo.getQueueAs<AMDGPUStreamTy *>();2822    if (!Stream)2823      return false;2824 2825    auto Query = Stream->query();2826    if (Query)2827      return !*Query;2828    return Query.takeError();2829  }2830 2831  Expected<bool> isEventCompleteImpl(void *EventPtr,2832                                     AsyncInfoWrapperTy &AsyncInfo) override {2833    AMDGPUEventTy *Event = reinterpret_cast<AMDGPUEventTy *>(EventPtr);2834    auto *Stream = AsyncInfo.getQueueAs<AMDGPUStreamTy *>();2835    return Stream && Stream->isEventComplete(*Event);2836  }2837 2838  /// Synchronize the current thread with the event.2839  Error syncEventImpl(void *EventPtr) override {2840    AMDGPUEventTy *Event = reinterpret_cast<AMDGPUEventTy *>(EventPtr);2841    return Event->sync();2842  }2843 2844  /// Print information about the device.2845  Expected<InfoTreeNode> obtainInfoImpl() override {2846    char TmpChar[1000];2847    const char *TmpCharPtr = "Unknown";2848    uint16_t Major, Minor;2849    uint32_t TmpUInt, TmpUInt2;2850    uint32_t CacheSize[4];2851    size_t TmpSt;2852    bool TmpBool;2853    uint16_t WorkgrpMaxDim[3];2854    hsa_dim3_t GridMaxDim;2855    hsa_status_t Status, Status2;2856    InfoTreeNode Info;2857 2858    Status = hsa_system_get_info(HSA_SYSTEM_INFO_VERSION_MAJOR, &Major);2859    Status2 = hsa_system_get_info(HSA_SYSTEM_INFO_VERSION_MINOR, &Minor);2860    if (Status == HSA_STATUS_SUCCESS && Status2 == HSA_STATUS_SUCCESS)2861      Info.add("HSA Runtime Version",2862               std::to_string(Major) + "." + std::to_string(Minor), "",2863               DeviceInfo::DRIVER_VERSION);2864 2865    Info.add("HSA OpenMP Device Number", DeviceId);2866 2867    Status = getDeviceAttrRaw(HSA_AMD_AGENT_INFO_PRODUCT_NAME, TmpChar);2868    if (Status == HSA_STATUS_SUCCESS)2869      Info.add("Product Name", TmpChar, "", DeviceInfo::PRODUCT_NAME);2870 2871    Status = getDeviceAttrRaw(HSA_AGENT_INFO_NAME, TmpChar);2872    if (Status == HSA_STATUS_SUCCESS)2873      Info.add("Device Name", TmpChar, "", DeviceInfo::NAME);2874 2875    Status = getDeviceAttrRaw(HSA_AGENT_INFO_VENDOR_NAME, TmpChar);2876    if (Status == HSA_STATUS_SUCCESS)2877      Info.add("Vendor Name", TmpChar, "", DeviceInfo::VENDOR);2878 2879    Info.add("Vendor ID", uint64_t{4130}, "", DeviceInfo::VENDOR_ID);2880 2881    hsa_machine_model_t MachineModel;2882    Status = getDeviceAttrRaw(HSA_AGENT_INFO_MACHINE_MODEL, MachineModel);2883    if (Status == HSA_STATUS_SUCCESS)2884      Info.add("Memory Address Size",2885               uint64_t{MachineModel == HSA_MACHINE_MODEL_SMALL ? 32u : 64u},2886               "bits", DeviceInfo::ADDRESS_BITS);2887 2888    hsa_device_type_t DevType;2889    Status = getDeviceAttrRaw(HSA_AGENT_INFO_DEVICE, DevType);2890    if (Status == HSA_STATUS_SUCCESS) {2891      switch (static_cast<int>(DevType)) {2892      case HSA_DEVICE_TYPE_CPU:2893        TmpCharPtr = "CPU";2894        break;2895      case HSA_DEVICE_TYPE_GPU:2896        TmpCharPtr = "GPU";2897        break;2898      case HSA_DEVICE_TYPE_DSP:2899        TmpCharPtr = "DSP";2900        break;2901      default:2902        TmpCharPtr = "Unknown";2903        break;2904      }2905      Info.add("Device Type", TmpCharPtr);2906    }2907 2908    Status = getDeviceAttrRaw(HSA_AGENT_INFO_QUEUES_MAX, TmpUInt);2909    if (Status == HSA_STATUS_SUCCESS)2910      Info.add("Max Queues", TmpUInt);2911 2912    Status = getDeviceAttrRaw(HSA_AGENT_INFO_QUEUE_MIN_SIZE, TmpUInt);2913    if (Status == HSA_STATUS_SUCCESS)2914      Info.add("Queue Min Size", TmpUInt);2915 2916    Status = getDeviceAttrRaw(HSA_AGENT_INFO_QUEUE_MAX_SIZE, TmpUInt);2917    if (Status == HSA_STATUS_SUCCESS)2918      Info.add("Queue Max Size", TmpUInt);2919 2920    // FIXME: This is deprecated according to HSA documentation. But using2921    // hsa_agent_iterate_caches and hsa_cache_get_info breaks execution during2922    // runtime.2923    Status = getDeviceAttrRaw(HSA_AGENT_INFO_CACHE_SIZE, CacheSize);2924    if (Status == HSA_STATUS_SUCCESS) {2925      auto &Cache = *Info.add("Cache");2926 2927      for (int I = 0; I < 4; I++)2928        if (CacheSize[I])2929          Cache.add("L" + std::to_string(I), CacheSize[I]);2930    }2931 2932    Status = getDeviceAttrRaw(HSA_AMD_AGENT_INFO_CACHELINE_SIZE, TmpUInt);2933    if (Status == HSA_STATUS_SUCCESS)2934      Info.add("Cacheline Size", TmpUInt);2935 2936    Info.add("Max Shared Memory per Work Group", MaxBlockSharedMemSize, "bytes",2937             DeviceInfo::WORK_GROUP_LOCAL_MEM_SIZE);2938 2939    Status = getDeviceAttrRaw(HSA_AMD_AGENT_INFO_MAX_CLOCK_FREQUENCY, TmpUInt);2940    if (Status == HSA_STATUS_SUCCESS)2941      Info.add("Max Clock Freq", TmpUInt, "MHz",2942               DeviceInfo::MAX_CLOCK_FREQUENCY);2943 2944    Status = getDeviceAttrRaw(HSA_AMD_AGENT_INFO_MEMORY_MAX_FREQUENCY, TmpUInt);2945    if (Status == HSA_STATUS_SUCCESS)2946      Info.add("Max Memory Clock Freq", TmpUInt, "MHz",2947               DeviceInfo::MEMORY_CLOCK_RATE);2948 2949    Status = getDeviceAttrRaw(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, TmpUInt);2950    if (Status == HSA_STATUS_SUCCESS)2951      Info.add("Compute Units", TmpUInt, "", DeviceInfo::NUM_COMPUTE_UNITS);2952 2953    Status = getDeviceAttrRaw(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU, TmpUInt);2954    if (Status == HSA_STATUS_SUCCESS)2955      Info.add("SIMD per CU", TmpUInt);2956 2957    Status = getDeviceAttrRaw(HSA_AGENT_INFO_FAST_F16_OPERATION, TmpBool);2958    if (Status == HSA_STATUS_SUCCESS)2959      Info.add("Fast F16 Operation", TmpBool);2960 2961    Status = getDeviceAttrRaw(HSA_AGENT_INFO_WAVEFRONT_SIZE, TmpUInt2);2962    if (Status == HSA_STATUS_SUCCESS)2963      Info.add("Wavefront Size", TmpUInt2);2964 2965    Status = getDeviceAttrRaw(HSA_AGENT_INFO_WORKGROUP_MAX_SIZE, TmpUInt);2966    if (Status == HSA_STATUS_SUCCESS)2967      Info.add("Workgroup Max Size", TmpUInt, "",2968               DeviceInfo::MAX_WORK_GROUP_SIZE);2969 2970    Status = getDeviceAttrRaw(HSA_AGENT_INFO_WORKGROUP_MAX_DIM, WorkgrpMaxDim);2971    if (Status == HSA_STATUS_SUCCESS) {2972      auto &MaxSize =2973          *Info.add("Workgroup Max Size per Dimension", std::monostate{}, "",2974                    DeviceInfo::MAX_WORK_GROUP_SIZE_PER_DIMENSION);2975      MaxSize.add("x", WorkgrpMaxDim[0]);2976      MaxSize.add("y", WorkgrpMaxDim[1]);2977      MaxSize.add("z", WorkgrpMaxDim[2]);2978    }2979 2980    Status = getDeviceAttrRaw(2981        (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU, TmpUInt);2982    if (Status == HSA_STATUS_SUCCESS) {2983      Info.add("Max Waves Per CU", TmpUInt);2984      Info.add("Max Work-item Per CU", TmpUInt * TmpUInt2);2985    }2986 2987    Status = getDeviceAttrRaw(HSA_AGENT_INFO_GRID_MAX_SIZE, TmpUInt);2988    if (Status == HSA_STATUS_SUCCESS)2989      Info.add("Grid Max Size", TmpUInt, "", DeviceInfo::MAX_WORK_SIZE);2990 2991    Status = getDeviceAttrRaw(HSA_AGENT_INFO_GRID_MAX_DIM, GridMaxDim);2992    if (Status == HSA_STATUS_SUCCESS) {2993      auto &MaxDim = *Info.add("Grid Max Size per Dimension", std::monostate{},2994                               "", DeviceInfo::MAX_WORK_SIZE_PER_DIMENSION);2995      MaxDim.add("x", GridMaxDim.x);2996      MaxDim.add("y", GridMaxDim.y);2997      MaxDim.add("z", GridMaxDim.z);2998    }2999 3000    Status = getDeviceAttrRaw(HSA_AGENT_INFO_FBARRIER_MAX_SIZE, TmpUInt);3001    if (Status == HSA_STATUS_SUCCESS)3002      Info.add("Max fbarriers/Workgrp", TmpUInt);3003 3004    auto &RootPool = *Info.add("Memory Pools");3005    for (AMDGPUMemoryPoolTy *Pool : AllMemoryPools) {3006      std::string TmpStr, TmpStr2;3007 3008      if (Pool->isGlobal())3009        TmpStr = "Global";3010      else if (Pool->isReadOnly())3011        TmpStr = "ReadOnly";3012      else if (Pool->isPrivate())3013        TmpStr = "Private";3014      else if (Pool->isGroup())3015        TmpStr = "Group";3016      else3017        TmpStr = "Unknown";3018 3019      auto &PoolNode = *RootPool.add(std::string("Pool ") + TmpStr);3020 3021      if (Pool->isGlobal()) {3022        if (Pool->isFineGrained())3023          TmpStr2 += "Fine Grained ";3024        if (Pool->isCoarseGrained())3025          TmpStr2 += "Coarse Grained ";3026        if (Pool->supportsKernelArgs())3027          TmpStr2 += "Kernarg ";3028 3029        PoolNode.add("Flags", TmpStr2);3030      }3031 3032      Status = Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_SIZE, TmpSt);3033      if (Status == HSA_STATUS_SUCCESS)3034        PoolNode.add(3035            "Size", TmpSt, "bytes",3036            (Pool->isGlobal() && Pool->isCoarseGrained())3037                ? std::optional<DeviceInfo>{DeviceInfo::GLOBAL_MEM_SIZE}3038                : std::nullopt);3039 3040      Status = Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED,3041                                TmpBool);3042      if (Status == HSA_STATUS_SUCCESS)3043        PoolNode.add("Allocatable", TmpBool);3044 3045      Status = Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,3046                                TmpSt);3047      if (Status == HSA_STATUS_SUCCESS)3048        PoolNode.add("Runtime Alloc Granule", TmpSt, "bytes");3049 3050      Status = Pool->getAttrRaw(3051          HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALIGNMENT, TmpSt);3052      if (Status == HSA_STATUS_SUCCESS)3053        PoolNode.add("Runtime Alloc Alignment", TmpSt, "bytes");3054 3055      Status =3056          Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_ACCESSIBLE_BY_ALL, TmpBool);3057      if (Status == HSA_STATUS_SUCCESS)3058        PoolNode.add("Accessible by all", TmpBool);3059    }3060 3061    auto &ISAs = *Info.add("ISAs");3062    auto Err = hsa_utils::iterateAgentISAs(getAgent(), [&](hsa_isa_t ISA) {3063      Status = hsa_isa_get_info_alt(ISA, HSA_ISA_INFO_NAME, TmpChar);3064      if (Status == HSA_STATUS_SUCCESS)3065        ISAs.add("Name", TmpChar);3066 3067      return Status;3068    });3069 3070    // Silently consume the error.3071    if (Err)3072      consumeError(std::move(Err));3073 3074    return Info;3075  }3076 3077  /// Returns true if auto zero-copy the best configuration for the current3078  /// arch.3079  /// On AMDGPUs, automatic zero-copy is turned on3080  /// when running on an APU with XNACK (unified memory) support3081  /// enabled. On discrete GPUs, automatic zero-copy is triggered3082  /// if the user sets the environment variable OMPX_APU_MAPS=13083  /// and if XNACK is enabled. The rationale is that zero-copy3084  /// is the best configuration (performance, memory footprint) on APUs,3085  /// while it is often not the best on discrete GPUs.3086  /// XNACK can be enabled with a kernel boot parameter or with3087  /// the HSA_XNACK environment variable.3088  bool useAutoZeroCopyImpl() override {3089    return ((IsAPU || OMPX_ApuMaps) && IsXnackEnabled);3090  }3091 3092  Expected<bool> isAccessiblePtrImpl(const void *Ptr, size_t Size) override {3093    hsa_amd_pointer_info_t Info;3094    Info.size = sizeof(hsa_amd_pointer_info_t);3095 3096    hsa_agent_t *Agents = nullptr;3097    uint32_t Count = 0;3098    hsa_status_t Status =3099        hsa_amd_pointer_info(Ptr, &Info, malloc, &Count, &Agents);3100 3101    if (auto Err = Plugin::check(Status, "error in hsa_amd_pointer_info: %s"))3102      return std::move(Err);3103 3104    // Checks if the pointer is known by HSA and accessible by the device3105    for (uint32_t i = 0; i < Count; i++) {3106      if (Agents[i].handle == getAgent().handle)3107        return Info.sizeInBytes >= Size;3108    }3109 3110    // If the pointer is unknown to HSA it's assumed a host pointer3111    // in that case the device can access it on unified memory support is3112    // enabled3113    return IsXnackEnabled;3114  }3115 3116  /// Getters and setters for stack and heap sizes.3117  Error getDeviceStackSize(uint64_t &Value) override {3118    Value = StackSize;3119    return Plugin::success();3120  }3121  Error setDeviceStackSize(uint64_t Value) override {3122    StackSize = Value;3123    return Plugin::success();3124  }3125  Error getDeviceMemorySize(uint64_t &Value) override {3126    for (AMDGPUMemoryPoolTy *Pool : AllMemoryPools) {3127      if (Pool->isGlobal()) {3128        hsa_status_t Status =3129            Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_SIZE, Value);3130        return Plugin::check(Status, "error in getting device memory size: %s");3131      }3132    }3133    return Plugin::error(ErrorCode::UNSUPPORTED,3134                         "getDeviceMemorySize:: no global pool");3135  }3136 3137  /// AMDGPU-specific function to get device attributes.3138  template <typename Ty> Error getDeviceAttr(uint32_t Kind, Ty &Value) {3139    hsa_status_t Status =3140        hsa_agent_get_info(Agent, (hsa_agent_info_t)Kind, &Value);3141    return Plugin::check(Status, "Error in hsa_agent_get_info: %s");3142  }3143 3144  template <typename Ty>3145  hsa_status_t getDeviceAttrRaw(uint32_t Kind, Ty &Value) {3146    return hsa_agent_get_info(Agent, (hsa_agent_info_t)Kind, &Value);3147  }3148 3149  /// Get the device agent.3150  hsa_agent_t getAgent() const override { return Agent; }3151 3152  /// Get the signal manager.3153  AMDGPUSignalManagerTy &getSignalManager() { return AMDGPUSignalManager; }3154 3155  /// Retrieve and construct all memory pools of the device agent.3156  Error retrieveAllMemoryPools() override {3157    // Iterate through the available pools of the device agent.3158    return hsa_utils::iterateAgentMemoryPools(3159        Agent, [&](hsa_amd_memory_pool_t HSAMemoryPool) {3160          AMDGPUMemoryPoolTy *MemoryPool =3161              Plugin.allocate<AMDGPUMemoryPoolTy>();3162          new (MemoryPool) AMDGPUMemoryPoolTy(HSAMemoryPool);3163          AllMemoryPools.push_back(MemoryPool);3164          return HSA_STATUS_SUCCESS;3165        });3166  }3167 3168  bool useMultipleSdmaEngines() const { return OMPX_UseMultipleSdmaEngines; }3169 3170private:3171  using AMDGPUEventRef = AMDGPUResourceRef<AMDGPUEventTy>;3172  using AMDGPUEventManagerTy = GenericDeviceResourceManagerTy<AMDGPUEventRef>;3173 3174  /// Common method to invoke a single threaded constructor or destructor3175  /// kernel by name.3176  Error callGlobalCtorDtorCommon(GenericPluginTy &Plugin, DeviceImageTy &Image,3177                                 bool IsCtor) {3178    const char *KernelName =3179        IsCtor ? "amdgcn.device.init" : "amdgcn.device.fini";3180    // Perform a quick check for the named kernel in the image. The kernel3181    // should be created by the 'amdgpu-lower-ctor-dtor' pass.3182    GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler();3183    if (!Handler.isSymbolInImage(*this, Image, KernelName))3184      return Plugin::success();3185 3186    // Allocate and construct the AMDGPU kernel.3187    AMDGPUKernelTy AMDGPUKernel(KernelName);3188    if (auto Err = AMDGPUKernel.init(*this, Image))3189      return Err;3190 3191    AsyncInfoWrapperTy AsyncInfoWrapper(*this, nullptr);3192 3193    KernelArgsTy KernelArgs = {};3194    uint32_t NumBlocksAndThreads[3] = {1u, 1u, 1u};3195    if (auto Err = AMDGPUKernel.launchImpl(3196            *this, NumBlocksAndThreads, NumBlocksAndThreads, KernelArgs,3197            KernelLaunchParamsTy{}, AsyncInfoWrapper))3198      return Err;3199 3200    Error Err = Plugin::success();3201    AsyncInfoWrapper.finalize(Err);3202 3203    return Err;3204  }3205 3206  /// Detect if current architecture is an APU.3207  Error checkIfAPU() {3208    // TODO: replace with ROCr API once it becomes available.3209    llvm::StringRef StrGfxName(ComputeUnitKind);3210    bool MayBeAPU = llvm::StringSwitch<bool>(StrGfxName)3211                        .Case("gfx942", true)3212                        .Default(false);3213    if (!MayBeAPU)3214      return Plugin::success();3215 3216    // can be MI300A or MI300X3217    uint32_t ChipID = 0;3218    if (auto Err = getDeviceAttr(HSA_AMD_AGENT_INFO_CHIP_ID, ChipID))3219      return Err;3220 3221    if (!(ChipID & 0x1)) {3222      IsAPU = true;3223      return Plugin::success();3224    }3225    return Plugin::success();3226  }3227 3228  bool checkIfCoarseGrainMemoryNearOrAbove64GB() {3229    for (AMDGPUMemoryPoolTy *Pool : AllMemoryPools) {3230      if (!Pool->isGlobal() || !Pool->isCoarseGrained())3231        continue;3232      uint64_t Value;3233      hsa_status_t Status =3234          Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_SIZE, Value);3235      if (Status != HSA_STATUS_SUCCESS)3236        continue;3237      constexpr uint64_t Almost64Gig = 0xFF0000000;3238      if (Value >= Almost64Gig)3239        return true;3240    }3241    return false; // CoarseGrain pool w/ 64GB or more capacity not found3242  }3243 3244  size_t getMemoryManagerSizeThreshold() override {3245    // Targeting high memory capacity GPUs such as3246    // data center GPUs.3247    if (checkIfCoarseGrainMemoryNearOrAbove64GB()) {3248      // Set GenericDeviceTy::MemoryManager's Threshold to 3GiB,3249      // if threshold is not already set by ENV var3250      // LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD.3251      // This MemoryManager is used for omp_target_alloc(), OpenMP3252      // (non-usm) map clause, etc.3253      //3254      // Ideally, this kind of pooling is best performed at3255      // a common level (e.g, user side of HSA) between OpenMP and HIP3256      // but that feature does not exist (yet).3257      return 3ul * 1024 * 1024 * 1024 /* 3 GiB */;3258    }3259    return 0;3260  }3261 3262  /// Envar for controlling the number of HSA queues per device. High number of3263  /// queues may degrade performance.3264  UInt32Envar OMPX_NumQueues;3265 3266  /// Envar for controlling the size of each HSA queue. The size is the number3267  /// of HSA packets a queue is expected to hold. It is also the number of HSA3268  /// packets that can be pushed into each queue without waiting the driver to3269  /// process them.3270  UInt32Envar OMPX_QueueSize;3271 3272  /// Envar for controlling the default number of teams relative to the number3273  /// of compute units (CUs) the device has:3274  ///   #default_teams = OMPX_DefaultTeamsPerCU * #CUs.3275  UInt32Envar OMPX_DefaultTeamsPerCU;3276 3277  /// Envar specifying the maximum size in bytes where the memory copies are3278  /// asynchronous operations. Up to this transfer size, the memory copies are3279  /// asynchronous operations pushed to the corresponding stream. For larger3280  /// transfers, they are synchronous transfers.3281  UInt32Envar OMPX_MaxAsyncCopyBytes;3282 3283  /// Envar controlling the initial number of HSA signals per device. There is3284  /// one manager of signals per device managing several pre-allocated signals.3285  /// These signals are mainly used by AMDGPU streams. If needed, more signals3286  /// will be created.3287  UInt32Envar OMPX_InitialNumSignals;3288 3289  /// Environment variables to set the time to wait in active state before3290  /// switching to blocked state. The default 2000000 busywaits for 2 seconds3291  /// before going into a blocking HSA wait state. The unit for these variables3292  /// are microseconds.3293  UInt32Envar OMPX_StreamBusyWait;3294 3295  /// Use ROCm 5.7 interface for multiple SDMA engines3296  BoolEnvar OMPX_UseMultipleSdmaEngines;3297 3298  /// Value of OMPX_APU_MAPS env var used to force3299  /// automatic zero-copy behavior on non-APU GPUs.3300  BoolEnvar OMPX_ApuMaps;3301 3302  /// Stream manager for AMDGPU streams.3303  AMDGPUStreamManagerTy AMDGPUStreamManager;3304 3305  /// Event manager for AMDGPU events.3306  AMDGPUEventManagerTy AMDGPUEventManager;3307 3308  /// Signal manager for AMDGPU signals.3309  AMDGPUSignalManagerTy AMDGPUSignalManager;3310 3311  /// The agent handler corresponding to the device.3312  hsa_agent_t Agent;3313 3314  /// The GPU architecture.3315  std::string ComputeUnitKind;3316 3317  /// The frequency of the steady clock inside the device.3318  uint64_t ClockFrequency;3319 3320  /// The total number of concurrent work items that can be running on the GPU.3321  uint64_t HardwareParallelism;3322 3323  /// Reference to the host device.3324  AMDHostDeviceTy &HostDevice;3325 3326  /// The current size of the stack that will be used in cases where it could3327  /// not be statically determined.3328  uint64_t StackSize = 16 * 1024 /* 16 KB */;3329 3330  /// Is the plugin associated with an APU?3331  bool IsAPU = false;3332 3333  /// True is the system is configured with XNACK-Enabled.3334  /// False otherwise.3335  bool IsXnackEnabled = false;3336};3337 3338Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) {3339  hsa_code_object_reader_t Reader;3340  hsa_status_t Status =3341      hsa_code_object_reader_create_from_memory(getStart(), getSize(), &Reader);3342  if (auto Err = Plugin::check(3343          Status, "error in hsa_code_object_reader_create_from_memory: %s"))3344    return Err;3345 3346  Status = hsa_executable_create_alt(3347      HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, "", &Executable);3348  if (auto Err =3349          Plugin::check(Status, "error in hsa_executable_create_alt: %s"))3350    return Err;3351 3352  hsa_loaded_code_object_t Object;3353  Status = hsa_executable_load_agent_code_object(Executable, Device.getAgent(),3354                                                 Reader, "", &Object);3355  if (auto Err = Plugin::check(3356          Status, "error in hsa_executable_load_agent_code_object: %s"))3357    return Err;3358 3359  Status = hsa_executable_freeze(Executable, "");3360  if (auto Err = Plugin::check(Status, "error in hsa_executable_freeze: %s"))3361    return Err;3362 3363  uint32_t Result;3364  Status = hsa_executable_validate(Executable, &Result);3365  if (auto Err = Plugin::check(Status, "error in hsa_executable_validate: %s"))3366    return Err;3367 3368  if (Result)3369    return Plugin::error(ErrorCode::INVALID_BINARY,3370                         "loaded HSA executable does not validate");3371 3372  Status = hsa_code_object_reader_destroy(Reader);3373  if (auto Err =3374          Plugin::check(Status, "error in hsa_code_object_reader_destroy: %s"))3375    return Err;3376 3377  if (auto Err = hsa_utils::readAMDGPUMetaDataFromImage(3378          getMemoryBuffer(), KernelInfoMap, ELFABIVersion))3379    return Err;3380 3381  return Plugin::success();3382}3383 3384Expected<hsa_executable_symbol_t>3385AMDGPUDeviceImageTy::findDeviceSymbol(GenericDeviceTy &Device,3386                                      StringRef SymbolName) const {3387 3388  AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(Device);3389  hsa_agent_t Agent = AMDGPUDevice.getAgent();3390 3391  hsa_executable_symbol_t Symbol;3392  hsa_status_t Status = hsa_executable_get_symbol_by_name(3393      Executable, SymbolName.data(), &Agent, &Symbol);3394  if (auto Err = Plugin::check(3395          Status, "error in hsa_executable_get_symbol_by_name(%s): %s",3396          SymbolName.data()))3397    return std::move(Err);3398 3399  return Symbol;3400}3401 3402template <typename ResourceTy>3403Error AMDGPUResourceRef<ResourceTy>::create(GenericDeviceTy &Device) {3404  if (Resource)3405    return Plugin::error(ErrorCode::INVALID_ARGUMENT,3406                         "creating an existing resource");3407 3408  AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(Device);3409 3410  Resource = new ResourceTy(AMDGPUDevice);3411 3412  return Resource->init();3413}3414 3415AMDGPUStreamTy::AMDGPUStreamTy(AMDGPUDeviceTy &Device)3416    : Agent(Device.getAgent()), Queue(nullptr),3417      SignalManager(Device.getSignalManager()), Device(Device),3418      // Initialize the std::deque with some empty positions.3419      Slots(32), NextSlot(0), SyncCycle(0),3420      StreamBusyWaitMicroseconds(Device.getStreamBusyWaitMicroseconds()),3421      UseMultipleSdmaEngines(Device.useMultipleSdmaEngines()) {}3422 3423/// Class implementing the AMDGPU-specific functionalities of the global3424/// handler.3425struct AMDGPUGlobalHandlerTy final : public GenericGlobalHandlerTy {3426  /// Get the metadata of a global from the device. The name and size of the3427  /// global is read from DeviceGlobal and the address of the global is written3428  /// to DeviceGlobal.3429  Error getGlobalMetadataFromDevice(GenericDeviceTy &Device,3430                                    DeviceImageTy &Image,3431                                    GlobalTy &DeviceGlobal) override {3432    AMDGPUDeviceImageTy &AMDImage = static_cast<AMDGPUDeviceImageTy &>(Image);3433 3434    // Find the symbol on the device executable.3435    auto SymbolOrErr =3436        AMDImage.findDeviceSymbol(Device, DeviceGlobal.getName());3437    if (!SymbolOrErr)3438      return SymbolOrErr.takeError();3439 3440    hsa_executable_symbol_t Symbol = *SymbolOrErr;3441    hsa_symbol_kind_t SymbolType;3442    hsa_status_t Status;3443    uint64_t SymbolAddr;3444    uint32_t SymbolSize;3445 3446    // Retrieve the type, address and size of the symbol.3447    std::pair<hsa_executable_symbol_info_t, void *> RequiredInfos[] = {3448        {HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &SymbolType},3449        {HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &SymbolAddr},3450        {HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &SymbolSize}};3451 3452    for (auto &Info : RequiredInfos) {3453      Status = hsa_executable_symbol_get_info(Symbol, Info.first, Info.second);3454      if (auto Err = Plugin::check(3455              Status, "error in hsa_executable_symbol_get_info: %s"))3456        return Err;3457    }3458 3459    // Check the size of the symbol.3460    if (DeviceGlobal.getSize() && SymbolSize != DeviceGlobal.getSize())3461      return Plugin::error(3462          ErrorCode::INVALID_BINARY,3463          "failed to load global '%s' due to size mismatch (%zu != %zu)",3464          DeviceGlobal.getName().data(), SymbolSize,3465          (size_t)DeviceGlobal.getSize());3466 3467    // Store the symbol address and size on the device global metadata.3468    DeviceGlobal.setPtr(reinterpret_cast<void *>(SymbolAddr));3469    DeviceGlobal.setSize(SymbolSize);3470 3471    return Plugin::success();3472  }3473};3474 3475/// Class implementing the AMDGPU-specific functionalities of the plugin.3476struct AMDGPUPluginTy final : public GenericPluginTy {3477  /// Create an AMDGPU plugin and initialize the AMDGPU driver.3478  AMDGPUPluginTy()3479      : GenericPluginTy(getTripleArch()), Initialized(false),3480        HostDevice(nullptr) {}3481 3482  /// This class should not be copied.3483  AMDGPUPluginTy(const AMDGPUPluginTy &) = delete;3484  AMDGPUPluginTy(AMDGPUPluginTy &&) = delete;3485 3486  /// Initialize the plugin and return the number of devices.3487  Expected<int32_t> initImpl() override {3488    hsa_status_t Status = hsa_init();3489    if (Status != HSA_STATUS_SUCCESS) {3490      // Cannot call hsa_success_string.3491      DP("Failed to initialize AMDGPU's HSA library\n");3492      return 0;3493    }3494 3495    // The initialization of HSA was successful. It should be safe to call3496    // HSA functions from now on, e.g., hsa_shut_down.3497    Initialized = true;3498 3499    // Register event handler to detect memory errors on the devices.3500    Status = hsa_amd_register_system_event_handler(eventHandler, this);3501    if (auto Err = Plugin::check(3502            Status, "error in hsa_amd_register_system_event_handler: %s"))3503      return std::move(Err);3504 3505    // List of host (CPU) agents.3506    llvm::SmallVector<hsa_agent_t> HostAgents;3507 3508    // Count the number of available agents.3509    auto Err = hsa_utils::iterateAgents([&](hsa_agent_t Agent) {3510      // Get the device type of the agent.3511      hsa_device_type_t DeviceType;3512      hsa_status_t Status =3513          hsa_agent_get_info(Agent, HSA_AGENT_INFO_DEVICE, &DeviceType);3514      if (Status != HSA_STATUS_SUCCESS)3515        return Status;3516 3517      // Classify the agents into kernel (GPU) and host (CPU) kernels.3518      if (DeviceType == HSA_DEVICE_TYPE_GPU) {3519        // Ensure that the GPU agent supports kernel dispatch packets.3520        hsa_agent_feature_t Features;3521        Status = hsa_agent_get_info(Agent, HSA_AGENT_INFO_FEATURE, &Features);3522        if (Features & HSA_AGENT_FEATURE_KERNEL_DISPATCH)3523          KernelAgents.push_back(Agent);3524      } else if (DeviceType == HSA_DEVICE_TYPE_CPU) {3525        HostAgents.push_back(Agent);3526      }3527      return HSA_STATUS_SUCCESS;3528    });3529 3530    if (Err)3531      return std::move(Err);3532 3533    int32_t NumDevices = KernelAgents.size();3534    if (NumDevices == 0) {3535      // Do not initialize if there are no devices.3536      DP("There are no devices supporting AMDGPU.\n");3537      return 0;3538    }3539 3540    // There are kernel agents but there is no host agent. That should be3541    // treated as an error.3542    if (HostAgents.empty())3543      return Plugin::error(ErrorCode::BACKEND_FAILURE, "no AMDGPU host agents");3544 3545    // Initialize the host device using host agents.3546    HostDevice = allocate<AMDHostDeviceTy>();3547    new (HostDevice) AMDHostDeviceTy(*this, HostAgents);3548 3549    // Setup the memory pools of available for the host.3550    if (auto Err = HostDevice->init())3551      return std::move(Err);3552 3553    return NumDevices;3554  }3555 3556  /// Deinitialize the plugin.3557  Error deinitImpl() override {3558    // The HSA runtime was not initialized, so nothing from the plugin was3559    // actually initialized.3560    if (!Initialized)3561      return Plugin::success();3562 3563    if (HostDevice)3564      if (auto Err = HostDevice->deinit())3565        return Err;3566 3567    // Finalize the HSA runtime.3568    hsa_status_t Status = hsa_shut_down();3569    return Plugin::check(Status, "error in hsa_shut_down: %s");3570  }3571 3572  /// Creates an AMDGPU device.3573  GenericDeviceTy *createDevice(GenericPluginTy &Plugin, int32_t DeviceId,3574                                int32_t NumDevices) override {3575    return new AMDGPUDeviceTy(Plugin, DeviceId, NumDevices, getHostDevice(),3576                              getKernelAgent(DeviceId));3577  }3578 3579  /// Creates an AMDGPU global handler.3580  GenericGlobalHandlerTy *createGlobalHandler() override {3581    return new AMDGPUGlobalHandlerTy();3582  }3583 3584  Triple::ArchType getTripleArch() const override { return Triple::amdgcn; }3585 3586  const char *getName() const override { return GETNAME(TARGET_NAME); }3587 3588  /// Get the ELF code for recognizing the compatible image binary.3589  uint16_t getMagicElfBits() const override { return ELF::EM_AMDGPU; }3590 3591  /// Check whether the image is compatible with an AMDGPU device.3592  Expected<bool> isELFCompatible(uint32_t DeviceId,3593                                 StringRef Image) const override {3594    // Get the associated architecture and flags from the ELF.3595    auto ElfOrErr = ELF64LEObjectFile::create(3596        MemoryBufferRef(Image, /*Identifier=*/""), /*InitContent=*/false);3597    if (!ElfOrErr)3598      return ElfOrErr.takeError();3599    std::optional<StringRef> Processor = ElfOrErr->tryGetCPUName();3600    if (!Processor)3601      return false;3602 3603    SmallVector<SmallString<32>> Targets;3604    if (auto Err = hsa_utils::getTargetTripleAndFeatures(3605            getKernelAgent(DeviceId), Targets))3606      return Err;3607    for (auto &Target : Targets)3608      if (offloading::amdgpu::isImageCompatibleWithEnv(3609              Processor ? *Processor : "", ElfOrErr->getPlatformFlags(),3610              Target.str()))3611        return true;3612    return false;3613  }3614 3615  bool isDataExchangable(int32_t SrcDeviceId, int32_t DstDeviceId) override {3616    return true;3617  }3618 3619  /// Get the host device instance.3620  AMDHostDeviceTy &getHostDevice() {3621    assert(HostDevice && "Host device not initialized");3622    return *HostDevice;3623  }3624 3625  /// Get the kernel agent with the corresponding agent id.3626  hsa_agent_t getKernelAgent(int32_t AgentId) const {3627    assert((uint32_t)AgentId < KernelAgents.size() && "Invalid agent id");3628    return KernelAgents[AgentId];3629  }3630 3631  /// Get the list of the available kernel agents.3632  const llvm::SmallVector<hsa_agent_t> &getKernelAgents() const {3633    return KernelAgents;3634  }3635 3636private:3637  /// Event handler that will be called by ROCr if an event is detected.3638  static hsa_status_t eventHandler(const hsa_amd_event_t *Event,3639                                   void *PluginPtr) {3640    if (Event->event_type != HSA_AMD_GPU_MEMORY_FAULT_EVENT)3641      return HSA_STATUS_SUCCESS;3642 3643    SmallVector<std::string> Reasons;3644    uint32_t ReasonsMask = Event->memory_fault.fault_reason_mask;3645    if (ReasonsMask & HSA_AMD_MEMORY_FAULT_PAGE_NOT_PRESENT)3646      Reasons.emplace_back("Page not present or supervisor privilege");3647    if (ReasonsMask & HSA_AMD_MEMORY_FAULT_READ_ONLY)3648      Reasons.emplace_back("Write access to a read-only page");3649    if (ReasonsMask & HSA_AMD_MEMORY_FAULT_NX)3650      Reasons.emplace_back("Execute access to a page marked NX");3651    if (ReasonsMask & HSA_AMD_MEMORY_FAULT_HOST_ONLY)3652      Reasons.emplace_back("GPU attempted access to a host only page");3653    if (ReasonsMask & HSA_AMD_MEMORY_FAULT_DRAMECC)3654      Reasons.emplace_back("DRAM ECC failure");3655    if (ReasonsMask & HSA_AMD_MEMORY_FAULT_IMPRECISE)3656      Reasons.emplace_back("Can't determine the exact fault address");3657    if (ReasonsMask & HSA_AMD_MEMORY_FAULT_SRAMECC)3658      Reasons.emplace_back("SRAM ECC failure (ie registers, no fault address)");3659    if (ReasonsMask & HSA_AMD_MEMORY_FAULT_HANG)3660      Reasons.emplace_back("GPU reset following unspecified hang");3661 3662    // If we do not know the reason, say so, otherwise remove the trailing comma3663    // and space.3664    if (Reasons.empty())3665      Reasons.emplace_back("Unknown (" + std::to_string(ReasonsMask) + ")");3666 3667    uint32_t Node = -1;3668    hsa_agent_get_info(Event->memory_fault.agent, HSA_AGENT_INFO_NODE, &Node);3669 3670    AMDGPUPluginTy &Plugin = *reinterpret_cast<AMDGPUPluginTy *>(PluginPtr);3671    for (uint32_t I = 0, E = Plugin.getNumDevices();3672         Node != uint32_t(-1) && I < E; ++I) {3673      AMDGPUDeviceTy &AMDGPUDevice =3674          reinterpret_cast<AMDGPUDeviceTy &>(Plugin.getDevice(I));3675      auto KernelTraceInfoRecord =3676          AMDGPUDevice.KernelLaunchTraces.getExclusiveAccessor();3677 3678      uint32_t DeviceNode = -1;3679      if (auto Err =3680              AMDGPUDevice.getDeviceAttr(HSA_AGENT_INFO_NODE, DeviceNode)) {3681        consumeError(std::move(Err));3682        continue;3683      }3684      if (DeviceNode != Node)3685        continue;3686      void *DevicePtr = (void *)Event->memory_fault.virtual_address;3687      std::string S;3688      llvm::raw_string_ostream OS(S);3689      OS << llvm::format("memory access fault by GPU %" PRIu323690                         " (agent 0x%" PRIx643691                         ") at virtual address %p. Reasons: %s",3692                         Node, Event->memory_fault.agent.handle,3693                         (void *)Event->memory_fault.virtual_address,3694                         llvm::join(Reasons, ", ").c_str());3695      ErrorReporter::reportKernelTraces(AMDGPUDevice, *KernelTraceInfoRecord);3696      ErrorReporter::reportMemoryAccessError(AMDGPUDevice, DevicePtr, S,3697                                             /*Abort*/ true);3698    }3699 3700    // Abort the execution since we do not recover from this error.3701    FATAL_MESSAGE(1,3702                  "memory access fault by GPU %" PRIu32 " (agent 0x%" PRIx643703                  ") at virtual address %p. Reasons: %s",3704                  Node, Event->memory_fault.agent.handle,3705                  (void *)Event->memory_fault.virtual_address,3706                  llvm::join(Reasons, ", ").c_str());3707 3708    return HSA_STATUS_ERROR;3709  }3710 3711  /// Indicate whether the HSA runtime was correctly initialized. Even if there3712  /// is no available devices this boolean will be true. It indicates whether3713  /// we can safely call HSA functions (e.g., hsa_shut_down).3714  bool Initialized;3715 3716  /// Arrays of the available GPU and CPU agents. These arrays of handles should3717  /// not be here but in the AMDGPUDeviceTy structures directly. However, the3718  /// HSA standard does not provide API functions to retirve agents directly,3719  /// only iterating functions. We cache the agents here for convenience.3720  llvm::SmallVector<hsa_agent_t> KernelAgents;3721 3722  /// The device representing all HSA host agents.3723  AMDHostDeviceTy *HostDevice;3724};3725 3726Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice,3727                                 uint32_t NumThreads[3], uint32_t NumBlocks[3],3728                                 KernelArgsTy &KernelArgs,3729                                 KernelLaunchParamsTy LaunchParams,3730                                 AsyncInfoWrapperTy &AsyncInfoWrapper) const {3731  AMDGPUPluginTy &AMDGPUPlugin =3732      static_cast<AMDGPUPluginTy &>(GenericDevice.Plugin);3733  AMDHostDeviceTy &HostDevice = AMDGPUPlugin.getHostDevice();3734  AMDGPUMemoryManagerTy &ArgsMemoryManager = HostDevice.getArgsMemoryManager();3735 3736  void *AllArgs = nullptr;3737  if (auto Err = ArgsMemoryManager.allocate(ArgsSize, &AllArgs))3738    return Err;3739 3740  // Account for user requested dynamic shared memory.3741  uint32_t GroupSize = getGroupSize();3742  if (uint32_t MaxDynCGroupMem = std::max(3743          KernelArgs.DynCGroupMem, GenericDevice.getDynamicMemorySize())) {3744    GroupSize += MaxDynCGroupMem;3745  }3746 3747  uint64_t StackSize;3748  if (auto Err = GenericDevice.getDeviceStackSize(StackSize))3749    return Err;3750 3751  // Copy the explicit arguments.3752  // TODO: We should expose the args memory manager alloc to the common part as3753  // 	   alternative to copying them twice.3754  if (LaunchParams.Size)3755    std::memcpy(AllArgs, LaunchParams.Data, LaunchParams.Size);3756 3757  AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(GenericDevice);3758 3759  AMDGPUStreamTy *Stream = nullptr;3760  if (auto Err = AMDGPUDevice.getStream(AsyncInfoWrapper, Stream))3761    return Err;3762 3763  uint64_t ImplArgsOffset = utils::roundUp(3764      LaunchParams.Size, alignof(hsa_utils::AMDGPUImplicitArgsTy));3765  if (ArgsSize > ImplArgsOffset) {3766    hsa_utils::AMDGPUImplicitArgsTy *ImplArgs =3767        reinterpret_cast<hsa_utils::AMDGPUImplicitArgsTy *>(3768            utils::advancePtr(AllArgs, ImplArgsOffset));3769 3770    // Set the COV5+ implicit arguments to the appropriate values if present.3771    uint64_t ImplArgsSize = ArgsSize - ImplArgsOffset;3772    std::memset(ImplArgs, 0, ImplArgsSize);3773 3774    using ImplArgsTy = hsa_utils::AMDGPUImplicitArgsTy;3775    hsa_utils::initImplArg(ImplArgs, &ImplArgsTy::BlockCountX, ImplArgsSize,3776                           NumBlocks[0]);3777    hsa_utils::initImplArg(ImplArgs, &ImplArgsTy::BlockCountY, ImplArgsSize,3778                           NumBlocks[1]);3779    hsa_utils::initImplArg(ImplArgs, &ImplArgsTy::BlockCountZ, ImplArgsSize,3780                           NumBlocks[2]);3781 3782    hsa_utils::initImplArg(ImplArgs, &ImplArgsTy::GroupSizeX, ImplArgsSize,3783                           NumThreads[0]);3784    hsa_utils::initImplArg(ImplArgs, &ImplArgsTy::GroupSizeY, ImplArgsSize,3785                           NumThreads[1]);3786    hsa_utils::initImplArg(ImplArgs, &ImplArgsTy::GroupSizeZ, ImplArgsSize,3787                           NumThreads[2]);3788 3789    hsa_utils::initImplArg(ImplArgs, &ImplArgsTy::GridDims, ImplArgsSize,3790                           NumBlocks[2] * NumThreads[2] > 13791                               ? 33792                               : 1 + (NumBlocks[1] * NumThreads[1] != 1));3793 3794    hsa_utils::initImplArg(ImplArgs, &ImplArgsTy::DynamicLdsSize, ImplArgsSize,3795                           KernelArgs.DynCGroupMem);3796  }3797 3798  // Push the kernel launch into the stream.3799  return Stream->pushKernelLaunch(*this, AllArgs, NumThreads, NumBlocks,3800                                  GroupSize, StackSize, ArgsMemoryManager);3801}3802 3803Error AMDGPUKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice,3804                                             KernelArgsTy &KernelArgs,3805                                             uint32_t NumThreads[3],3806                                             uint32_t NumBlocks[3]) const {3807  // Only do all this when the output is requested3808  if (!(getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL))3809    return Plugin::success();3810 3811  // We don't have data to print additional info, but no hard error3812  if (!KernelInfo.has_value())3813    return Plugin::success();3814 3815  // General Info3816  auto *NumGroups = NumBlocks;3817  auto *ThreadsPerGroup = NumThreads;3818 3819  // Kernel Arguments Info3820  auto ArgNum = KernelArgs.NumArgs;3821  auto LoopTripCount = KernelArgs.Tripcount;3822 3823  // Details for AMDGPU kernels (read from image)3824  // https://www.llvm.org/docs/AMDGPUUsage.html#code-object-v4-metadata3825  auto GroupSegmentSize = (*KernelInfo).GroupSegmentList;3826  auto SGPRCount = (*KernelInfo).SGPRCount;3827  auto VGPRCount = (*KernelInfo).VGPRCount;3828  auto SGPRSpillCount = (*KernelInfo).SGPRSpillCount;3829  auto VGPRSpillCount = (*KernelInfo).VGPRSpillCount;3830  auto MaxFlatWorkgroupSize = (*KernelInfo).MaxFlatWorkgroupSize;3831 3832  // Prints additional launch info that contains the following.3833  // Num Args: The number of kernel arguments3834  // Teams x Thrds: The number of teams and the number of threads actually3835  // running.3836  // MaxFlatWorkgroupSize: Maximum flat work-group size supported by the3837  // kernel in work-items3838  // LDS Usage: Amount of bytes used in LDS storage3839  // S/VGPR Count: the number of S/V GPRs occupied by the kernel3840  // S/VGPR Spill Count: how many S/VGPRs are spilled by the kernel3841  // Tripcount: loop tripcount for the kernel3842  INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(),3843       "#Args: %d Teams x Thrds: %4ux%4u (MaxFlatWorkGroupSize: %u) LDS "3844       "Usage: %uB #SGPRs/VGPRs: %u/%u #SGPR/VGPR Spills: %u/%u Tripcount: "3845       "%lu\n",3846       ArgNum, NumGroups[0] * NumGroups[1] * NumGroups[2],3847       ThreadsPerGroup[0] * ThreadsPerGroup[1] * ThreadsPerGroup[2],3848       MaxFlatWorkgroupSize, GroupSegmentSize, SGPRCount, VGPRCount,3849       SGPRSpillCount, VGPRSpillCount, LoopTripCount);3850 3851  return Plugin::success();3852}3853 3854template <typename... ArgsTy>3855static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) {3856  hsa_status_t ResultCode = static_cast<hsa_status_t>(Code);3857  if (ResultCode == HSA_STATUS_SUCCESS || ResultCode == HSA_STATUS_INFO_BREAK)3858    return Plugin::success();3859 3860  const char *Desc = "unknown error";3861  hsa_status_t Ret = hsa_status_string(ResultCode, &Desc);3862  if (Ret != HSA_STATUS_SUCCESS)3863    REPORT("Unrecognized " GETNAME(TARGET_NAME) " error code %d\n", Code);3864 3865  // TODO: Add more entries to this switch3866  ErrorCode OffloadErrCode;3867  switch (ResultCode) {3868  case HSA_STATUS_ERROR_INVALID_SYMBOL_NAME:3869    OffloadErrCode = ErrorCode::NOT_FOUND;3870    break;3871  case HSA_STATUS_ERROR_INVALID_CODE_OBJECT:3872    OffloadErrCode = ErrorCode::INVALID_BINARY;3873    break;3874  default:3875    OffloadErrCode = ErrorCode::UNKNOWN;3876  }3877 3878  return Plugin::error(OffloadErrCode, ErrFmt, Args..., Desc);3879}3880 3881Expected<void *> AMDGPUMemoryManagerTy::allocate(size_t Size, void *HstPtr,3882                                                 TargetAllocTy Kind) {3883  // Allocate memory from the pool.3884  void *Ptr = nullptr;3885  if (auto Err = MemoryPool->allocate(Size, &Ptr))3886    return std::move(Err);3887 3888  assert(Ptr && "Invalid pointer");3889 3890  // Get a list of agents that can access this memory pool.3891  llvm::SmallVector<hsa_agent_t> Agents;3892  llvm::copy_if(3893      Plugin.getKernelAgents(), std::back_inserter(Agents),3894      [&](hsa_agent_t Agent) { return MemoryPool->canAccess(Agent); });3895 3896  // Allow all valid kernel agents to access the allocation.3897  if (auto Err = MemoryPool->enableAccess(Ptr, Size, Agents))3898    return std::move(Err);3899  return Ptr;3900}3901 3902Expected<void *> AMDGPUDeviceTy::allocate(size_t Size, void *,3903                                          TargetAllocTy Kind) {3904  if (Size == 0)3905    return nullptr;3906 3907  // Find the correct memory pool.3908  AMDGPUMemoryPoolTy *MemoryPool = nullptr;3909  switch (Kind) {3910  case TARGET_ALLOC_DEFAULT:3911  case TARGET_ALLOC_DEVICE:3912    MemoryPool = CoarseGrainedMemoryPools[0];3913    break;3914  case TARGET_ALLOC_HOST:3915    MemoryPool = &HostDevice.getFineGrainedMemoryPool();3916    break;3917  case TARGET_ALLOC_SHARED:3918    MemoryPool = &HostDevice.getFineGrainedMemoryPool();3919    break;3920  }3921 3922  if (!MemoryPool)3923    return Plugin::error(ErrorCode::UNSUPPORTED,3924                         "no memory pool for the specified allocation kind");3925 3926  // Allocate from the corresponding memory pool.3927  void *Alloc = nullptr;3928  if (auto Err = MemoryPool->allocate(Size, &Alloc))3929    return std::move(Err);3930 3931  if (Alloc) {3932    // Get a list of agents that can access this memory pool. Inherently3933    // necessary for host or shared allocations Also enabled for device memory3934    // to allow device to device memcpy3935    llvm::SmallVector<hsa_agent_t> Agents;3936    llvm::copy_if(static_cast<AMDGPUPluginTy &>(Plugin).getKernelAgents(),3937                  std::back_inserter(Agents), [&](hsa_agent_t Agent) {3938                    return MemoryPool->canAccess(Agent);3939                  });3940 3941    // Enable all valid kernel agents to access the buffer.3942    if (auto Err = MemoryPool->enableAccess(Alloc, Size, Agents))3943      return std::move(Err);3944  }3945 3946  return Alloc;3947}3948 3949void AMDGPUQueueTy::callbackError(hsa_status_t Status, hsa_queue_t *Source,3950                                  void *Data) {3951  auto &AMDGPUDevice = *reinterpret_cast<AMDGPUDeviceTy *>(Data);3952 3953  if (Status == HSA_STATUS_ERROR_EXCEPTION) {3954    auto KernelTraceInfoRecord =3955        AMDGPUDevice.KernelLaunchTraces.getExclusiveAccessor();3956    std::function<bool(__tgt_async_info &)> AsyncInfoWrapperMatcher =3957        [=](__tgt_async_info &AsyncInfo) {3958          auto *Stream = reinterpret_cast<AMDGPUStreamTy *>(AsyncInfo.Queue);3959          if (!Stream || !Stream->getQueue())3960            return false;3961          return Stream->getQueue()->Queue == Source;3962        };3963    ErrorReporter::reportTrapInKernel(AMDGPUDevice, *KernelTraceInfoRecord,3964                                      AsyncInfoWrapperMatcher);3965  }3966 3967  auto Err = Plugin::check(Status, "received error in queue %p: %s", Source);3968  FATAL_MESSAGE(1, "%s", toString(std::move(Err)).data());3969}3970 3971} // namespace plugin3972} // namespace target3973} // namespace omp3974} // namespace llvm3975 3976extern "C" {3977llvm::omp::target::plugin::GenericPluginTy *createPlugin_amdgpu() {3978  return new llvm::omp::target::plugin::AMDGPUPluginTy();3979}3980}3981