//===- PluginInterface.cpp - Target independent plugin device interface ---===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
//===----------------------------------------------------------------------===//

#include "PluginInterface.h"

#include "Shared/APITypes.h"
#include "Shared/Debug.h"
#include "Shared/Environment.h"

#include "ErrorReporting.h"
#include "GlobalHandler.h"
#include "JIT.h"
#include "Shared/Utils.h"
#include "Utils/ELF.h"
#include "omptarget.h"

#ifdef OMPT_SUPPORT
#include "OpenMP/OMPT/Callback.h"
#include "omp-tools.h"
#endif

#include "llvm/Bitcode/BitcodeReader.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
#include "llvm/Support/Error.h"
#include "llvm/Support/MathExtras.h"
#include "llvm/Support/MemoryBuffer.h"
#include "llvm/Support/Signals.h"
#include "llvm/Support/raw_ostream.h"

#include <cstdint>
#include <limits>

using namespace llvm;
using namespace omp;
using namespace target;
using namespace plugin;
using namespace error;
using namespace llvm::offload::debug;

AsyncInfoWrapperTy::AsyncInfoWrapperTy(GenericDeviceTy &Device,
                                       __tgt_async_info *AsyncInfoPtr)
    : Device(Device),
      AsyncInfoPtr(AsyncInfoPtr ? AsyncInfoPtr : &LocalAsyncInfo) {}

Error AsyncInfoWrapperTy::synchronize() {
  assert(AsyncInfoPtr && "AsyncInfoWrapperTy already finalized");

  // Synchronize with the async info's operations without releasing the queue.
  return Device.synchronize(AsyncInfoPtr, /*ReleaseQueue=*/false);
}

void AsyncInfoWrapperTy::finalize(Error &Err) {
  assert(AsyncInfoPtr && "AsyncInfoWrapperTy already finalized");

  // If we used a local async info object we want synchronous behavior. In that
  // case, and assuming the current status code is correct, we will synchronize
  // explicitly when the object is deleted. Update the error with the result of
  // the synchronize operation.
  if (AsyncInfoPtr == &LocalAsyncInfo && LocalAsyncInfo.Queue && !Err)
    Err = Device.synchronize(&LocalAsyncInfo);

  // Invalidate the wrapper object.
  AsyncInfoPtr = nullptr;
}

Error GenericKernelTy::init(GenericDeviceTy &GenericDevice,
                            DeviceImageTy &Image) {

  ImagePtr = &Image;

  // Retrieve kernel environment object for the kernel.
  std::string EnvironmentName = std::string(Name) + "_kernel_environment";
  GenericGlobalHandlerTy &GHandler = GenericDevice.Plugin.getGlobalHandler();
  if (GHandler.isSymbolInImage(GenericDevice, Image, EnvironmentName)) {
    GlobalTy KernelEnv(EnvironmentName, sizeof(KernelEnvironment),
                       &KernelEnvironment);
    if (auto Err =
            GHandler.readGlobalFromImage(GenericDevice, *ImagePtr, KernelEnv))
      return Err;
  } else {
    KernelEnvironment = KernelEnvironmentTy{};
    ODBG(OLDT_Kernel) << "Failed to read kernel environment for '" << getName()
                      << "' Using default Bare (0) execution mode";
  }

  // Max = Config.Max > 0 ? min(Config.Max, Device.Max) : Device.Max;
  MaxNumThreads = KernelEnvironment.Configuration.MaxThreads > 0
                      ? std::min(KernelEnvironment.Configuration.MaxThreads,
                                 int32_t(GenericDevice.getThreadLimit()))
                      : GenericDevice.getThreadLimit();

  // Pref = Config.Pref > 0 ? max(Config.Pref, Device.Pref) : Device.Pref;
  PreferredNumThreads =
      KernelEnvironment.Configuration.MinThreads > 0
          ? std::max(KernelEnvironment.Configuration.MinThreads,
                     int32_t(GenericDevice.getDefaultNumThreads()))
          : GenericDevice.getDefaultNumThreads();

  return initImpl(GenericDevice, Image);
}

Expected<KernelLaunchEnvironmentTy *>
GenericKernelTy::getKernelLaunchEnvironment(
    GenericDeviceTy &GenericDevice, const KernelArgsTy &KernelArgs,
    const DynBlockMemConfTy &DynBlockMemConf,
    AsyncInfoWrapperTy &AsyncInfoWrapper) const {
  // Ctor/Dtor have no arguments, replaying uses the original kernel launch
  // environment. Older versions of the compiler do not generate a kernel
  // launch environment.
  if ((GenericDevice.getRecordReplay() &&
       GenericDevice.getRecordReplay()->isReplaying()) ||
      KernelArgs.Version < OMP_KERNEL_ARG_MIN_VERSION_WITH_DYN_PTR)
    return nullptr;

  if ((!KernelEnvironment.Configuration.ReductionDataSize ||
       !KernelEnvironment.Configuration.ReductionBufferLength) &&
      KernelArgs.DynCGroupMem == 0)
    return reinterpret_cast<KernelLaunchEnvironmentTy *>(~0);

  auto AllocOrErr = GenericDevice.dataAlloc(sizeof(KernelLaunchEnvironmentTy),
                                            /*HostPtr=*/nullptr,
                                            TargetAllocTy::TARGET_ALLOC_DEVICE);
  if (!AllocOrErr)
    return AllocOrErr.takeError();

  // Remember to free the memory later.
  AsyncInfoWrapper.freeAllocationAfterSynchronization(*AllocOrErr);

  /// Use the KLE in the __tgt_async_info to ensure a stable address for the
  /// async data transfer.
  auto &LocalKLE = (*AsyncInfoWrapper).KernelLaunchEnvironment;
  LocalKLE = KernelLaunchEnvironment;

  LocalKLE.DynCGroupMemSize = DynBlockMemConf.Size;
  LocalKLE.DynCGroupMemFbPtr = DynBlockMemConf.FallbackPtr;
  LocalKLE.DynCGroupMemFb = DynBlockMemConf.Fallback;
  LocalKLE.ReductionBuffer = nullptr;

  if (KernelEnvironment.Configuration.ReductionDataSize &&
      KernelEnvironment.Configuration.ReductionBufferLength) {
    auto AllocOrErr = GenericDevice.dataAlloc(
        KernelEnvironment.Configuration.ReductionDataSize *
            KernelEnvironment.Configuration.ReductionBufferLength,
        /*HostPtr=*/nullptr, TargetAllocTy::TARGET_ALLOC_DEVICE);
    if (!AllocOrErr)
      return AllocOrErr.takeError();
    LocalKLE.ReductionBuffer = *AllocOrErr;
    // Remember to free the memory later.
    AsyncInfoWrapper.freeAllocationAfterSynchronization(*AllocOrErr);
  }

  INFO(OMP_INFOTYPE_DATA_TRANSFER, GenericDevice.getDeviceId(),
       "Copying data from host to device, HstPtr=" DPxMOD ", TgtPtr=" DPxMOD
       ", Size=%" PRId64 ", Name=KernelLaunchEnv\n",
       DPxPTR(&LocalKLE), DPxPTR(*AllocOrErr),
       sizeof(KernelLaunchEnvironmentTy));

  auto Err = GenericDevice.dataSubmit(*AllocOrErr, &LocalKLE,
                                      sizeof(KernelLaunchEnvironmentTy),
                                      AsyncInfoWrapper);
  if (Err)
    return Err;
  return static_cast<KernelLaunchEnvironmentTy *>(*AllocOrErr);
}

Error GenericKernelTy::printLaunchInfo(GenericDeviceTy &GenericDevice,
                                       KernelArgsTy &KernelArgs,
                                       uint32_t NumThreads[3],
                                       uint32_t NumBlocks[3]) const {
  INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(),
       "Launching kernel %s with [%u,%u,%u] blocks and [%u,%u,%u] threads in "
       "%s mode\n",
       getName(), NumBlocks[0], NumBlocks[1], NumBlocks[2], NumThreads[0],
       NumThreads[1], NumThreads[2], getExecutionModeName());
  return printLaunchInfoDetails(GenericDevice, KernelArgs, NumThreads,
                                NumBlocks);
}

Error GenericKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice,
                                              KernelArgsTy &KernelArgs,
                                              uint32_t NumThreads[3],
                                              uint32_t NumBlocks[3]) const {
  return Plugin::success();
}

Expected<DynBlockMemConfTy>
GenericKernelTy::prepareBlockMemory(GenericDeviceTy &GenericDevice,
                                    KernelArgsTy &KernelArgs,
                                    uint32_t NumBlocks) const {
  uint32_t MaxBlockMemSize = GenericDevice.getMaxBlockSharedMemSize();
  uint32_t DynBlockMemSize = KernelArgs.DynCGroupMem;
  uint32_t TotalBlockMemSize = StaticBlockMemSize + DynBlockMemSize;
  uint32_t DynNativeBlockMemSize = DynBlockMemSize;
  void *DynFallbackPtr = nullptr;

  // No enough block memory to cover the static one. Cannot run the kernel.
  if (StaticBlockMemSize > MaxBlockMemSize)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "Static block memory size exceeds maximum");
  // No enough block memory to cover dynamic one, and the fallback is aborting.
  if (static_cast<DynCGroupMemFallbackType>(
          KernelArgs.Flags.DynCGroupMemFallback) ==
          DynCGroupMemFallbackType::Abort &&
      TotalBlockMemSize > MaxBlockMemSize)
    return Plugin::error(
        ErrorCode::INVALID_ARGUMENT,
        "Requested block memory size (static + dynamic) exceeds maximum");

  DynCGroupMemFallbackType DynFallback = DynCGroupMemFallbackType::None;
  if (DynBlockMemSize && TotalBlockMemSize > MaxBlockMemSize) {
    // Launch without native dynamic block memory.
    DynNativeBlockMemSize = 0;
    DynFallback = static_cast<DynCGroupMemFallbackType>(
        KernelArgs.Flags.DynCGroupMemFallback);
    if (DynFallback != DynCGroupMemFallbackType::DefaultMem) {
      // Do not provide any memory as fallback.
      DynBlockMemSize = 0;
    } else {
      // Get global memory as fallback.
      auto AllocOrErr = GenericDevice.dataAlloc(
          NumBlocks * DynBlockMemSize,
          /*HostPtr=*/nullptr, TargetAllocTy::TARGET_ALLOC_DEVICE);
      if (!AllocOrErr)
        return AllocOrErr.takeError();
      DynFallbackPtr = *AllocOrErr;
    }
  }
  return DynBlockMemConfTy{DynBlockMemSize, DynNativeBlockMemSize, DynFallback,
                           DynFallbackPtr};
}

Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
                              ptrdiff_t *ArgOffsets, KernelArgsTy &KernelArgs,
                              KernelExtraArgsTy *KernelExtraArgs,
                              AsyncInfoWrapperTy &AsyncInfoWrapper) const {
  llvm::SmallVector<void *, 16> Args;
  llvm::SmallVector<void *, 16> Ptrs;

  uint32_t EffectiveNumThreads[3] = {KernelArgs.UserThreadLimit[0],
                                     KernelArgs.UserThreadLimit[1],
                                     KernelArgs.UserThreadLimit[2]};
  uint32_t EffectiveNumBlocks[3] = {KernelArgs.UserNumBlocks[0],
                                    KernelArgs.UserNumBlocks[1],
                                    KernelArgs.UserNumBlocks[2]};
  if (!isBareMode()) {
    assert(
        EffectiveNumThreads[1] == 1 && EffectiveNumThreads[2] == 1 &&
        EffectiveNumBlocks[1] == 1 && EffectiveNumBlocks[2] == 1 &&
        "Non-bare mode should only use the first thread and block dimensions");
    EffectiveNumThreads[0] =
        getEffectiveNumThreads(GenericDevice, EffectiveNumThreads);
    EffectiveNumBlocks[0] = getEffectiveNumBlocks(
        GenericDevice, EffectiveNumBlocks, KernelArgs.Tripcount,
        EffectiveNumThreads[0], KernelArgs.UserThreadLimit[0] > 0);
  }

  auto DynBlockMemConfOrErr = prepareBlockMemory(
      GenericDevice, KernelArgs,
      EffectiveNumBlocks[0] * EffectiveNumBlocks[1] * EffectiveNumBlocks[2]);
  if (!DynBlockMemConfOrErr)
    return DynBlockMemConfOrErr.takeError();

  DynBlockMemConfTy &DynBlockMemConf = *DynBlockMemConfOrErr;
  if (DynBlockMemConf.FallbackPtr)
    AsyncInfoWrapper.freeAllocationAfterSynchronization(
        DynBlockMemConf.FallbackPtr);

  auto KernelLaunchEnvOrErr = getKernelLaunchEnvironment(
      GenericDevice, KernelArgs, DynBlockMemConf, AsyncInfoWrapper);
  if (!KernelLaunchEnvOrErr)
    return KernelLaunchEnvOrErr.takeError();

  KernelLaunchParamsTy LaunchParams;

  // Kernel languages don't use indirection.
  if (KernelArgs.Flags.IsCUDA) {
    LaunchParams =
        *reinterpret_cast<KernelLaunchParamsTy *>(KernelArgs.ArgPtrs);
  } else {
    LaunchParams =
        prepareArgs(GenericDevice, ArgPtrs, ArgOffsets, KernelArgs.NumArgs,
                    Args, Ptrs, *KernelLaunchEnvOrErr, KernelArgs.Version);
  }

  if (auto Err = printLaunchInfo(GenericDevice, KernelArgs, EffectiveNumThreads,
                                 EffectiveNumBlocks))
    return Err;

  RecordReplayTy::HandleTy RRHandle;
  RecordReplayTy *RecordReplay = GenericDevice.getRecordReplay();
  if (RecordReplay) {
    // Record replay requires synchronization of any previous operation.
    if (auto Err = AsyncInfoWrapper.synchronize())
      return Err;

    // Record the kernel prologue data before kernel launch.
    auto RRHandleOrErr = RecordReplay->recordPrologue(
        *this, KernelArgs, KernelExtraArgs, LaunchParams, EffectiveNumBlocks,
        EffectiveNumThreads, DynBlockMemConf.NativeSize);
    if (!RRHandleOrErr)
      return RRHandleOrErr.takeError();
    RRHandle = *RRHandleOrErr;
  }

  if (auto Err = launchImpl(GenericDevice, EffectiveNumThreads,
                            EffectiveNumBlocks, DynBlockMemConf.NativeSize,
                            KernelArgs, LaunchParams, AsyncInfoWrapper))
    return Err;

  if (RecordReplay) {
    // Record replay requires synchronization.
    if (auto Err = AsyncInfoWrapper.synchronize())
      return Err;

    // Record the epilogue data after kernel synchronization.
    return RecordReplay->recordEpilogue(*this, RRHandle);
  }
  return Plugin::success();
}

KernelLaunchParamsTy
GenericKernelTy::prepareArgs(GenericDeviceTy &GenericDevice, void **ArgPtrs,
                             ptrdiff_t *ArgOffsets, uint32_t &NumArgs,
                             llvm::SmallVectorImpl<void *> &Args,
                             llvm::SmallVectorImpl<void *> &Ptrs,
                             KernelLaunchEnvironmentTy *KernelLaunchEnvironment,
                             uint32_t Version) const {
  if (NumArgs == 0)
    return KernelLaunchParamsTy{};

  // The argument arrays already include the dyn_ptr slot at the end (appended
  // by the host for version >= 4, or by upgradeKernelArgs for version 3).
  Args.resize(NumArgs);
  Ptrs.resize(NumArgs);

  for (uint32_t I = 0; I < NumArgs; ++I)
    Args[I] = reinterpret_cast<void *>(reinterpret_cast<intptr_t>(ArgPtrs[I]) +
                                       ArgOffsets[I]);

  // Optionally assign the KernelLaunchEnvironment to the last slot (dyn_ptr).
  if (KernelLaunchEnvironment)
    Args[NumArgs - 1] = KernelLaunchEnvironment;

  // Version 3 device kernels have dyn_ptr baked in at position 0. Rotate the
  // last element to the front to match the device ABI.
  if (Version <= OMP_KERNEL_ARG_MIN_VERSION_WITH_DYN_PTR &&
      KernelLaunchEnvironment)
    std::rotate(Args.begin(), Args.end() - 1, Args.end());

  for (uint32_t I = 0; I < NumArgs; ++I)
    Ptrs[I] = &Args[I];

  return KernelLaunchParamsTy{sizeof(void *) * NumArgs, &Args[0], &Ptrs[0]};
}

uint32_t
GenericKernelTy::getEffectiveNumThreads(GenericDeviceTy &GenericDevice,
                                        uint32_t UserThreadLimit[3]) const {
  assert(!isBareMode() && "bare kernel should not call this function");

  assert(UserThreadLimit[1] == 1 && UserThreadLimit[2] == 1 &&
         "Multi dimensional launch not supported yet.");

  if (UserThreadLimit[0] > 0 && isGenericMode())
    UserThreadLimit[0] += GenericDevice.getWarpSize();

  return std::min(MaxNumThreads, (UserThreadLimit[0] > 0)
                                     ? UserThreadLimit[0]
                                     : PreferredNumThreads);
}

uint32_t GenericKernelTy::getEffectiveNumBlocks(
    GenericDeviceTy &GenericDevice, uint32_t UserNumBlocks[3],
    uint64_t LoopTripCount, uint32_t &EffectiveNumThreads,
    bool IsNumThreadsFromUser) const {
  assert(!isBareMode() && "bare kernel should not call this function");

  assert(UserNumBlocks[1] == 1 && UserNumBlocks[2] == 1 &&
         "Multi dimensional launch not supported yet.");

  if (UserNumBlocks[0] > 0) {
    // TODO: We need to honor any value and consequently allow more than the
    // block limit. For this we might need to start multiple kernels or let the
    // blocks start again until the requested number has been started.
    return std::min(UserNumBlocks[0], GenericDevice.getBlockLimit());
  }

  // Return the number of blocks required to cover the loop iterations.
  if (isNoLoopMode())
    return LoopTripCount > 0 ? (((LoopTripCount - 1) / EffectiveNumThreads) + 1)
                             : 1;

  uint64_t DefaultNumBlocks = GenericDevice.getDefaultNumBlocks();
  uint64_t TripCountNumBlocks = std::numeric_limits<uint64_t>::max();
  if (LoopTripCount > 0) {
    if (isSPMDMode()) {
      // We have a combined construct, i.e. `target teams distribute
      // parallel for [simd]`. We launch so many blocks so that each thread
      // will execute one iteration of the loop; rounded up to the nearest
      // integer. However, if that results in too few blocks, we artificially
      // reduce the thread count per block to increase the outer parallelism.
      auto MinThreads = GenericDevice.getMinThreadsForLowTripCountLoop();
      MinThreads = std::min(MinThreads, EffectiveNumThreads);

      // Honor the thread_limit clause; only lower the number of threads.
      [[maybe_unused]] auto OldNumThreads = EffectiveNumThreads;
      if (LoopTripCount >= DefaultNumBlocks * EffectiveNumThreads ||
          IsNumThreadsFromUser) {
        // Enough parallelism for blocks and threads.
        TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
        assert(IsNumThreadsFromUser ||
               TripCountNumBlocks >= DefaultNumBlocks &&
                   "Expected sufficient outer parallelism.");
      } else if (LoopTripCount >= DefaultNumBlocks * MinThreads) {
        // Enough parallelism for blocks, limit threads.

        // This case is hard; for now, we force "full warps":
        // First, compute a thread count assuming DefaultNumBlocks.
        auto NumThreadsDefaultBlocks =
            (LoopTripCount + DefaultNumBlocks - 1) / DefaultNumBlocks;
        // Now get a power of two that is larger or equal.
        auto NumThreadsDefaultBlocksP2 =
            llvm::PowerOf2Ceil(NumThreadsDefaultBlocks);
        // Do not increase a thread limit given be the user.
        EffectiveNumThreads =
            std::min(EffectiveNumThreads, uint32_t(NumThreadsDefaultBlocksP2));
        assert(EffectiveNumThreads >= MinThreads &&
               "Expected sufficient inner parallelism.");
        TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
      } else {
        // Not enough parallelism for blocks and threads, limit both.
        EffectiveNumThreads = std::min(EffectiveNumThreads, MinThreads);
        TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
      }

      assert(EffectiveNumThreads * TripCountNumBlocks >= LoopTripCount &&
             "Expected sufficient parallelism");
      assert(OldNumThreads >= EffectiveNumThreads &&
             "Number of threads cannot be increased!");
    } else {
      assert((isGenericMode() || isGenericSPMDMode()) &&
             "Unexpected execution mode!");
      // If we reach this point, then we have a non-combined construct, i.e.
      // `teams distribute` with a nested `parallel for` and each block is
      // assigned one iteration of the `distribute` loop. E.g.:
      //
      // #pragma omp target teams distribute
      // for(...loop_tripcount...) {
      //   #pragma omp parallel for
      //   for(...) {}
      // }
      //
      // Threads within a block will execute the iterations of the `parallel`
      // loop.
      TripCountNumBlocks = LoopTripCount;
    }
  }

  uint32_t PreferredNumBlocks = TripCountNumBlocks;
  // If the loops are long running we rather reuse blocks than spawn too many.
  if (GenericDevice.getReuseBlocksForHighTripCount())
    PreferredNumBlocks = std::min(TripCountNumBlocks, DefaultNumBlocks);
  return std::min(PreferredNumBlocks, GenericDevice.getBlockLimit());
}

GenericDeviceTy::GenericDeviceTy(GenericPluginTy &Plugin, int32_t DeviceId,
                                 int32_t NumDevices,
                                 const llvm::omp::GV &OMPGridValues)
    : Plugin(Plugin), MemoryManager(nullptr), OMP_TeamLimit("OMP_TEAM_LIMIT"),
      OMP_NumTeams("OMP_NUM_TEAMS"),
      OMP_TeamsThreadLimit("OMP_TEAMS_THREAD_LIMIT"),
      OMPX_DebugKind("LIBOMPTARGET_DEVICE_RTL_DEBUG"),
      // Do not initialize the following two envars since they depend on the
      // device initialization. These cannot be consulted until the device is
      // initialized correctly. We initialize them in GenericDeviceTy::init().
      OMPX_TargetStackSize(), OMPX_TargetHeapSize(),
      // By default, the initial number of streams and events is 1.
      OMPX_InitialNumStreams("LIBOMPTARGET_NUM_INITIAL_STREAMS", 1),
      OMPX_InitialNumEvents("LIBOMPTARGET_NUM_INITIAL_EVENTS", 1),
      DeviceId(DeviceId), GridValues(OMPGridValues),
      PeerAccesses(NumDevices, PeerAccessState::PENDING), PeerAccessesLock(),
      PinnedAllocs(*this), RPCServer(nullptr) {
  // Conservative fall-back to the plugin's device uid for the case that no real
  // vendor (u)uid will become available later.
  setDeviceUidFromVendorUid(std::to_string(static_cast<uint64_t>(DeviceId)));

#ifdef OMPT_SUPPORT
  OmptInitialized.store(false);
  // Bind the callbacks to this device's member functions
#define bindOmptCallback(Name, Type, Code)                                     \
  if (ompt::Initialized && ompt::lookupCallbackByCode) {                       \
    ompt::lookupCallbackByCode((ompt_callbacks_t)(Code),                       \
                               ((ompt_callback_t *)&(Name##_fn)));             \
    ODBG(OLDT_Tool) << "OMPT: class bound " << #Name << "="                    \
                    << ((void *)(uint64_t)Name##_fn);                          \
  }

  FOREACH_OMPT_DEVICE_EVENT(bindOmptCallback);
#undef bindOmptCallback

#endif

  // Envar that indicates whether mapped host buffers should be locked
  // automatically. The possible values are boolean (on/off) and a special:
  //   off:       Mapped host buffers are not locked.
  //   on:        Mapped host buffers are locked in a best-effort approach.
  //              Failure to lock the buffers are silent.
  //   mandatory: Mapped host buffers are always locked and failures to lock
  //              a buffer results in a fatal error.
  StringEnvar OMPX_LockMappedBuffers("LIBOMPTARGET_LOCK_MAPPED_HOST_BUFFERS",
                                     "off");

  bool Enabled;
  if (StringParser::parse(OMPX_LockMappedBuffers.get().data(), Enabled)) {
    // Parsed as a boolean value. Enable the feature if necessary.
    LockMappedBuffers = Enabled;
    IgnoreLockMappedFailures = true;
  } else if (OMPX_LockMappedBuffers.get() == "mandatory") {
    // Enable the feature and failures are fatal.
    LockMappedBuffers = true;
    IgnoreLockMappedFailures = false;
  } else {
    // Disable by default.
    ODBG(OLDT_Alloc) << "Invalid value LIBOMPTARGET_LOCK_MAPPED_HOST_BUFFERS="
                     << OMPX_LockMappedBuffers.get();
    LockMappedBuffers = false;
  }
}

Error GenericDeviceTy::init(GenericPluginTy &Plugin) {
  if (auto Err = initImpl(Plugin))
    return Err;

#ifdef OMPT_SUPPORT
  if (ompt::Initialized) {
    bool ExpectedStatus = false;
    if (OmptInitialized.compare_exchange_strong(ExpectedStatus, true))
      performOmptCallback(device_initialize, Plugin.getUserId(DeviceId),
                          /*type=*/getComputeUnitKind().c_str(),
                          /*device=*/reinterpret_cast<ompt_device_t *>(this),
                          /*lookup=*/ompt::lookupCallbackByName,
                          /*documentation=*/nullptr);
  }
#endif

  // Read and reinitialize the envars that depend on the device initialization.
  // Notice these two envars may change the stack size and heap size of the
  // device, so they need the device properly initialized.
  auto StackSizeEnvarOrErr = UInt64Envar::create(
      "LIBOMPTARGET_STACK_SIZE",
      [this](uint64_t &V) -> Error { return getDeviceStackSize(V); },
      [this](uint64_t V) -> Error { return setDeviceStackSize(V); });
  if (!StackSizeEnvarOrErr)
    return StackSizeEnvarOrErr.takeError();
  OMPX_TargetStackSize = std::move(*StackSizeEnvarOrErr);

  if (hasDeviceHeapSize()) {
    auto HeapSizeEnvarOrErr = UInt64Envar::create(
        "LIBOMPTARGET_HEAP_SIZE",
        [this](uint64_t &V) -> Error { return getDeviceHeapSize(V); },
        [this](uint64_t V) -> Error { return setDeviceHeapSize(V); });
    if (!HeapSizeEnvarOrErr)
      return HeapSizeEnvarOrErr.takeError();
    OMPX_TargetHeapSize = std::move(*HeapSizeEnvarOrErr);
  }

  // Update the maximum number of teams and threads after the device
  // initialization sets the corresponding hardware limit.
  if (OMP_NumTeams > 0)
    GridValues.GV_Max_Teams =
        std::min(GridValues.GV_Max_Teams, uint32_t(OMP_NumTeams));

  if (OMP_TeamsThreadLimit > 0)
    GridValues.GV_Max_WG_Size =
        std::min(GridValues.GV_Max_WG_Size, uint32_t(OMP_TeamsThreadLimit));

  // Enable the memory manager if required.
  auto [ThresholdMM, EnableMM] = MemoryManagerTy::getSizeThresholdFromEnv();
  if (EnableMM) {
    if (ThresholdMM == 0)
      ThresholdMM = getMemoryManagerSizeThreshold();
    MemoryManager = new MemoryManagerTy(*this, ThresholdMM);
  }

  return Plugin::success();
}

Error GenericDeviceTy::unloadBinary(DeviceImageTy *Image) {
  if (auto Err = callGlobalDestructors(Plugin, *Image))
    return Err;

  GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler();
  auto ProfOrErr = Handler.readProfilingGlobals(*this, *Image);
  if (!ProfOrErr)
    return ProfOrErr.takeError();

  if (!ProfOrErr->empty()) {
    // Dump out profdata
    if ((OMPX_DebugKind.get() & uint32_t(DeviceDebugKind::PGODump)) ==
        uint32_t(DeviceDebugKind::PGODump))
      ProfOrErr->dump();

    // Write data to profiling file
    if (auto Err = ProfOrErr->write())
      return Err;
  }

  return unloadBinaryImpl(Image);
}

Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
  for (auto &I : LoadedImages)
    if (auto Err = unloadBinary(I))
      return Err;
  LoadedImages.clear();

  // Delete the memory manager before deinitializing the device. Otherwise,
  // we may delete device allocations after the device is deinitialized.
  if (MemoryManager)
    delete MemoryManager;
  MemoryManager = nullptr;

  if (RecordReplay) {
    if (auto Err = RecordReplay->deinit())
      return Err;
    delete RecordReplay;
    RecordReplay = nullptr;
  }

  if (RPCServer)
    if (auto Err = RPCServer->deinitDevice(*this))
      return Err;

#ifdef OMPT_SUPPORT
  if (ompt::Initialized) {
    bool ExpectedStatus = true;
    if (OmptInitialized.compare_exchange_strong(ExpectedStatus, false))
      performOmptCallback(device_finalize, Plugin.getUserId(DeviceId));
  }
#endif

  return deinitImpl();
}
Expected<DeviceImageTy *> GenericDeviceTy::loadBinary(GenericPluginTy &Plugin,
                                                      StringRef InputTgtImage) {
  ODBG(OLDT_Init) << "Load data from image "
                  << static_cast<const void *>(InputTgtImage.bytes_begin());

  std::unique_ptr<MemoryBuffer> Buffer;
  if (identify_magic(InputTgtImage) == file_magic::bitcode) {
    auto CompiledImageOrErr = Plugin.getJIT().process(InputTgtImage, *this);
    if (!CompiledImageOrErr) {
      return Plugin::error(ErrorCode::COMPILE_FAILURE,
                           CompiledImageOrErr.takeError(),
                           "failure to jit IR image");
    }
    Buffer = std::move(*CompiledImageOrErr);
  } else {
    Buffer = MemoryBuffer::getMemBufferCopy(InputTgtImage);
  }

  // Load the binary and allocate the image object. Use the next available id
  // for the image id, which is the number of previously loaded images.
  auto ImageOrErr = loadBinaryImpl(std::move(Buffer), LoadedImages.size());
  if (!ImageOrErr)
    return ImageOrErr.takeError();
  DeviceImageTy *Image = *ImageOrErr;

  // Add the image to list.
  LoadedImages.push_back(Image);

  if (auto Err = setupRPCServer(Plugin, *Image))
    return std::move(Err);

#ifdef OMPT_SUPPORT
  if (ompt::Initialized) {
    size_t Bytes = InputTgtImage.size();
    performOmptCallback(
        device_load, Plugin.getUserId(DeviceId),
        /*FileName=*/nullptr, /*FileOffset=*/0, /*VmaInFile=*/nullptr,
        /*ImgSize=*/Bytes,
        /*HostAddr=*/const_cast<unsigned char *>(InputTgtImage.bytes_begin()),
        /*DeviceAddr=*/nullptr, /* FIXME: ModuleId */ 0);
  }
#endif

  // Call any global constructors present on the device.
  if (auto Err = callGlobalConstructors(Plugin, *Image))
    return std::move(Err);

  // Return the pointer to the table of entries.
  return Image;
}

Error GenericDeviceTy::setupRPCServer(GenericPluginTy &Plugin,
                                      DeviceImageTy &Image) {
  // The plugin either does not need an RPC server or it is unavailable.
  if (!shouldSetupRPCServer())
    return Plugin::success();

  // Check if this device needs to run an RPC server.
  RPCServerTy &Server = Plugin.getRPCServer();
  auto UsingOrErr =
      Server.isDeviceUsingRPC(*this, Plugin.getGlobalHandler(), Image);
  if (!UsingOrErr)
    return UsingOrErr.takeError();

  if (!UsingOrErr.get())
    return Plugin::success();

  if (auto Err = Server.initDevice(*this, Plugin.getGlobalHandler(), Image))
    return Err;

  if (auto Err = Server.startThread())
    return Err;

  RPCServer = &Server;
  ODBG(OLDT_Init) << "Running an RPC server on device " << getDeviceId();
  return Plugin::success();
}

Error PinnedAllocationMapTy::insertEntry(void *HstPtr, void *DevAccessiblePtr,
                                         size_t Size, bool ExternallyLocked) {
  // Insert the new entry into the map.
  auto Res = Allocs.insert({HstPtr, DevAccessiblePtr, Size, ExternallyLocked});
  if (!Res.second)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "cannot insert locked buffer entry");

  // Check whether the next entry overlaps with the inserted entry.
  auto It = std::next(Res.first);
  if (It == Allocs.end())
    return Plugin::success();

  const EntryTy *NextEntry = &(*It);
  if (intersects(NextEntry->HstPtr, NextEntry->Size, HstPtr, Size))
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "partial overlapping not allowed in locked buffers");

  return Plugin::success();
}

Error PinnedAllocationMapTy::eraseEntry(const EntryTy &Entry) {
  // Erase the existing entry. Notice this requires an additional map lookup,
  // but this should not be a performance issue. Using iterators would make
  // the code more difficult to read.
  size_t Erased = Allocs.erase({Entry.HstPtr});
  if (!Erased)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "cannot erase locked buffer entry");
  return Plugin::success();
}

Error PinnedAllocationMapTy::registerEntryUse(const EntryTy &Entry,
                                              void *HstPtr, size_t Size) {
  if (!contains(Entry.HstPtr, Entry.Size, HstPtr, Size))
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "partial overlapping not allowed in locked buffers");

  ++Entry.References;
  return Plugin::success();
}

Expected<bool> PinnedAllocationMapTy::unregisterEntryUse(const EntryTy &Entry) {
  if (Entry.References == 0)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "invalid number of references");

  // Return whether this was the last user.
  return (--Entry.References == 0);
}

Error PinnedAllocationMapTy::registerHostBuffer(void *HstPtr,
                                                void *DevAccessiblePtr,
                                                size_t Size) {
  assert(HstPtr && "Invalid pointer");
  assert(DevAccessiblePtr && "Invalid pointer");
  assert(Size && "Invalid size");

  std::lock_guard<std::shared_mutex> Lock(Mutex);

  // No pinned allocation should intersect.
  const EntryTy *Entry = findIntersecting(HstPtr);
  if (Entry)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "cannot insert entry due to an existing one");

  // Now insert the new entry.
  return insertEntry(HstPtr, DevAccessiblePtr, Size);
}

Error PinnedAllocationMapTy::unregisterHostBuffer(void *HstPtr) {
  assert(HstPtr && "Invalid pointer");

  std::lock_guard<std::shared_mutex> Lock(Mutex);

  const EntryTy *Entry = findIntersecting(HstPtr);
  if (!Entry)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "cannot find locked buffer");

  // The address in the entry should be the same we are unregistering.
  if (Entry->HstPtr != HstPtr)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "unexpected host pointer in locked buffer entry");

  // Unregister from the entry.
  auto LastUseOrErr = unregisterEntryUse(*Entry);
  if (!LastUseOrErr)
    return LastUseOrErr.takeError();

  // There should be no other references to the pinned allocation.
  if (!(*LastUseOrErr))
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "the locked buffer is still being used");

  // Erase the entry from the map.
  return eraseEntry(*Entry);
}

Expected<void *> PinnedAllocationMapTy::registerMemory(void *HstPtr,
                                                       size_t Size,
                                                       bool LockMemory) {
  assert(HstPtr && "Invalid pointer");
  assert(Size && "Invalid size");

  std::lock_guard<std::shared_mutex> Lock(Mutex);

  const EntryTy *Entry = findIntersecting(HstPtr);

  if (Entry) {
    // An already registered intersecting buffer was found. Register a new use.
    if (auto Err = registerEntryUse(*Entry, HstPtr, Size))
      return std::move(Err);

    // Return the device accessible pointer with the correct offset.
    return utils::advancePtr(Entry->DevAccessiblePtr,
                             utils::getPtrDiff(HstPtr, Entry->HstPtr));
  }

  size_t BaseSize;
  void *BaseHstPtr, *BaseDevAccessiblePtr;

  // Check if it was externally pinned by a vendor-specific API.
  auto IsPinnedOrErr = Device.isPinnedPtrImpl(HstPtr, BaseHstPtr,
                                              BaseDevAccessiblePtr, BaseSize);
  if (!IsPinnedOrErr)
    return IsPinnedOrErr.takeError();

  // If pinned, just insert the entry representing the whole pinned buffer.
  if (*IsPinnedOrErr) {
    if (auto Err = insertEntry(BaseHstPtr, BaseDevAccessiblePtr, BaseSize,
                               /*Externallylocked=*/true))
      return std::move(Err);
    return BaseDevAccessiblePtr;
  }

  // Not externally pinned. Do nothing if locking of mapped buffers is disabled.
  if (!LockMemory)
    return nullptr;

  // No intersecting registered allocation found in the map. First, lock the
  // host buffer and retrieve the device accessible pointer.
  auto DevAccessiblePtrOrErr = Device.dataLockImpl(HstPtr, Size);
  if (!DevAccessiblePtrOrErr)
    return DevAccessiblePtrOrErr.takeError();

  // Now insert the new entry into the map.
  if (auto Err = insertEntry(HstPtr, *DevAccessiblePtrOrErr, Size))
    return std::move(Err);

  // Return the device accessible pointer.
  return *DevAccessiblePtrOrErr;
}

Error PinnedAllocationMapTy::unregisterMemory(void *HstPtr, bool UnlockMemory) {
  assert(HstPtr && "Invalid pointer");

  std::lock_guard<std::shared_mutex> Lock(Mutex);

  const EntryTy *Entry = findIntersecting(HstPtr);

  // No entry but automatic locking of mapped buffers is disabled, so
  // nothing to do.
  if (!Entry && !UnlockMemory)
    return Plugin::success();

  if (!Entry)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "cannot find locked buffer");

  // Unregister from the locked buffer. No need to do anything if there are
  // others using the allocation.
  auto LastUseOrErr = unregisterEntryUse(*Entry);
  if (!LastUseOrErr)
    return LastUseOrErr.takeError();

  // No need to do anything if there are others using the allocation.
  if (!(*LastUseOrErr))
    return Plugin::success();

  // This was the last user of the allocation. Unlock the original locked buffer
  // if it was locked by the plugin. Do not unlock it if it was locked by an
  // external entity. Unlock the buffer using the host pointer of the entry.
  if (!Entry->ExternallyLocked)
    if (auto Err = Device.dataUnlockImpl(Entry->HstPtr))
      return Err;

  // Erase the entry from the map.
  return eraseEntry(*Entry);
}

Error GenericDeviceTy::synchronize(__tgt_async_info *AsyncInfo,
                                   bool ReleaseQueue) {
  if (!AsyncInfo)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "invalid async info queue");

  SmallVector<void *> AllocsToDelete{};
  {
    std::lock_guard<std::mutex> AllocationGuard{AsyncInfo->Mutex};

    // This can be false when no work has been added to the AsyncInfo. In which
    // case, the device has nothing to synchronize.
    if (AsyncInfo->Queue)
      if (auto Err = synchronizeImpl(*AsyncInfo, ReleaseQueue))
        return Err;

    std::swap(AllocsToDelete, AsyncInfo->AssociatedAllocations);
  }

  for (auto *Ptr : AllocsToDelete)
    if (auto Err = dataDelete(Ptr, TargetAllocTy::TARGET_ALLOC_DEVICE))
      return Err;

  return Plugin::success();
}

Error GenericDeviceTy::queryAsync(__tgt_async_info *AsyncInfo,
                                  bool ReleaseQueue,
                                  bool *IsQueueWorkCompleted) {
  if (!AsyncInfo || !AsyncInfo->Queue)
    return Plugin::error(ErrorCode::INVALID_ARGUMENT,
                         "invalid async info queue");

  return queryAsyncImpl(*AsyncInfo, ReleaseQueue, IsQueueWorkCompleted);
}

Error GenericDeviceTy::memoryVAMap(void **Addr, void *VAddr, size_t *RSize) {
  return Plugin::error(ErrorCode::UNSUPPORTED,
                       "device does not support VA Management");
}

Error GenericDeviceTy::memoryVAUnMap(void *VAddr, size_t Size) {
  return Plugin::error(ErrorCode::UNSUPPORTED,
                       "device does not support VA Management");
}

Error GenericDeviceTy::getDeviceMemorySize(uint64_t &DSize) {
  return Plugin::error(
      ErrorCode::UNIMPLEMENTED,
      "missing getDeviceMemorySize implementation (required by RR-heuristic");
}

Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
                                            TargetAllocTy Kind) {
  void *Alloc = nullptr;

  if (RecordReplay && RecordReplay->isRecordingOrReplaying())
    return RecordReplay->allocate(Size);

  switch (Kind) {
  case TARGET_ALLOC_DEFAULT:
  case TARGET_ALLOC_DEVICE:
    if (MemoryManager) {
      auto AllocOrErr = MemoryManager->allocate(Size, HostPtr);
      if (!AllocOrErr)
        return AllocOrErr.takeError();
      Alloc = *AllocOrErr;
      if (!Alloc)
        return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
                             "failed to allocate from memory manager");
      break;
    }
    [[fallthrough]];
  case TARGET_ALLOC_HOST:
  case TARGET_ALLOC_SHARED: {
    auto AllocOrErr = allocate(Size, HostPtr, Kind);
    if (!AllocOrErr)
      return AllocOrErr.takeError();
    Alloc = *AllocOrErr;
    if (!Alloc)
      return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
                           "failed to allocate from device allocator");
  }
  }

  // Report error if the memory manager or the device allocator did not return
  // any memory buffer.
  if (!Alloc)
    return Plugin::error(ErrorCode::UNIMPLEMENTED,
                         "invalid target data allocation kind or requested "
                         "allocator not implemented yet");

  // Register allocated buffer as pinned memory if the type is host memory.
  if (Kind == TARGET_ALLOC_HOST)
    if (auto Err = PinnedAllocs.registerHostBuffer(Alloc, Alloc, Size))
      return std::move(Err);

  // Keep track of the allocation stack if we track allocation traces.
  if (OMPX_TrackAllocationTraces) {
    std::string StackTrace;
    llvm::raw_string_ostream OS(StackTrace);
    llvm::sys::PrintStackTrace(OS);

    AllocationTraceInfoTy *ATI = new AllocationTraceInfoTy();
    ATI->AllocationTrace = std::move(StackTrace);
    ATI->DevicePtr = Alloc;
    ATI->HostPtr = HostPtr;
    ATI->Size = Size;
    ATI->Kind = Kind;

    auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor();
    auto *&MapATI = (*AllocationTraceMap)[Alloc];
    ATI->LastAllocationInfo = MapATI;
    MapATI = ATI;
  }

  return Alloc;
}

Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) {
  // Free is a noop when recording or replaying.
  if (RecordReplay && RecordReplay->isRecordingOrReplaying())
    return RecordReplay->deallocate(TgtPtr);

  // Keep track of the deallocation stack if we track allocation traces.
  if (OMPX_TrackAllocationTraces) {
    AllocationTraceInfoTy *ATI = nullptr;
    {
      auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor();
      ATI = (*AllocationTraceMap)[TgtPtr];
    }

    std::string StackTrace;
    llvm::raw_string_ostream OS(StackTrace);
    llvm::sys::PrintStackTrace(OS);

    if (!ATI)
      ErrorReporter::reportDeallocationOfNonAllocatedPtr(TgtPtr, Kind, ATI,
                                                         StackTrace);

    // ATI is not null, thus we can lock it to inspect and modify it further.
    std::lock_guard<std::mutex> LG(ATI->Lock);
    if (!ATI->DeallocationTrace.empty())
      ErrorReporter::reportDeallocationOfDeallocatedPtr(TgtPtr, Kind, ATI,
                                                        StackTrace);

    if (ATI->Kind != Kind)
      ErrorReporter::reportDeallocationOfWrongPtrKind(TgtPtr, Kind, ATI,
                                                      StackTrace);

    ATI->DeallocationTrace = StackTrace;
  }

  switch (Kind) {
  case TARGET_ALLOC_DEFAULT:
  case TARGET_ALLOC_DEVICE:
    if (MemoryManager) {
      if (auto Err = MemoryManager->free(TgtPtr))
        return Err;
      break;
    }
    [[fallthrough]];
  case TARGET_ALLOC_HOST:
  case TARGET_ALLOC_SHARED:
    if (auto Err = free(TgtPtr, Kind))
      return Err;
  }

  // Unregister deallocated pinned memory buffer if the type is host memory.
  if (Kind == TARGET_ALLOC_HOST)
    if (auto Err = PinnedAllocs.unregisterHostBuffer(TgtPtr))
      return Err;

  return Plugin::success();
}

Error GenericDeviceTy::dataSubmit(void *TgtPtr, const void *HstPtr,
                                  int64_t Size, __tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);

  auto Err = dataSubmitImpl(TgtPtr, HstPtr, Size, AsyncInfoWrapper);
  AsyncInfoWrapper.finalize(Err);
  return Err;
}

Error GenericDeviceTy::dataRetrieve(void *HstPtr, const void *TgtPtr,
                                    int64_t Size, __tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);

  auto Err = dataRetrieveImpl(HstPtr, TgtPtr, Size, AsyncInfoWrapper);
  AsyncInfoWrapper.finalize(Err);
  return Err;
}

Error GenericDeviceTy::dataExchange(const void *SrcPtr, GenericDeviceTy &DstDev,
                                    void *DstPtr, int64_t Size,
                                    __tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);

  auto Err = dataExchangeImpl(SrcPtr, DstDev, DstPtr, Size, AsyncInfoWrapper);
  AsyncInfoWrapper.finalize(Err);
  return Err;
}

Error GenericDeviceTy::dataFill(void *TgtPtr, const void *PatternPtr,
                                int64_t PatternSize, int64_t Size,
                                __tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);
  auto Err =
      dataFillImpl(TgtPtr, PatternPtr, PatternSize, Size, AsyncInfoWrapper);
  AsyncInfoWrapper.finalize(Err);
  return Err;
}

Error GenericDeviceTy::launchKernel(void *EntryPtr, void **ArgPtrs,
                                    ptrdiff_t *ArgOffsets,
                                    KernelArgsTy &KernelArgs,
                                    KernelExtraArgsTy *KernelExtraArgs,
                                    __tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);

  GenericKernelTy &GenericKernel =
      *reinterpret_cast<GenericKernelTy *>(EntryPtr);

  {
    std::string StackTrace;
    if (OMPX_TrackNumKernelLaunches) {
      llvm::raw_string_ostream OS(StackTrace);
      llvm::sys::PrintStackTrace(OS);
    }

    auto KernelTraceInfoRecord = KernelLaunchTraces.getExclusiveAccessor();
    (*KernelTraceInfoRecord)
        .emplace(&GenericKernel, std::move(StackTrace), AsyncInfo);
  }

  auto Err = GenericKernel.launch(*this, ArgPtrs, ArgOffsets, KernelArgs,
                                  KernelExtraArgs, AsyncInfoWrapper);

  AsyncInfoWrapper.finalize(Err);

  return Err;
}

Error GenericDeviceTy::initAsyncInfo(__tgt_async_info **AsyncInfoPtr) {
  assert(AsyncInfoPtr && "Invalid async info");

  *AsyncInfoPtr = new __tgt_async_info();

  AsyncInfoWrapperTy AsyncInfoWrapper(*this, *AsyncInfoPtr);

  auto Err = initAsyncInfoImpl(AsyncInfoWrapper);
  AsyncInfoWrapper.finalize(Err);
  return Err;
}

Error GenericDeviceTy::enqueueHostCall(void (*Callback)(void *), void *UserData,
                                       __tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);

  auto Err = enqueueHostCallImpl(Callback, UserData, AsyncInfoWrapper);
  AsyncInfoWrapper.finalize(Err);
  return Err;
}

Expected<InfoTreeNode> GenericDeviceTy::obtainInfo() {
  auto InfoOrErr = obtainInfoImpl();
  if (InfoOrErr)
    InfoOrErr->add("UID", getDeviceUid(), "", DeviceInfo::UID);
  return InfoOrErr;
}

Error GenericDeviceTy::printInfo() {
  auto InfoOrErr = obtainInfo();

  // Get the vendor-specific info entries describing the device properties.
  if (auto Err = InfoOrErr.takeError())
    return Err;

  // Print all info entries.
  InfoOrErr->print();

  return Plugin::success();
}

Error GenericDeviceTy::createEvent(void **EventPtrStorage,
                                   bool EnableProfiling) {
  return createEventImpl(EventPtrStorage, EnableProfiling);
}

Error GenericDeviceTy::destroyEvent(void *EventPtr, bool EnableProfiling) {
  return destroyEventImpl(EventPtr, EnableProfiling);
}

Error GenericDeviceTy::recordEvent(void *EventPtr, __tgt_async_info *AsyncInfo,
                                   bool EnableProfiling) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);

  auto Err = recordEventImpl(EventPtr, AsyncInfoWrapper, EnableProfiling);
  AsyncInfoWrapper.finalize(Err);
  return Err;
}

Error GenericDeviceTy::waitEvent(void *EventPtr, __tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);

  auto Err = waitEventImpl(EventPtr, AsyncInfoWrapper);
  AsyncInfoWrapper.finalize(Err);
  return Err;
}

Expected<bool> GenericDeviceTy::hasPendingWork(__tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);
  auto Res = hasPendingWorkImpl(AsyncInfoWrapper);
  if (auto Err = Res.takeError()) {
    AsyncInfoWrapper.finalize(Err);
    return Err;
  }

  auto Err = Plugin::success();
  AsyncInfoWrapper.finalize(Err);
  if (Err)
    return Err;
  return Res;
}

Expected<bool> GenericDeviceTy::isEventComplete(void *Event,
                                                __tgt_async_info *AsyncInfo) {
  AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);
  auto Res = isEventCompleteImpl(Event, AsyncInfoWrapper);
  if (auto Err = Res.takeError()) {
    AsyncInfoWrapper.finalize(Err);
    return Err;
  }

  auto Err = Plugin::success();
  AsyncInfoWrapper.finalize(Err);
  if (Err)
    return Err;
  return Res;
}

Error GenericDeviceTy::syncEvent(void *EventPtr) {
  return syncEventImpl(EventPtr);
}

Expected<float> GenericDeviceTy::getEventElapsedTime(void *StartEventPtr,
                                                     void *EndEventPtr) {
  return getEventElapsedTimeImpl(StartEventPtr, EndEventPtr);
}

bool GenericDeviceTy::useAutoZeroCopy() { return useAutoZeroCopyImpl(); }

Expected<bool> GenericDeviceTy::isAccessiblePtr(const void *Ptr, size_t Size) {
  return isAccessiblePtrImpl(Ptr, Size);
}

void GenericDeviceTy::setDeviceUidFromVendorUid(StringRef VendorUid) {
  DeviceUid = std::string(Plugin.getName()) + "-" + std::string(VendorUid);
}

Error GenericPluginTy::init() {
  if (Initialized)
    return Plugin::success();

  auto NumDevicesOrErr = initImpl();
  if (!NumDevicesOrErr)
    return NumDevicesOrErr.takeError();
  Initialized = true;

  NumDevices = *NumDevicesOrErr;
  if (NumDevices == 0)
    return Plugin::success();

  assert(Devices.size() == 0 && "Plugin already initialized");
  Devices.resize(NumDevices, nullptr);

  GlobalHandler = createGlobalHandler();
  assert(GlobalHandler && "Invalid global handler");

  RPCServer = new RPCServerTy(*this);
  assert(RPCServer && "Invalid RPC server");

  return Plugin::success();
}

Error GenericPluginTy::deinit() {
  assert(Initialized && "Plugin was not initialized!");

  // Deinitialize all active devices.
  for (int32_t DeviceId = 0; DeviceId < NumDevices; ++DeviceId) {
    if (Devices[DeviceId]) {
      if (auto Err = deinitDevice(DeviceId))
        return Err;
    }
    assert(!Devices[DeviceId] && "Device was not deinitialized");
  }

  // There is no global handler if no device is available.
  if (GlobalHandler)
    delete GlobalHandler;

  if (RPCServer) {
    if (Error Err = RPCServer->shutDown(*this))
      return Err;
    delete RPCServer;
  }

  // Perform last deinitializations on the plugin.
  if (Error Err = deinitImpl())
    return Err;
  Initialized = false;

  return Plugin::success();
}

Error GenericPluginTy::initDevice(int32_t DeviceId) {
  assert(!Devices[DeviceId] && "Device already initialized");

  // Create the device and save the reference.
  GenericDeviceTy *Device = createDevice(*this, DeviceId, NumDevices);
  assert(Device && "Invalid device");

  // Save the device reference into the list.
  Devices[DeviceId] = Device;

  // Initialize the device and its resources.
  return Device->init(*this);
}

Error GenericPluginTy::deinitDevice(int32_t DeviceId) {
  // The device may be already deinitialized.
  if (Devices[DeviceId] == nullptr)
    return Plugin::success();

  // Deinitialize the device and release its resources.
  if (auto Err = Devices[DeviceId]->deinit(*this))
    return Err;

  // Delete the device and invalidate its reference.
  delete Devices[DeviceId];
  Devices[DeviceId] = nullptr;

  return Plugin::success();
}

Expected<bool> GenericPluginTy::checkELFImage(StringRef Image) const {
  // First check if this image is a regular ELF file.
  if (!utils::elf::isELF(Image))
    return false;

  // Check if this image is an ELF with a matching machine value.
  auto MachineOrErr = utils::elf::checkMachine(Image, getMagicElfBits());
  if (!MachineOrErr)
    return MachineOrErr.takeError();

  return MachineOrErr;
}

Expected<bool> GenericPluginTy::checkBitcodeImage(StringRef Image) const {
  if (identify_magic(Image) != file_magic::bitcode)
    return false;

  LLVMContext Context;
  auto ModuleOrErr = getLazyBitcodeModule(MemoryBufferRef(Image, ""), Context,
                                          /*ShouldLazyLoadMetadata=*/true);
  if (!ModuleOrErr)
    return ModuleOrErr.takeError();
  Module &M = **ModuleOrErr;

  return M.getTargetTriple().getArch() == getTripleArch();
}

int32_t GenericPluginTy::is_initialized() const { return Initialized; }

int32_t GenericPluginTy::isPluginCompatible(StringRef Image) {
  auto HandleError = [&](Error Err) -> bool {
    std::string ErrStr = toString(std::move(Err));
    ODBG(OLDT_Init) << "Failure to check validity of image "
                    << static_cast<const void *>(Image.data()) << ": "
                    << ErrStr;
    return false;
  };
  switch (identify_magic(Image)) {
  case file_magic::elf:
  case file_magic::elf_relocatable:
  case file_magic::elf_executable:
  case file_magic::elf_shared_object:
  case file_magic::elf_core: {
    auto MatchOrErr = checkELFImage(Image);
    if (Error Err = MatchOrErr.takeError())
      return HandleError(std::move(Err));
    return *MatchOrErr;
  }
  case file_magic::bitcode: {
    auto MatchOrErr = checkBitcodeImage(Image);
    if (Error Err = MatchOrErr.takeError())
      return HandleError(std::move(Err));
    return *MatchOrErr;
  }
  default:
    auto MatchOrErr = isImageCompatible(Image);
    if (Error Err = MatchOrErr.takeError())
      return HandleError(std::move(Err));
    return *MatchOrErr;
  }
}

int32_t GenericPluginTy::isDeviceCompatible(int32_t DeviceId, StringRef Image) {
  auto HandleError = [&](Error Err) -> bool {
    std::string ErrStr = toString(std::move(Err));
    ODBG(OLDT_Init) << "Failure to check validity of image "
                    << static_cast<const void *>(Image.data()) << ": "
                    << ErrStr;
    return false;
  };
  switch (identify_magic(Image)) {
  case file_magic::elf:
  case file_magic::elf_relocatable:
  case file_magic::elf_executable:
  case file_magic::elf_shared_object:
  case file_magic::elf_core: {
    auto MatchOrErr = checkELFImage(Image);
    if (Error Err = MatchOrErr.takeError())
      return HandleError(std::move(Err));
    if (!*MatchOrErr)
      return false;

    // Perform plugin-dependent checks for the specific architecture if needed.
    auto CompatibleOrErr = isELFCompatible(DeviceId, Image);
    if (Error Err = CompatibleOrErr.takeError())
      return HandleError(std::move(Err));
    return *CompatibleOrErr;
  }
  case file_magic::bitcode: {
    auto MatchOrErr = checkBitcodeImage(Image);
    if (Error Err = MatchOrErr.takeError())
      return HandleError(std::move(Err));
    return *MatchOrErr;
  }
  default:
    auto MatchOrErr = isImageCompatible(DeviceId, Image);
    if (Error Err = MatchOrErr.takeError())
      return HandleError(std::move(Err));
    return *MatchOrErr;
  }
}

int32_t GenericPluginTy::is_device_initialized(int32_t DeviceId) const {
  return isValidDeviceId(DeviceId) && Devices[DeviceId] != nullptr;
}

int32_t GenericPluginTy::init_device(int32_t DeviceId) {
  auto Err = initDevice(DeviceId);
  if (Err) {
    REPORT() << "Failure to initialize device " << DeviceId << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::number_of_devices() { return getNumDevices(); }

int32_t GenericPluginTy::is_data_exchangable(int32_t SrcDeviceId,
                                             int32_t DstDeviceId) {
  return isDataExchangable(SrcDeviceId, DstDeviceId);
}

int32_t GenericPluginTy::initialize_record_replay(
    int32_t DeviceId, int64_t MemorySize, void *VAddr, bool IsRecord,
    bool IsNative, bool SaveOutput, bool EmitReport,
    const char *OutputDirPath) {
  GenericDeviceTy &Device = getDevice(DeviceId);

  if (auto Err =
          Device.initRecordReplay(MemorySize, VAddr, IsRecord, IsNative,
                                  SaveOutput, EmitReport, OutputDirPath)) {
    REPORT() << "Failure to initialize RR with " << MemorySize
             << " bytes on device " << DeviceId << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }
  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::load_binary(int32_t DeviceId,
                                     __tgt_device_image *TgtImage,
                                     __tgt_device_binary *Binary) {
  GenericDeviceTy &Device = getDevice(DeviceId);

  StringRef Buffer(reinterpret_cast<const char *>(TgtImage->ImageStart),
                   utils::getPtrDiff(TgtImage->ImageEnd, TgtImage->ImageStart));
  auto ImageOrErr = Device.loadBinary(*this, Buffer);
  if (!ImageOrErr) {
    auto Err = ImageOrErr.takeError();
    REPORT() << "Failure to load binary image " << TgtImage << " on device "
             << DeviceId << ": " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  DeviceImageTy *Image = *ImageOrErr;
  assert(Image != nullptr && "Invalid Image");

  *Binary = __tgt_device_binary{reinterpret_cast<uint64_t>(Image)};

  return OFFLOAD_SUCCESS;
}

void *GenericPluginTy::data_alloc(int32_t DeviceId, int64_t Size, void *HostPtr,
                                  int32_t Kind) {
  auto AllocOrErr =
      getDevice(DeviceId).dataAlloc(Size, HostPtr, (TargetAllocTy)Kind);
  if (!AllocOrErr) {
    auto Err = AllocOrErr.takeError();
    REPORT() << "Failure to allocate device memory: "
             << toString(std::move(Err));
    return nullptr;
  }
  assert(*AllocOrErr && "Null pointer upon successful allocation");

  return *AllocOrErr;
}

int32_t GenericPluginTy::data_delete(int32_t DeviceId, void *TgtPtr,
                                     int32_t Kind) {
  auto Err =
      getDevice(DeviceId).dataDelete(TgtPtr, static_cast<TargetAllocTy>(Kind));
  if (Err) {
    REPORT() << "Failure to deallocate device pointer " << TgtPtr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::data_lock(int32_t DeviceId, void *Ptr, int64_t Size,
                                   void **LockedPtr) {
  auto LockedPtrOrErr = getDevice(DeviceId).registerMemory(Ptr, Size);
  if (!LockedPtrOrErr) {
    auto Err = LockedPtrOrErr.takeError();
    REPORT() << "Failure to lock memory " << Ptr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  if (!(*LockedPtrOrErr)) {
    REPORT() << "Failure to lock memory " << Ptr
             << ": obtained a null locked pointer";
    return OFFLOAD_FAIL;
  }
  *LockedPtr = *LockedPtrOrErr;

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::data_unlock(int32_t DeviceId, void *Ptr) {
  auto Err = getDevice(DeviceId).unregisterMemory(Ptr);
  if (Err) {
    REPORT() << "Failure to unlock memory " << Ptr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::data_notify_mapped(int32_t DeviceId, void *HstPtr,
                                            int64_t Size) {
  auto Err = getDevice(DeviceId).notifyDataMapped(HstPtr, Size);
  if (Err) {
    REPORT() << "Failure to notify data mapped " << HstPtr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::data_notify_unmapped(int32_t DeviceId, void *HstPtr) {
  auto Err = getDevice(DeviceId).notifyDataUnmapped(HstPtr);
  if (Err) {
    REPORT() << "Failure to notify data unmapped " << HstPtr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::data_submit(int32_t DeviceId, void *TgtPtr,
                                     void *HstPtr, int64_t Size) {
  return data_submit_async(DeviceId, TgtPtr, HstPtr, Size,
                           /*AsyncInfoPtr=*/nullptr);
}

int32_t GenericPluginTy::data_submit_async(int32_t DeviceId, void *TgtPtr,
                                           void *HstPtr, int64_t Size,
                                           __tgt_async_info *AsyncInfoPtr) {
  auto Err = getDevice(DeviceId).dataSubmit(TgtPtr, HstPtr, Size, AsyncInfoPtr);
  if (Err) {
    REPORT() << "Failure to copy data from host to device. Pointers: host "
             << "= " << HstPtr << ", device = " << TgtPtr << ", size = " << Size
             << ": " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::data_retrieve(int32_t DeviceId, void *HstPtr,
                                       void *TgtPtr, int64_t Size) {
  return data_retrieve_async(DeviceId, HstPtr, TgtPtr, Size,
                             /*AsyncInfoPtr=*/nullptr);
}

int32_t GenericPluginTy::data_retrieve_async(int32_t DeviceId, void *HstPtr,
                                             void *TgtPtr, int64_t Size,
                                             __tgt_async_info *AsyncInfoPtr) {
  auto Err =
      getDevice(DeviceId).dataRetrieve(HstPtr, TgtPtr, Size, AsyncInfoPtr);
  if (Err) {
    REPORT() << "Failure to copy data from device to host. Pointers: host "
             << "= " << HstPtr << ", device = " << TgtPtr << ", size = " << Size
             << ": " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::data_exchange(int32_t SrcDeviceId, void *SrcPtr,
                                       int32_t DstDeviceId, void *DstPtr,
                                       int64_t Size) {
  return data_exchange_async(SrcDeviceId, SrcPtr, DstDeviceId, DstPtr, Size,
                             /*AsyncInfoPtr=*/nullptr);
}

int32_t GenericPluginTy::data_exchange_async(int32_t SrcDeviceId, void *SrcPtr,
                                             int DstDeviceId, void *DstPtr,
                                             int64_t Size,
                                             __tgt_async_info *AsyncInfo) {
  GenericDeviceTy &SrcDevice = getDevice(SrcDeviceId);
  GenericDeviceTy &DstDevice = getDevice(DstDeviceId);
  auto Err = SrcDevice.dataExchange(SrcPtr, DstDevice, DstPtr, Size, AsyncInfo);
  if (Err) {
    REPORT() << "Failure to copy data from device (" << SrcDeviceId
             << ") to device (" << DstDeviceId
             << "). Pointers: host = " << SrcPtr << ", device = " << DstPtr
             << ", size = " << Size << ": " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::launch_kernel(int32_t DeviceId, void *TgtEntryPtr,
                                       void **TgtArgs, ptrdiff_t *TgtOffsets,
                                       KernelArgsTy *KernelArgs,
                                       KernelExtraArgsTy *KernelExtraArgs,
                                       __tgt_async_info *AsyncInfoPtr) {
  auto Err = getDevice(DeviceId).launchKernel(TgtEntryPtr, TgtArgs, TgtOffsets,
                                              *KernelArgs, KernelExtraArgs,
                                              AsyncInfoPtr);
  if (Err) {
    REPORT() << "Failure to run target region " << TgtEntryPtr << " in device "
             << DeviceId << ": " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::synchronize(int32_t DeviceId,
                                     __tgt_async_info *AsyncInfoPtr) {
  auto Err = getDevice(DeviceId).synchronize(AsyncInfoPtr);
  if (Err) {
    REPORT() << "Failure to synchronize stream " << AsyncInfoPtr->Queue << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::query_async(int32_t DeviceId,
                                     __tgt_async_info *AsyncInfoPtr) {
  auto Err = getDevice(DeviceId).queryAsync(AsyncInfoPtr);
  if (Err) {
    REPORT() << "Failure to query stream " << AsyncInfoPtr->Queue << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

InfoTreeNode GenericPluginTy::obtain_device_info(int32_t DeviceId) {
  auto InfoOrErr = getDevice(DeviceId).obtainInfo();
  if (auto Err = InfoOrErr.takeError()) {
    REPORT() << "Failure to obtain device " << DeviceId
             << " info: " << toString(std::move(Err));
    return InfoTreeNode{};
  }
  return std::move(*InfoOrErr);
}

void GenericPluginTy::print_device_info(int32_t DeviceId) {
  if (auto Err = getDevice(DeviceId).printInfo())
    REPORT() << "Failure to print device " << DeviceId
             << " info: " << toString(std::move(Err));
}

int32_t GenericPluginTy::create_event(int32_t DeviceId, void **EventPtr) {
  auto Err = getDevice(DeviceId).createEvent(EventPtr);
  if (Err) {
    REPORT() << "Failure to create event: " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::record_event(int32_t DeviceId, void *EventPtr,
                                      __tgt_async_info *AsyncInfoPtr) {
  auto Err = getDevice(DeviceId).recordEvent(EventPtr, AsyncInfoPtr);
  if (Err) {
    REPORT() << "Failure to record event " << EventPtr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::wait_event(int32_t DeviceId, void *EventPtr,
                                    __tgt_async_info *AsyncInfoPtr) {
  auto Err = getDevice(DeviceId).waitEvent(EventPtr, AsyncInfoPtr);
  if (Err) {
    REPORT() << "Failure to wait event " << EventPtr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::sync_event(int32_t DeviceId, void *EventPtr) {
  auto Err = getDevice(DeviceId).syncEvent(EventPtr);
  if (Err) {
    REPORT() << "Failure to synchronize event " << EventPtr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::get_event_elapsed_time(int32_t DeviceId,
                                                void *StartEventPtr,
                                                void *EndEventPtr,
                                                float *ElapsedTime) {
  auto ElapsedTimeOrErr =
      getDevice(DeviceId).getEventElapsedTime(StartEventPtr, EndEventPtr);
  if (!ElapsedTimeOrErr) {
    REPORT() << "Failure to get elapsed time between events " << StartEventPtr
             << " and " << EndEventPtr << ": "
             << toString(ElapsedTimeOrErr.takeError());
    return OFFLOAD_FAIL;
  }

  *ElapsedTime = *ElapsedTimeOrErr;
  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::destroy_event(int32_t DeviceId, void *EventPtr) {
  auto Err = getDevice(DeviceId).destroyEvent(EventPtr);
  if (Err) {
    REPORT() << "Failure to destroy event " << EventPtr << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

void GenericPluginTy::set_info_flag(uint32_t NewInfoLevel) {
  std::atomic<uint32_t> &InfoLevel = getInfoLevelInternal();
  InfoLevel.store(NewInfoLevel);
}

int32_t GenericPluginTy::init_async_info(int32_t DeviceId,
                                         __tgt_async_info **AsyncInfoPtr) {
  assert(AsyncInfoPtr && "Invalid async info");

  auto Err = getDevice(DeviceId).initAsyncInfo(AsyncInfoPtr);
  if (Err) {
    REPORT() << "Failure to initialize async info at " << *AsyncInfoPtr
             << " on device " << DeviceId << ": " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::set_device_identifier(int32_t UserId,
                                               int32_t DeviceId) {
  UserDeviceIds[DeviceId] = UserId;

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::use_auto_zero_copy(int32_t DeviceId) {
  return getDevice(DeviceId).useAutoZeroCopy();
}

int32_t GenericPluginTy::is_accessible_ptr(int32_t DeviceId, const void *Ptr,
                                           size_t Size) {
  auto HandleError = [&](Error Err) -> bool {
    std::string ErrStr = toString(std::move(Err));
    ODBG(OLDT_Device) << "Failure while checking accessibility of pointer "
                      << Ptr << " for device " << DeviceId << ": " << ErrStr;
    return false;
  };

  auto AccessibleOrErr = getDevice(DeviceId).isAccessiblePtr(Ptr, Size);
  if (Error Err = AccessibleOrErr.takeError())
    return HandleError(std::move(Err));

  return *AccessibleOrErr;
}

int32_t GenericPluginTy::get_global(__tgt_device_binary Binary, uint64_t Size,
                                    const char *Name, void **DevicePtr) {
  assert(Binary.handle && "Invalid device binary handle");
  DeviceImageTy &Image = *reinterpret_cast<DeviceImageTy *>(Binary.handle);

  GenericDeviceTy &Device = Image.getDevice();

  GlobalTy DeviceGlobal(Name, Size);
  GenericGlobalHandlerTy &GHandler = getGlobalHandler();
  if (auto Err =
          GHandler.getGlobalMetadataFromDevice(Device, Image, DeviceGlobal)) {
    consumeError(std::move(Err));
    return OFFLOAD_FAIL;
  }

  *DevicePtr = DeviceGlobal.getPtr();
  assert(DevicePtr && "Invalid device global's address");

  // Save the loaded globals if we are recording.
  RecordReplayTy *RecordReplay = Device.getRecordReplay();
  if (RecordReplay && RecordReplay->isRecording())
    RecordReplay->addGlobal(Name, Size, *DevicePtr);

  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::get_function(__tgt_device_binary Binary,
                                      const char *Name, void **KernelPtr) {
  assert(Binary.handle && "Invalid device binary handle");
  DeviceImageTy &Image = *reinterpret_cast<DeviceImageTy *>(Binary.handle);

  GenericDeviceTy &Device = Image.getDevice();

  auto KernelOrErr = Device.constructKernel(Name);
  if (Error Err = KernelOrErr.takeError()) {
    REPORT() << "Failure to look up kernel: " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  GenericKernelTy &Kernel = *KernelOrErr;
  if (auto Err = Kernel.init(Device, Image)) {
    REPORT() << "Failure to init kernel: " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  // Note that this is not the kernel's device address.
  *KernelPtr = &Kernel;
  return OFFLOAD_SUCCESS;
}

/// Create OpenMP interop with the given interop context
omp_interop_val_t *
GenericPluginTy::create_interop(int32_t ID, int32_t InteropContext,
                                interop_spec_t *InteropSpec) {
  assert(InteropSpec && "Interop spec is null");
  auto &Device = getDevice(ID);
  auto InteropOrErr = Device.createInterop(InteropContext, *InteropSpec);
  if (!InteropOrErr) {
    REPORT() << "Failure to create interop object for device " << InteropSpec
             << ": " << toString(InteropOrErr.takeError());
    return nullptr;
  }
  return *InteropOrErr;
}

/// Release OpenMP interop object
int32_t GenericPluginTy::release_interop(int32_t ID,
                                         omp_interop_val_t *Interop) {
  assert(Interop && "Interop is null");
  assert(Interop->device_id == ID && "Interop does not match device id");
  auto &Device = getDevice(ID);
  auto Err = Device.releaseInterop(Interop);
  if (Err) {
    REPORT() << "Failure to release interop object " << Interop << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }
  return OFFLOAD_SUCCESS;
}

/// Flush the queue associated with the interop object if necessary
int32_t GenericPluginTy::flush_queue(omp_interop_val_t *Interop) {
  assert(Interop && "Interop is null");
  auto Err = flushQueueImpl(Interop);
  if (Err) {
    REPORT() << "Failure to flush interop object " << Interop
             << " queue: " << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }
  return OFFLOAD_SUCCESS;
}

/// Perform a host synchronization with the queue associated with the interop
/// object and wait for it to complete.
int32_t GenericPluginTy::sync_barrier(omp_interop_val_t *Interop) {
  assert(Interop && "Interop is null");
  auto Err = syncBarrierImpl(Interop);
  if (Err) {
    REPORT() << "Failure to synchronize interop object " << Interop << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }
  return OFFLOAD_SUCCESS;
}

/// Queue an asynchronous barrier in the queue associated with the interop
/// object and return immediately.
int32_t GenericPluginTy::async_barrier(omp_interop_val_t *Interop) {
  assert(Interop && "Interop is null");
  auto Err = asyncBarrierImpl(Interop);
  if (Err) {
    REPORT() << "Failure to queue barrier in interop object " << Interop << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }
  return OFFLOAD_SUCCESS;
}

int32_t GenericPluginTy::data_fence(int32_t DeviceId,
                                    __tgt_async_info *AsyncInfo) {
  auto Err = getDevice(DeviceId).dataFence(AsyncInfo);
  if (Err) {
    REPORT() << "Failure to place data fence on device " << DeviceId << ": "
             << toString(std::move(Err));
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}