//===-- OpenMP/Mapping.cpp - OpenMP/OpenACC pointer mapping impl. ---------===//
//
// 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 "OpenMP/Mapping.h"

#include "PluginManager.h"
#include "Shared/Debug.h"
#include "Shared/Requirements.h"
#include "device.h"

using namespace llvm::omp::target::debug;

/// Dump a table of all the host-target pointer pairs on failure
void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device,
                               bool toStdOut) {
  MappingInfoTy::HDTTMapAccessorTy HDTTMap =
      Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor();
  if (HDTTMap->empty()) {
    DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID,
              "OpenMP Host-Device pointer mappings table empty\n");
    return;
  }

  SourceInfo Kernel(Loc);
  DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID,
            "OpenMP Host-Device pointer mappings after block at %s:%d:%d:\n",
            Kernel.getFilename(), Kernel.getLine(), Kernel.getColumn());
  DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID,
            "%-18s %-18s %s %s %s %s\n", "Host Ptr", "Target Ptr", "Size (B)",
            "DynRefCount", "HoldRefCount", "Declaration");
  for (const auto &It : *HDTTMap) {
    HostDataToTargetTy &HDTT = *It.HDTT;
    SourceInfo Info(HDTT.HstPtrName);
    DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID,
              DPxMOD " " DPxMOD " %-8" PRIuPTR " %-11s %-12s %s at %s:%d:%d\n",
              DPxPTR(HDTT.HstPtrBegin), DPxPTR(HDTT.TgtPtrBegin),
              HDTT.HstPtrEnd - HDTT.HstPtrBegin,
              HDTT.dynRefCountToStr().c_str(), HDTT.holdRefCountToStr().c_str(),
              Info.getName(), Info.getFilename(), Info.getLine(),
              Info.getColumn());
  }
}

int MappingInfoTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin,
                                int64_t Size) {
  HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();

  // Check if entry exists
  auto It = HDTTMap->find(HstPtrBegin);
  if (It != HDTTMap->end()) {
    HostDataToTargetTy &HDTT = *It->HDTT;
    std::lock_guard<HostDataToTargetTy> LG(HDTT);
    // Mapping already exists
    bool IsValid = HDTT.HstPtrEnd == (uintptr_t)HstPtrBegin + Size &&
                   HDTT.TgtPtrBegin == (uintptr_t)TgtPtrBegin;
    if (IsValid) {
      ODBG(ODT_MappingExists) << "Attempt to re-associate the same device "
                              << "ptr+offset with the same "
                              << "host ptr, nothing to do";
      return OFFLOAD_SUCCESS;
    }
    REPORT() << "Not allowed to re-associate a different device ptr+offset "
             << "with the same host ptr";
    return OFFLOAD_FAIL;
  }

  // Mapping does not exist, allocate it with refCount=INF
  const HostDataToTargetTy &NewEntry =
      *HDTTMap
           ->emplace(new HostDataToTargetTy(
               /*HstPtrBase=*/(uintptr_t)HstPtrBegin,
               /*HstPtrBegin=*/(uintptr_t)HstPtrBegin,
               /*HstPtrEnd=*/(uintptr_t)HstPtrBegin + Size,
               /*TgtAllocBegin=*/(uintptr_t)TgtPtrBegin,
               /*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin,
               /*UseHoldRefCount=*/false, /*Name=*/nullptr,
               /*IsRefCountINF=*/true))
           .first->HDTT;
  ODBG(ODT_Mapping) << "Creating new map entry: HstBase="
                    << reinterpret_cast<void *>(NewEntry.HstPtrBase)
                    << ", HstBegin="
                    << reinterpret_cast<void *>(NewEntry.HstPtrBegin)
                    << ", HstEnd="
                    << reinterpret_cast<void *>(NewEntry.HstPtrEnd)
                    << ", TgtBegin="
                    << reinterpret_cast<void *>(NewEntry.TgtPtrBegin)
                    << ", DynRefCount=" << NewEntry.dynRefCountToStr()
                    << ", HoldRefCount=" << NewEntry.holdRefCountToStr();
  (void)NewEntry;

  // Notify the plugin about the new mapping.
  return Device.notifyDataMapped(HstPtrBegin, Size);
}

int MappingInfoTy::disassociatePtr(void *HstPtrBegin) {
  HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();

  auto It = HDTTMap->find(HstPtrBegin);
  if (It == HDTTMap->end()) {
    REPORT() << "Association not found";
    return OFFLOAD_FAIL;
  }
  // Mapping exists
  HostDataToTargetTy &HDTT = *It->HDTT;
  std::lock_guard<HostDataToTargetTy> LG(HDTT);

  if (HDTT.getHoldRefCount()) {
    // This is based on OpenACC 3.1, sec 3.2.33 "acc_unmap_data", L3656-3657:
    // "It is an error to call acc_unmap_data if the structured reference
    // count for the pointer is not zero."
    REPORT() << "Trying to disassociate a pointer with a non-zero "
             << "hold reference count";
    return OFFLOAD_FAIL;
  }

  if (HDTT.isDynRefCountInf()) {
    ODBG(ODT_Mapping) << "Association found, removing it";
    void *Event = HDTT.getEvent();
    delete &HDTT;
    if (Event)
      Device.destroyEvent(Event);
    HDTTMap->erase(It);
    return Device.notifyDataUnmapped(HstPtrBegin);
  }

  REPORT() << "Trying to disassociate a pointer which was not mapped via "
           << "omp_target_associate_ptr";
  return OFFLOAD_FAIL;
}

LookupResult MappingInfoTy::lookupMapping(HDTTMapAccessorTy &HDTTMap,
                                          void *HstPtrBegin, int64_t Size,
                                          HostDataToTargetTy *OwnedTPR) {

  uintptr_t HP = (uintptr_t)HstPtrBegin;
  LookupResult LR;

  ODBG(ODT_Mapping) << "Looking up mapping(HstPtrBegin=" << HstPtrBegin
                    << ", Size=" << Size << ")...";

  if (HDTTMap->empty())
    return LR;

  // HDTTMap is std::set, ordered by HstPtrBegin.
  // Upper is the first element whose HstPtrBegin > HP.
  auto Upper = HDTTMap->upper_bound(HP);

  if (Size == 0) {
    // HP satisfies
    //   std::prev(Upper)->HDTT.HstPtrBegin <= HP < Upper->HDTT.HstPtrBegin
    if (Upper != HDTTMap->begin()) {
      LR.TPR.setEntry(std::prev(Upper)->HDTT, OwnedTPR);
      // We know that HP >= LR.TPR.getEntry()->HstPtrBegin
      LR.Flags.IsContained = HP < LR.TPR.getEntry()->HstPtrEnd;
    }

    if (!LR.Flags.IsContained && Upper != HDTTMap->end()) {
      LR.TPR.setEntry(Upper->HDTT, OwnedTPR);
      // This is a special case: HP is not really contained in the mapped
      // address range, but it's contained in the extended address range,
      // which suffices to get the mapping of the base pointer.
      // We know that HP < LR.TPR.getEntry()->HstPtrBegin
      LR.Flags.IsContained = HP >= LR.TPR.getEntry()->HstPtrBase;
    }
  } else {
    if (Upper != HDTTMap->begin()) {
      LR.TPR.setEntry(std::prev(Upper)->HDTT, OwnedTPR);
      // We know that HP >= LR.TPR.getEntry()->HstPtrBegin
      LR.Flags.IsContained = HP < LR.TPR.getEntry()->HstPtrEnd &&
                             (HP + Size) <= LR.TPR.getEntry()->HstPtrEnd;
      // Does it extend beyond the mapped address range?
      LR.Flags.ExtendsAfter = HP < LR.TPR.getEntry()->HstPtrEnd &&
                              (HP + Size) > LR.TPR.getEntry()->HstPtrEnd;
    }

    if (!(LR.Flags.IsContained || LR.Flags.ExtendsAfter) &&
        Upper != HDTTMap->end()) {
      LR.TPR.setEntry(Upper->HDTT, OwnedTPR);
      // Does it extend into an already mapped address range?
      // We know that HP < LR.TPR.getEntry()->HstPtrBegin
      LR.Flags.ExtendsBefore = (HP + Size) > LR.TPR.getEntry()->HstPtrBegin;
      // Does it extend beyond the mapped address range?
      LR.Flags.ExtendsAfter = HP < LR.TPR.getEntry()->HstPtrEnd &&
                              (HP + Size) > LR.TPR.getEntry()->HstPtrEnd;
    }

    if (LR.Flags.ExtendsBefore) {
      ODBG(ODT_Mapping) << "WARNING: Pointer is not mapped but section extends "
                        << "into already mapped data";
    }
    if (LR.Flags.ExtendsAfter) {
      ODBG(ODT_Mapping) << "WARNING: Pointer is already mapped but section "
                        << "extends beyond mapped region";
    }
  }

  return LR;
}

TargetPointerResultTy MappingInfoTy::getTargetPointer(
    HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin, void *HstPtrBase,
    int64_t TgtPadding, int64_t Size, map_var_info_t HstPtrName, bool HasFlagTo,
    bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
    bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
    AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap,
    StateInfoTy *StateInfo) {

  LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size, OwnedTPR);
  LR.TPR.Flags.IsPresent = true;

  // Release the mapping table lock only after the entry is locked by
  // attaching it to TPR. Once TPR is destroyed it will release the lock
  // on entry. If it is returned the lock will move to the returned object.
  // If LR.Entry is already owned/locked we avoid trying to lock it again.

  // Check if the pointer is contained.
  // If a variable is mapped to the device manually by the user - which would
  // lead to the IsContained flag to be true - then we must ensure that the
  // device address is returned even under unified memory conditions.
  if (LR.Flags.IsContained ||
      ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && IsImplicit)) {
    const char *RefCountAction;
    if (UpdateRefCount) {
      // After this, reference count >= 1. If the reference count was 0 but the
      // entry was still there we can reuse the data on the device and avoid a
      // new submission.
      LR.TPR.getEntry()->incRefCount(HasHoldModifier);
      RefCountAction = " (incremented)";
    } else {
      // It might have been allocated with the parent, but it's still new.
      LR.TPR.Flags.IsNewEntry = LR.TPR.getEntry()->getTotalRefCount() == 1;
      RefCountAction = " (update suppressed)";
    }
    const char *DynRefCountAction = HasHoldModifier ? "" : RefCountAction;
    const char *HoldRefCountAction = HasHoldModifier ? RefCountAction : "";
    uintptr_t Ptr = LR.TPR.getEntry()->TgtPtrBegin +
                    ((uintptr_t)HstPtrBegin - LR.TPR.getEntry()->HstPtrBegin);
    INFO(OMP_INFOTYPE_MAPPING_EXISTS, Device.DeviceID,
         "Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
         ", Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n",
         (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(Ptr),
         Size, LR.TPR.getEntry()->dynRefCountToStr().c_str(), DynRefCountAction,
         LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction,
         (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
    LR.TPR.TargetPointer = (void *)Ptr;
  } else if ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && !IsImplicit) {
    // Explicit extension of mapped data - not allowed.
    MESSAGE("explicit extension not allowed: host address specified is " DPxMOD
            " (%" PRId64
            " bytes), but device allocation maps to host at " DPxMOD
            " (%" PRId64 " bytes)",
            DPxPTR(HstPtrBegin), Size, DPxPTR(LR.TPR.getEntry()->HstPtrBegin),
            LR.TPR.getEntry()->HstPtrEnd - LR.TPR.getEntry()->HstPtrBegin);
    if (HasPresentModifier)
      MESSAGE("device mapping required by 'present' map type modifier does not "
              "exist for host address " DPxMOD " (%" PRId64 " bytes)",
              DPxPTR(HstPtrBegin), Size);
  } else if ((PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY &&
              !HasCloseModifier) ||
             (PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY)) {

    // If unified shared memory is active, implicitly mapped variables that are
    // not privatized use host address. Any explicitly mapped variables also use
    // host address where correctness is not impeded. In all other cases maps
    // are respected.
    // In addition to the mapping rules above, the close map modifier forces the
    // mapping of the variable to the device.
    if (Size) {
      INFO(OMP_INFOTYPE_MAPPING_CHANGED, Device.DeviceID,
           "Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
           "memory\n",
           DPxPTR((uintptr_t)HstPtrBegin), Size);
      ODBG(ODT_Mapping) << "Return HstPtrBegin " << HstPtrBegin
                        << " Size=" << Size << " for unified shared memory";
      LR.TPR.Flags.IsPresent = false;
      LR.TPR.Flags.IsHostPointer = true;
      LR.TPR.TargetPointer = HstPtrBegin;
    }
  } else if (HasPresentModifier) {
    ODBG(ODT_Mapping) << "Mapping required by 'present' map type modifier does "
                      << "not exist for HstPtrBegin=" << HstPtrBegin
                      << ", Size=" << Size;
    MESSAGE("device mapping required by 'present' map type modifier does not "
            "exist for host address " DPxMOD " (%" PRId64 " bytes)",
            DPxPTR(HstPtrBegin), Size);
  } else if (Size) {
    // If it is not contained and Size > 0, we should create a new entry for it.
    LR.TPR.Flags.IsNewEntry = true;
    uintptr_t TgtAllocBegin =
        (uintptr_t)Device.allocData(TgtPadding + Size, HstPtrBegin);
    uintptr_t TgtPtrBegin = TgtAllocBegin + TgtPadding;
    // Release the mapping table lock only after the entry is locked by
    // attaching it to TPR.
    LR.TPR.setEntry(HDTTMap
                        ->emplace(new HostDataToTargetTy(
                            (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
                            (uintptr_t)HstPtrBegin + Size, TgtAllocBegin,
                            TgtPtrBegin, HasHoldModifier, HstPtrName))
                        .first->HDTT);
    INFO(OMP_INFOTYPE_MAPPING_CHANGED, Device.DeviceID,
         "Creating new map entry with HstPtrBase=" DPxMOD
         ", HstPtrBegin=" DPxMOD ", TgtAllocBegin=" DPxMOD
         ", TgtPtrBegin=" DPxMOD
         ", Size=%ld, DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
         DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(TgtAllocBegin),
         DPxPTR(TgtPtrBegin), Size,
         LR.TPR.getEntry()->dynRefCountToStr().c_str(),
         LR.TPR.getEntry()->holdRefCountToStr().c_str(),
         (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
    LR.TPR.TargetPointer = (void *)TgtPtrBegin;

    // Notify the plugin about the new mapping.
    if (Device.notifyDataMapped(HstPtrBegin, Size))
      return TargetPointerResultTy{};
  } else {
    // This entry is not present and we did not create a new entry for it.
    LR.TPR.Flags.IsPresent = false;
  }

  // All mapping table modifications have been made. If the user requested it we
  // give up the lock.
  if (ReleaseHDTTMap)
    HDTTMap.destroy();

  // Lambda to check if this pointer was newly allocated on the current region.
  // This is needed to handle cases when the TO entry is encountered after an
  // alloc entry for the same pointer. In such cases, the ref-count is already
  // non-zero when TO is encountered, but we still need to do a transfer. e.g.
  //
  // struct S {
  //   int *p;
  // };
  // #pragma omp declare mapper(id : S s) map(to: s.p, s.p[0 : 10])
  //
  // S s1;
  // ...
  // #pragma omp target map(alloc : s1.p[0 : 10]) map(mapper(id), to : s1)
  auto WasNewlyAllocatedForCurrentRegion = [&]() {
    if (!StateInfo)
      return false;
    bool WasNewlyAllocated =
        StateInfo->wasNewlyAllocated(HstPtrBegin).has_value();
    if (WasNewlyAllocated)
      ODBG(ODT_Mapping) << "HstPtrBegin " << HstPtrBegin
                        << " was newly allocated for the current region";
    return WasNewlyAllocated;
  };

  // Even if this isn't a new entry, we still need to do a data-transfer if
  // the pointer was newly allocated on the current target region.
  if (LR.TPR.TargetPointer && !LR.TPR.Flags.IsHostPointer && HasFlagTo &&
      (LR.TPR.Flags.IsNewEntry || HasFlagAlways ||
       WasNewlyAllocatedForCurrentRegion()) &&
      Size != 0) {

    // If we have something like:
    //   #pragma omp target map(to: s.myarr[0:10]) map(to: s.myarr[0:10])
    // then we see two "new" mappings of the struct member s.myarr here --
    // and both have the "IsNewEntry" flag set, so trigger the copy to device
    // below.  But, the shadow pointer is only initialised on the target for
    // the first copy, and the second copy clobbers it.  So, this condition
    // avoids the (second) copy here if we have already set shadow pointer info.
    auto FailOnPtrFound = [HstPtrBegin, Size](ShadowPtrInfoTy &SP) {
      if (SP.HstPtrAddr >= HstPtrBegin &&
          SP.HstPtrAddr < (void *)((char *)HstPtrBegin + Size))
        return OFFLOAD_FAIL;
      return OFFLOAD_SUCCESS;
    };
    if (LR.TPR.getEntry()->foreachShadowPointerInfo(FailOnPtrFound) ==
        OFFLOAD_FAIL) {
      ODBG(ODT_Mapping) << "Multiple new mappings of " << Size
                        << " bytes detected (hst:" << HstPtrBegin
                        << ") -> (tgt:" << LR.TPR.TargetPointer << ")";
      return std::move(LR.TPR);
    }

    ODBG(ODT_Mapping) << "Moving " << Size << " bytes (hst:" << HstPtrBegin
                      << ") -> (tgt:" << LR.TPR.TargetPointer << ")";

    int Ret = Device.submitData(LR.TPR.TargetPointer, HstPtrBegin, Size,
                                AsyncInfo, LR.TPR.getEntry());
    if (Ret != OFFLOAD_SUCCESS) {
      REPORT() << "Copying data to device failed.";
      // We will also return nullptr if the data movement fails because that
      // pointer points to a corrupted memory region so it doesn't make any
      // sense to continue to use it.
      LR.TPR.TargetPointer = nullptr;
    } else if (LR.TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
               OFFLOAD_SUCCESS)
      return TargetPointerResultTy{};
  } else {
    // If not a host pointer and no present modifier, we need to wait for the
    // event if it exists.
    // Note: Entry might be nullptr because of zero length array section.
    if (LR.TPR.getEntry() && !LR.TPR.Flags.IsHostPointer &&
        !HasPresentModifier) {
      void *Event = LR.TPR.getEntry()->getEvent();
      if (Event) {
        int Ret = Device.waitEvent(Event, AsyncInfo);
        if (Ret != OFFLOAD_SUCCESS) {
          // If it fails to wait for the event, we need to return nullptr in
          // case of any data race.
          REPORT() << "Failed to wait for event " << Event << ".";
          return TargetPointerResultTy{};
        }
      }
    }
  }

  return std::move(LR.TPR);
}

TargetPointerResultTy MappingInfoTy::getTgtPtrBegin(
    void *HstPtrBegin, int64_t Size, bool UpdateRefCount, bool UseHoldRefCount,
    bool MustContain, bool ForceDelete, bool FromDataEnd) {
  HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();

  LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);

  LR.TPR.Flags.IsPresent = true;

  if (LR.Flags.IsContained ||
      (!MustContain && (LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter))) {
    LR.TPR.Flags.IsLast =
        LR.TPR.getEntry()->decShouldRemove(UseHoldRefCount, ForceDelete);

    if (ForceDelete) {
      LR.TPR.getEntry()->resetRefCount(UseHoldRefCount);
      assert(LR.TPR.Flags.IsLast ==
                 LR.TPR.getEntry()->decShouldRemove(UseHoldRefCount) &&
             "expected correct IsLast prediction for reset");
    }

    // Increment the number of threads that is using the entry on a
    // targetDataEnd, tracking the number of possible "deleters". A thread may
    // come to own the entry deletion even if it was not the last one querying
    // for it. Thus, we must track every query on targetDataEnds to ensure only
    // the last thread that holds a reference to an entry actually deletes it.
    if (FromDataEnd)
      LR.TPR.getEntry()->incDataEndThreadCount();

    const char *RefCountAction;
    if (!UpdateRefCount) {
      RefCountAction = " (update suppressed)";
    } else if (LR.TPR.Flags.IsLast) {
      LR.TPR.getEntry()->decRefCount(UseHoldRefCount);
      assert(LR.TPR.getEntry()->getTotalRefCount() == 0 &&
             "Expected zero reference count when deletion is scheduled");
      if (ForceDelete)
        RefCountAction = " (reset, delayed deletion)";
      else
        RefCountAction = " (decremented, delayed deletion)";
    } else {
      LR.TPR.getEntry()->decRefCount(UseHoldRefCount);
      RefCountAction = " (decremented)";
    }
    const char *DynRefCountAction = UseHoldRefCount ? "" : RefCountAction;
    const char *HoldRefCountAction = UseHoldRefCount ? RefCountAction : "";
    uintptr_t TP = LR.TPR.getEntry()->TgtPtrBegin +
                   ((uintptr_t)HstPtrBegin - LR.TPR.getEntry()->HstPtrBegin);
    INFO(OMP_INFOTYPE_MAPPING_EXISTS, Device.DeviceID,
         "Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
         "Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s\n",
         DPxPTR(HstPtrBegin), DPxPTR(TP), Size,
         LR.TPR.getEntry()->dynRefCountToStr().c_str(), DynRefCountAction,
         LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction);
    LR.TPR.TargetPointer = (void *)TP;
  } else if (PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY ||
             PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY) {
    // If the value isn't found in the mapping and unified shared memory
    // is on then it means we have stumbled upon a value which we need to
    // use directly from the host.
    ODBG(ODT_Mapping) << "Get HstPtrBegin " << HstPtrBegin << " Size=" << Size
                      << " for unified shared memory";
    LR.TPR.Flags.IsPresent = false;
    LR.TPR.Flags.IsHostPointer = true;
    LR.TPR.TargetPointer = HstPtrBegin;
  } else {
    // OpenMP Specification v5.2: if a matching list item is not found, the
    // pointer retains its original value as per firstprivate semantics.
    LR.TPR.Flags.IsPresent = false;
    LR.TPR.Flags.IsHostPointer = false;
    LR.TPR.TargetPointer = HstPtrBegin;
  }

  return std::move(LR.TPR);
}

// Return the target pointer begin (where the data will be moved).
void *MappingInfoTy::getTgtPtrBegin(HDTTMapAccessorTy &HDTTMap,
                                    void *HstPtrBegin, int64_t Size) {
  uintptr_t HP = (uintptr_t)HstPtrBegin;
  LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
  if (LR.Flags.IsContained || LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) {
    uintptr_t TP =
        LR.TPR.getEntry()->TgtPtrBegin + (HP - LR.TPR.getEntry()->HstPtrBegin);
    return (void *)TP;
  }

  return NULL;
}

int MappingInfoTy::eraseMapEntry(HDTTMapAccessorTy &HDTTMap,
                                 HostDataToTargetTy *Entry, int64_t Size) {
  assert(Entry && "Trying to delete a null entry from the HDTT map.");
  assert(Entry->getTotalRefCount() == 0 &&
         Entry->getDataEndThreadCount() == 0 &&
         "Trying to delete entry that is in use or owned by another thread.");

  INFO(OMP_INFOTYPE_MAPPING_CHANGED, Device.DeviceID,
       "Removing map entry with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
       ", Size=%" PRId64 ", Name=%s\n",
       DPxPTR(Entry->HstPtrBegin), DPxPTR(Entry->TgtPtrBegin), Size,
       (Entry->HstPtrName) ? getNameFromMapping(Entry->HstPtrName).c_str()
                           : "unknown");

  if (HDTTMap->erase(Entry) == 0) {
    REPORT() << "Trying to remove a non-existent map entry";
    return OFFLOAD_FAIL;
  }

  return OFFLOAD_SUCCESS;
}

int MappingInfoTy::deallocTgtPtrAndEntry(HostDataToTargetTy *Entry,
                                         int64_t Size) {
  assert(Entry && "Trying to deallocate a null entry.");

  ODBG(ODT_Mapping) << "Deleting tgt data "
                    << reinterpret_cast<void *>(Entry->TgtPtrBegin)
                    << " of size " << Size << " by freeing allocation "
                    << "starting at "
                    << reinterpret_cast<void *>(Entry->TgtAllocBegin);

  void *Event = Entry->getEvent();
  if (Event && Device.destroyEvent(Event) != OFFLOAD_SUCCESS) {
    REPORT() << "Failed to destroy event " << Event;
    return OFFLOAD_FAIL;
  }

  int Ret = Device.deleteData((void *)Entry->TgtAllocBegin);

  // Notify the plugin about the unmapped memory.
  Ret |= Device.notifyDataUnmapped((void *)Entry->HstPtrBegin);

  delete Entry;

  return Ret;
}

static void printCopyInfoImpl(int DeviceId, bool H2D, void *SrcPtrBegin,
                              void *DstPtrBegin, int64_t Size,
                              HostDataToTargetTy *HT) {

  INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceId,
       "Copying data from %s to %s, %sPtr=" DPxMOD ", %sPtr=" DPxMOD
       ", Size=%" PRId64 ", Name=%s\n",
       H2D ? "host" : "device", H2D ? "device" : "host", H2D ? "Hst" : "Tgt",
       DPxPTR(H2D ? SrcPtrBegin : DstPtrBegin), H2D ? "Tgt" : "Hst",
       DPxPTR(H2D ? DstPtrBegin : SrcPtrBegin), Size,
       (HT && HT->HstPtrName) ? getNameFromMapping(HT->HstPtrName).c_str()
                              : "unknown");
}

void MappingInfoTy::printCopyInfo(
    void *TgtPtrBegin, void *HstPtrBegin, int64_t Size, bool H2D,
    HostDataToTargetTy *Entry, MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr) {
  auto HDTTMap =
      HostDataToTargetMap.getExclusiveAccessor(!!Entry || !!HDTTMapPtr);
  LookupResult LR;
  if (!Entry) {
    LR = lookupMapping(HDTTMapPtr ? *HDTTMapPtr : HDTTMap, HstPtrBegin, Size);
    Entry = LR.TPR.getEntry();
  }
  printCopyInfoImpl(Device.DeviceID, H2D, HstPtrBegin, TgtPtrBegin, Size,
                    Entry);
}