[OpenMP][NFCI] Cleanup APIs and improve object encapsulation

[jdoerfert]

Summary

As a preparation patch this one cleans up the API for getTargetPointer
and getTgtPtrBegin. It also encapsulates TargetPointerResultTy properly.
The latter will become more important as we lock the entries
consistently. The former helps us to avoid inconsistencies, e.g., we
returned IsHostPointer in the Flags of the TargetPointerResultTy but
also via the boolean reference. This patch eliminates both reference
arguments.

[ye-luo]

openmp/libomptarget/include/device.h

285–288

three 0?

openmp/libomptarget/src/device.cpp

241

Can we avoid all the setXXX functions?
I feel better keep the local IsNew and then return by calling TPR::TPR(Entry, TargetPointer, IsNew...)

[jdoerfert]

openmp/libomptarget/src/device.cpp

241

But then I need to do manual locking again and I would like us to avoid that.

[ye-luo]

openmp/libomptarget/src/device.cpp

241

Do you have a patch with locking added? Try to understand better your intention.

[jdoerfert]

openmp/libomptarget/src/device.cpp

241

I do not. Every setEntry will lock the entry and unlock the lat entry, if any. Every time TPR goes out of scope, e.g., when we give up and return an dummy object, you unlock the entry. If you do manual locking the caller needs to also unlock when TPR is destroyed, various places all across the code base.

[ye-luo]

openmp/libomptarget/src/device.cpp

241

Assuming the lock is in the entry, it makes sense.

[jdoerfert]

openmp/libomptarget/src/device.cpp

241

Yes. Entry is the one with the lock. TPR is bacically a exclusive accessor with some extra information. @tianshilei1992 asked for more documentation, I will add that before I merge.

[jdoerfert]

openmp/libomptarget/src/device.cpp

391

This belongs into the next patch of this series.

[jdoerfert]

Address comments

[jdoerfert]

Fix rebase errors, address comments

[jdoerfert]

dropped in favor of https://reviews.llvm.org/D123446

openmp/libomptarget/include/device.h

struct LookupResult {
   /// The corresponding map table entry which is stable.
   HostDataToTargetTy *Entry = nullptr;
 
   LookupResult() : Flags({0, 0, 0}), Entry() {}
 };
 
 /// This struct will be returned by \p DeviceTy::getTargetPointer which provides
 /// more data than just a target pointer.
-struct TargetPointerResultTy {
-  struct {
-    /// If the map table entry is just created
+class TargetPointerResultTy {
+  struct FlagTy {
+    /// Flag indicating that the map table entry was just created.
     unsigned IsNewEntry : 1;
-    /// If the pointer is actually a host pointer (when unified memory enabled)
+    /// Flag indicating that the pointer is actually a host pointer, used when
+    /// unified memory enabled.
     unsigned IsHostPointer : 1;
-  } Flags = {0, 0};
+    /// Flag indicating that this was the last user of the entry and the ref
+    /// count is now 0.
+    unsigned IsLast : 1;
+  } Flags = {0, 0, 0};
 
   /// The corresponding map table entry which is stable.
   HostDataToTargetTy *Entry = nullptr;
 
   /// The corresponding target pointer
   void *TargetPointer = nullptr;
+
+public:
+  TargetPointerResultTy() {}
+
+  TargetPointerResultTy(FlagTy Flags, HostDataToTargetTy *Entry,
+                        void *TargetPointer)
+      : Flags(Flags), Entry(Entry), TargetPointer(TargetPointer) {
+  }
+
+  bool isHostPtr() const { return Flags.IsHostPointer; }
+  void setIsHostPtr(bool IHP) { Flags.IsHostPointer = IHP; }
+
+  bool isNew() const { return Flags.IsNewEntry; }
+  void setIsNew(bool IN) { Flags.IsNewEntry = IN; }
+
+  bool isLast() const { return Flags.IsLast; }
+  void setIsLast(bool IL) { Flags.IsLast = IL; }
+
+  void *getTargetPointer() const { return TargetPointer; }
+  void setTargetPointer(void *TP) { TargetPointer = TP; }
+
+  HostDataToTargetTy *getEntry() const { return Entry; }
+  void setEntry(HostDataToTargetTy *HDTTT) { Entry = HDTTT; }
 };
 
 /// Map for shadow pointers
 struct ShadowPtrValTy {
   void *HstPtrVal;
   void *TgtPtrAddr;
   void *TgtPtrVal;
 };
struct DeviceTy {
   /// \p Size will also be transferred to the device. If the mapping doesn't
   /// exist, and if unified shared memory is not enabled, a new mapping will be
   /// created and the data will also be transferred accordingly. nullptr will be
   /// returned because of any of following reasons:
   /// - Data allocation failed;
   /// - The user tried to do an illegal mapping;
   /// - Data transfer issue fails.
   TargetPointerResultTy
-  getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
-                   map_var_info_t HstPtrName, bool HasFlagTo,
-                   bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
-                   bool HasCloseModifier, bool HasPresentModifier,
-                   bool HasHoldModifier, AsyncInfoTy &AsyncInfo);
+  getTargetPointer(HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin,
+                   void *HstPtrBase, int64_t Size, map_var_info_t HstPtrName,
+                   bool HasFlagTo, bool HasFlagAlways, bool IsImplicit,
+                   bool UpdateRefCount, bool HasCloseModifier,
+                   bool HasPresentModifier, bool HasHoldModifier,
+                   AsyncInfoTy &AsyncInfo);
 
   /// Return the target pointer for \p HstPtrBegin in \p HDTTMap. The accessor
   /// ensures exclusive access to the HDTT map.
   void *getTgtPtrBegin(HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin,
                        int64_t Size);
 
   TargetPointerResultTy getTgtPtrBegin(void *HstPtrBegin, int64_t Size,
-                                       bool &IsLast, bool UpdateRefCount,
-                                       bool UseHoldRefCount, bool &IsHostPtr,
+                                       bool UpdateRefCount,
+                                       bool UseHoldRefCount,
                                        bool MustContain = false,
                                        bool ForceDelete = false);
 
   /// Deallocate \p LR and remove the entry. Assume the total reference count is
   /// zero and the calling thread is the deleting thread for \p LR. \p HDTTMap
   /// ensure the caller holds exclusive access and can modify the map. Return \c
   /// OFFLOAD_SUCCESS if the map entry existed, and return \c OFFLOAD_FAIL if
   /// not. It is the caller's responsibility to skip calling this function if

openmp/libomptarget/src/api.cpp

EXTERN int omp_target_is_present(const void *ptr, int device_num) {
   PM->RTLsMtx.unlock();
   if (DevicesSize <= (size_t)device_num) {
     DP("Call to omp_target_is_present with invalid device ID, returning "
        "false\n");
     return false;
   }
 
   DeviceTy &Device = *PM->Devices[device_num];
-  bool IsLast; // not used
-  bool IsHostPtr;
-  TargetPointerResultTy TPR =
-      Device.getTgtPtrBegin(const_cast<void *>(ptr), 0, IsLast,
+  TargetPointerResultTy TPR = Device.getTgtPtrBegin(const_cast<void *>(ptr), 0,
                                                     /*UpdateRefCount=*/false,
-                            /*UseHoldRefCount=*/false, IsHostPtr);
-  int rc = (TPR.TargetPointer != NULL);
+                                                    /*UseHoldRefCount=*/false);
+  int rc = (TPR.getTargetPointer() != NULL);
   // Under unified memory the host pointer can be returned by the
   // getTgtPtrBegin() function which means that there is no device
   // corresponding point for ptr. This function should return false
   // in that situation.
   if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)
-    rc = !IsHostPtr;
+    rc = !TPR.isHostPtr();
   DP("Call to omp_target_is_present returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_target_memcpy(void *dst, const void *src, size_t length,
                              size_t dst_offset, size_t src_offset,
                              int dst_device, int src_device) {
   TIMESCOPE();

openmp/libomptarget/src/device.cpp

if (lr.Flags.ExtendsAfter) {
          "region\n");
     }
   }
 
   return lr;
 }
 
 TargetPointerResultTy DeviceTy::getTargetPointer(
-    void *HstPtrBegin, void *HstPtrBase, int64_t Size,
-    map_var_info_t HstPtrName, bool HasFlagTo, bool HasFlagAlways,
+    HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin, void *HstPtrBase,
+    int64_t Size, map_var_info_t HstPtrName, bool HasFlagTo, bool HasFlagAlways,
     bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier,
     bool HasPresentModifier, bool HasHoldModifier, AsyncInfoTy &AsyncInfo) {
-  HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();
-
-  void *TargetPointer = nullptr;
-  bool IsHostPtr = false;
-  bool IsNew = false;
+  TargetPointerResultTy TPR;
 
   LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
-  auto *Entry = LR.Entry;
+
+  TPR.setEntry(LR.Entry);
+
+  // Not that the entry can be null.
+  HostDataToTargetTy *Entry = TPR.getEntry();
 
   // 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)) {
     auto &HT = *LR.Entry;
     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.
       HT.incRefCount(HasHoldModifier);
       RefCountAction = " (incremented)";
     } else {
       // It might have been allocated with the parent, but it's still new.
-      IsNew = HT.getTotalRefCount() == 1;
+      TPR.setIsNew(HT.getTotalRefCount() == 1);
       RefCountAction = " (update suppressed)";
     }
     const char *DynRefCountAction = HasHoldModifier ? "" : RefCountAction;
     const char *HoldRefCountAction = HasHoldModifier ? RefCountAction : "";
     uintptr_t Ptr = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     INFO(OMP_INFOTYPE_MAPPING_EXISTS, 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, HT.dynRefCountToStr().c_str(), DynRefCountAction,
          HT.holdRefCountToStr().c_str(), HoldRefCountAction,
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
-    TargetPointer = (void *)Ptr;
+    TPR.setTargetPointer((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(Entry->HstPtrBegin),
             Entry->HstPtrEnd - Entry->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->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
              !HasCloseModifier) {
     // 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) {
       DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
          "memory\n",
          DPxPTR((uintptr_t)HstPtrBegin), Size);
-      IsHostPtr = true;
-      TargetPointer = HstPtrBegin;
+      TPR.setIsHostPtr(true);
+      TPR.setTargetPointer(HstPtrBegin);
     }
   } else if (HasPresentModifier) {
     DP("Mapping required by 'present' map type modifier does not exist for "
        "HstPtrBegin=" DPxMOD ", Size=%" PRId64 "\n",
        DPxPTR(HstPtrBegin), 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.
-    IsNew = true;
+    TPR.setIsNew(true);
     uintptr_t Ptr = (uintptr_t)allocData(Size, HstPtrBegin);
     Entry = HDTTMap
                 ->emplace(new HostDataToTargetTy(
                     (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
                     (uintptr_t)HstPtrBegin + Size, Ptr, HasHoldModifier,
                     HstPtrName))
                 .first->HDTT;
+    TPR.setEntry(Entry);
+
     INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
          "Creating new map entry with HstPtrBase= " DPxMOD
          ", HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, "
          "DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
          DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size,
          Entry->dynRefCountToStr().c_str(), Entry->holdRefCountToStr().c_str(),
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
-    TargetPointer = (void *)Ptr;
+    TPR.setTargetPointer((void *)Ptr);
   }
 
   // If the target pointer is valid, and we need to transfer data, issue the
   // data transfer.
-  if (TargetPointer && !IsHostPtr && HasFlagTo && (IsNew || HasFlagAlways)) {
-    // Lock the entry before releasing the mapping table lock such that another
-    // thread that could issue data movement will get the right result.
-    std::lock_guard<decltype(*Entry)> LG(*Entry);
-    // Release the mapping table lock right after the entry is locked.
-    HDTTMap.destroy();
-
+  if (TPR.getTargetPointer() && !TPR.isHostPtr() && HasFlagTo &&
+      (TPR.isNew() || HasFlagAlways)) {
     DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n", Size,
-       DPxPTR(HstPtrBegin), DPxPTR(TargetPointer));
+       DPxPTR(HstPtrBegin), DPxPTR(TPR.getTargetPointer()));
 
-    int Ret = submitData(TargetPointer, HstPtrBegin, Size, AsyncInfo);
+    int Ret = submitData(TPR.getTargetPointer(), HstPtrBegin, Size, AsyncInfo);
     if (Ret != OFFLOAD_SUCCESS) {
       REPORT("Copying data to device failed.\n");
       // 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.
-      TargetPointer = nullptr;
-    } else if (Entry->addEventIfNecessary(*this, AsyncInfo) != OFFLOAD_SUCCESS)
+      TPR.setTargetPointer(nullptr);
+    } else if (Entry->addEventIfNecessary(*this, AsyncInfo) !=
+               OFFLOAD_SUCCESS) {
       return {{false /* IsNewEntry */, false /* IsHostPointer */},
               nullptr /* Entry */,
               nullptr /* TargetPointer */};
+    }
   } else {
-    // Release the mapping table lock directly.
-    HDTTMap.destroy();
     // 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 (Entry && !IsHostPtr && !HasPresentModifier) {
-      std::lock_guard<decltype(*Entry)> LG(*Entry);
+    if (Entry && !TPR.isHostPtr() && !HasPresentModifier) {
       void *Event = Entry->getEvent();
       if (Event) {
         int Ret = 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 " DPxMOD ".\n", DPxPTR(Event));
           return {{false /* IsNewEntry */, false /* IsHostPointer */},
                   nullptr /* Entry */,
                   nullptr /* TargetPointer */};
         }
       }
     }
   }
 
-  return {{IsNew, IsHostPtr}, Entry, TargetPointer};
+  return TPR;
 }
 
 // Used by targetDataBegin, targetDataEnd, targetDataUpdate and target.
 // Return the target pointer begin (where the data will be moved).
 // Decrement the reference counter if called from targetDataEnd.
-TargetPointerResultTy
-DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-                         bool UpdateRefCount, bool UseHoldRefCount,
-                         bool &IsHostPtr, bool MustContain, bool ForceDelete) {
-  HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();
+TargetPointerResultTy DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size,
+                                               bool UpdateRefCount,
+                                               bool UseHoldRefCount,
+                                               bool MustContain,
+                                               bool ForceDelete) {
 
-  void *TargetPointer = NULL;
-  bool IsNew = false;
-  IsHostPtr = false;
-  IsLast = false;
-  LookupResult lr = lookupMapping(HDTTMap, HstPtrBegin, Size);
+  TargetPointerResultTy TPR;
+  LookupResult LR;
+  {
+    HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();
+    LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
+    TPR.setEntry(LR.Entry);
+  }
 
-  if (lr.Flags.IsContained ||
-      (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
-    auto &HT = *lr.Entry;
-    IsLast = HT.decShouldRemove(UseHoldRefCount, ForceDelete);
+  if (LR.Flags.IsContained ||
+      (!MustContain && (LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter))) {
+    auto &HT = *LR.Entry;
+    TPR.setIsLast(HT.decShouldRemove(UseHoldRefCount, ForceDelete));
 
     if (ForceDelete) {
       HT.resetRefCount(UseHoldRefCount);
-      assert(IsLast == HT.decShouldRemove(UseHoldRefCount) &&
+      assert(TPR.isLast() == HT.decShouldRemove(UseHoldRefCount) &&
              "expected correct IsLast prediction for reset");
     }
 
     const char *RefCountAction;
     if (!UpdateRefCount) {
       RefCountAction = " (update suppressed)";
-    } else if (IsLast) {
+    } else if (TPR.isLast()) {
       // Mark the entry as to be deleted by this thread. Another thread might
       // reuse the entry and take "ownership" for the deletion while this thread
       // is waiting for data transfers. That is fine and the current thread will
       // simply skip the deletion step then.
       HT.setDeleteThreadId();
       HT.decRefCount(UseHoldRefCount);
       assert(HT.getTotalRefCount() == 0 &&
              "Expected zero reference count when deletion is scheduled");
       if (ForceDelete)
         RefCountAction = " (reset, delayed deletion)";
       else
         RefCountAction = " (decremented, delayed deletion)";
     } else {
       HT.decRefCount(UseHoldRefCount);
       RefCountAction = " (decremented)";
     }
     const char *DynRefCountAction = UseHoldRefCount ? "" : RefCountAction;
     const char *HoldRefCountAction = UseHoldRefCount ? RefCountAction : "";
     uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
          "Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
          "Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s\n",
          DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.dynRefCountToStr().c_str(),
          DynRefCountAction, HT.holdRefCountToStr().c_str(), HoldRefCountAction);
-    TargetPointer = (void *)tp;
+    TPR.setTargetPointer((void *)tp);
   } else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
     // 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.
     DP("Get HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
        "memory\n",
        DPxPTR((uintptr_t)HstPtrBegin), Size);
-    IsHostPtr = true;
-    TargetPointer = HstPtrBegin;
+    TPR.setIsHostPtr(true);
+    TPR.setTargetPointer(HstPtrBegin);
   }
 
-  return {{IsNew, IsHostPtr}, lr.Entry, TargetPointer};
+  return TPR;
 }
 
 // Return the target pointer begin (where the data will be moved).
 void *DeviceTy::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) {

openmp/libomptarget/src/omptarget.cpp

if (getParentIndex(arg_types[i]) < 0 && next_i < arg_num &&
            "\n",
            padding, DPxPTR(HstPtrBegin));
         HstPtrBegin = (char *)HstPtrBegin - padding;
         data_size += padding;
       }
     }
 
     // Address of pointer on the host and device, respectively.
-    void *Pointer_HstPtrBegin, *PointerTgtPtrBegin;
+    void **PointerHstPtrBegin, **PointerTgtPtrBegin;
     TargetPointerResultTy Pointer_TPR;
-    bool IsHostPtr = false;
     bool IsImplicit = arg_types[i] & OMP_TGT_MAPTYPE_IMPLICIT;
     // Force the creation of a device side copy of the data when:
     // a close map modifier was associated with a map that contained a to.
     bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE;
     bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT;
     bool HasHoldModifier = arg_types[i] & OMP_TGT_MAPTYPE_OMPX_HOLD;
     // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we
     // have reached this point via __tgt_target_data_begin and not __tgt_target
     // then no argument is marked as TARGET_PARAM ("omp target data map" is not
     // associated with a target region, so there are no target parameters). This
     // may be considered a hack, we could revise the scheme in the future.
     bool UpdateRef =
         !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && i == 0);
+
+    DeviceTy::HDTTMapAccessorTy HDTTMap =
+        Device.HostDataToTargetMap.getExclusiveAccessor();
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       DP("Has a pointer entry: \n");
       // Base is address of pointer.
       //
       // Usually, the pointer is already allocated by this time.  For example:
       //
       //   #pragma omp target map(s.p[0:N])
       //
       // The map entry for s comes first, and the PTR_AND_OBJ entry comes
       // afterward, so the pointer is already allocated by the time the
       // PTR_AND_OBJ entry is handled below, and PointerTgtPtrBegin is thus
       // non-null.  However, "declare target link" can produce a PTR_AND_OBJ
       // entry for a global that might not already be allocated by the time the
       // PTR_AND_OBJ entry is handled below, and so the allocation might fail
       // when HasPresentModifier.
       Pointer_TPR = Device.getTargetPointer(
-          HstPtrBase, HstPtrBase, sizeof(void *), /*HstPtrName=*/nullptr,
+          HDTTMap, HstPtrBase, HstPtrBase, sizeof(void *),
+          /*HstPtrName=*/nullptr,
           /*HasFlagTo=*/false, /*HasFlagAlways=*/false, IsImplicit, UpdateRef,
           HasCloseModifier, HasPresentModifier, HasHoldModifier, AsyncInfo);
-      PointerTgtPtrBegin = Pointer_TPR.TargetPointer;
-      IsHostPtr = Pointer_TPR.Flags.IsHostPointer;
+      PointerTgtPtrBegin = (void **)Pointer_TPR.getTargetPointer();
       if (!PointerTgtPtrBegin) {
         REPORT("Call to getTargetPointer returned null pointer (%s).\n",
                HasPresentModifier ? "'present' map type modifier"
                                   : "device failure or illegal mapping");
         return OFFLOAD_FAIL;
       }
       DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new"
          "\n",
          sizeof(void *), DPxPTR(PointerTgtPtrBegin),
-         (Pointer_TPR.Flags.IsNewEntry ? "" : " not"));
-      Pointer_HstPtrBegin = HstPtrBase;
+         (Pointer_TPR.isNew() ? "" : " not"));
+      PointerHstPtrBegin = (void **)HstPtrBase;
       // modify current entry.
-      HstPtrBase = *(void **)HstPtrBase;
+      HstPtrBase = *PointerHstPtrBegin;
       // No need to update pointee ref count for the first element of the
       // subelement that comes from mapper.
       UpdateRef =
           (!FromMapper || i != 0); // subsequently update ref count of pointee
     }
 
     const bool HasFlagTo = arg_types[i] & OMP_TGT_MAPTYPE_TO;
     const bool HasFlagAlways = arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS;
     auto TPR = Device.getTargetPointer(
-        HstPtrBegin, HstPtrBase, data_size, HstPtrName, HasFlagTo,
+        HDTTMap, HstPtrBegin, HstPtrBase, data_size, HstPtrName, HasFlagTo,
         HasFlagAlways, IsImplicit, UpdateRef, HasCloseModifier,
         HasPresentModifier, HasHoldModifier, AsyncInfo);
-    void *TgtPtrBegin = TPR.TargetPointer;
-    IsHostPtr = TPR.Flags.IsHostPointer;
+
+    void *TgtPtrBegin = TPR.getTargetPointer();
     // If data_size==0, then the argument could be a zero-length pointer to
     // NULL, so getOrAlloc() returning NULL is not an error.
     if (!TgtPtrBegin && (data_size || HasPresentModifier)) {
       REPORT("Call to getTargetPointer returned null pointer (%s).\n",
              HasPresentModifier ? "'present' map type modifier"
                                 : "device failure or illegal mapping");
       return OFFLOAD_FAIL;
     }
     DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
        " - is%s new\n",
-       data_size, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));
+       data_size, DPxPTR(TgtPtrBegin), (TPR.isNew() ? "" : " not"));
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
       uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
       void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
       DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
       args_base[i] = TgtPtrBase;
     }
 
-    if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
+    if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !TPR.isHostPtr()) {
       // Check whether we need to update the pointer on the device
       bool UpdateDevPtr = false;
 
       uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
       void *ExpectedTgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta);
 
       Device.ShadowMtx.lock();
-      auto Entry = Device.ShadowPtrMap.find(Pointer_HstPtrBegin);
+      auto Entry = Device.ShadowPtrMap.find(PointerHstPtrBegin);
       // If this pointer is not in the map we need to insert it. If the map
       // contains a stale entry, we need to update it (e.g. if the pointee was
       // deallocated and later on is reallocated at another device address). The
       // latter scenario is the subject of LIT test env/base_ptr_ref_count.c. An
       // entry is removed from ShadowPtrMap only when the PTR of a PTR_AND_OBJ
       // pair is deallocated, not when the OBJ is deallocated. In
       // env/base_ptr_ref_count.c the PTR is a global "declare target" pointer,
       // so it stays in the map for the lifetime of the application. When the
       // OBJ is deallocated and later on allocated again (at a different device
-      // address), ShadowPtrMap still contains an entry for Pointer_HstPtrBegin
+      // address), ShadowPtrMap still contains an entry for PointerHstPtrBegin
       // which is stale, pointing to the old ExpectedTgtPtrBase of the OBJ.
       if (Entry == Device.ShadowPtrMap.end() ||
           Entry->second.TgtPtrVal != ExpectedTgtPtrBase) {
         // create or update shadow pointers for this entry
-        Device.ShadowPtrMap[Pointer_HstPtrBegin] = {
+        Device.ShadowPtrMap[PointerHstPtrBegin] = {
             HstPtrBase, PointerTgtPtrBegin, ExpectedTgtPtrBase};
-        Pointer_TPR.Entry->setMayContainAttachedPointers();
+        Pointer_TPR.getEntry()->setMayContainAttachedPointers();
         UpdateDevPtr = true;
       }
 
       if (UpdateDevPtr) {
-        std::lock_guard<decltype(*Pointer_TPR.Entry)> LG(*Pointer_TPR.Entry);
+        std::lock_guard<decltype(*Pointer_TPR.getEntry())> LG(
+            *Pointer_TPR.getEntry());
         Device.ShadowMtx.unlock();
 
         DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n",
            DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin));
 
         void *&TgtPtrBase = AsyncInfo.getVoidPtrLocation();
         TgtPtrBase = ExpectedTgtPtrBase;
 
         int Ret = Device.submitData(PointerTgtPtrBegin, &TgtPtrBase,
                                     sizeof(void *), AsyncInfo);
         if (Ret != OFFLOAD_SUCCESS) {
           REPORT("Copying data to device failed.\n");
           return OFFLOAD_FAIL;
         }
-        if (Pointer_TPR.Entry->addEventIfNecessary(Device, AsyncInfo) !=
+        if (Pointer_TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
             OFFLOAD_SUCCESS)
           return OFFLOAD_FAIL;
       } else
         Device.ShadowMtx.unlock();
     }
   }
 
   return OFFLOAD_SUCCESS;
static void applyToShadowMapEntries(DeviceTy &Device, CBTy CB, void *Begin,
                                     const TargetPointerResultTy &TPR) {
   // If we have an object that is too small to hold a pointer subobject, no need
   // to do any checking.
   if (Size < sizeof(void *))
     return;
 
   // If the map entry for the object was never marked as containing attached
   // pointers, no need to do any checking.
-  if (!TPR.Entry || !TPR.Entry->getMayContainAttachedPointers())
+  if (!TPR.getEntry() || !TPR.getEntry()->getMayContainAttachedPointers())
     return;
 
   uintptr_t LB = (uintptr_t)Begin;
   uintptr_t UB = LB + Size;
   // Now we are looking into the shadow map so we need to lock it.
   std::lock_guard<decltype(Device.ShadowMtx)> LG(Device.ShadowMtx);
   for (ShadowPtrListTy::iterator Itr = Device.ShadowPtrMap.begin();
        Itr != Device.ShadowPtrMap.end();) {
if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
         DP("Using a Padding of %" PRId64 " bytes for begin address " DPxMOD
            "\n",
            Padding, DPxPTR(HstPtrBegin));
         HstPtrBegin = (char *)HstPtrBegin - Padding;
         DataSize += Padding;
       }
     }
 
-    bool IsLast, IsHostPtr;
     bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;
     bool UpdateRef = (!(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
                       (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) &&
                      !(FromMapper && I == 0);
     bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
     bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
     bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
-    TargetPointerResultTy TPR = Device.getTgtPtrBegin(
-        HstPtrBegin, DataSize, IsLast, UpdateRef, HasHoldModifier, IsHostPtr,
+    TargetPointerResultTy TPR =
+        Device.getTgtPtrBegin(HstPtrBegin, DataSize, UpdateRef, HasHoldModifier,
                               !IsImplicit, ForceDelete);
-    void *TgtPtrBegin = TPR.TargetPointer;
+    void *TgtPtrBegin = TPR.getTargetPointer();
     if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
       DP("Mapping does not exist (%s)\n",
          (HasPresentModifier ? "'present' map type modifier" : "ignored"));
       if (HasPresentModifier) {
         // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:
         // "If a map clause appears on a target, target data, target enter data
         // or target exit data construct with a present map-type-modifier then
         // on entry to the region if the corresponding list item does not appear
if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
         MESSAGE("device mapping required by 'present' map type modifier does "
                 "not exist for host address " DPxMOD " (%" PRId64 " bytes)",
                 DPxPTR(HstPtrBegin), DataSize);
         return OFFLOAD_FAIL;
       }
     } else {
       DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
          " - is%s last\n",
-         DataSize, DPxPTR(TgtPtrBegin), (IsLast ? "" : " not"));
+         DataSize, DPxPTR(TgtPtrBegin), (TPR.isLast() ? "" : " not"));
     }
 
     // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
     // "If the map clause appears on a target, target data, or target exit data
     // construct and a corresponding list item of the original list item is not
     // present in the device data environment on exit from the region then the
     // list item is ignored."
     if (!TgtPtrBegin)
       continue;
 
-    bool DelEntry = IsLast;
+    bool DelEntry = TPR.isLast();
 
     // If the last element from the mapper (for end transfer args comes in
     // reverse order), do not remove the partial entry, the parent struct still
     // exists.
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
         !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
       DelEntry = false; // protect parent struct from being deallocated
     }
 
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
       // Move data back to the host
       if (ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) {
         bool Always = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
-        if ((Always || IsLast) && !IsHostPtr) {
+        if ((Always || TPR.isLast()) && !TPR.isHostPtr()) {
           DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
              DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
 
-          std::lock_guard<decltype(*TPR.Entry)> LG(*TPR.Entry);
           // Wait for any previous transfer if an event is present.
-          if (void *Event = TPR.Entry->getEvent()) {
+          if (void *Event = TPR.getEntry()->getEvent()) {
             if (Device.waitEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {
               REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));
               return OFFLOAD_FAIL;
             }
           }
 
           Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize,
                                     AsyncInfo);
           if (Ret != OFFLOAD_SUCCESS) {
             REPORT("Copying data from device failed.\n");
             return OFFLOAD_FAIL;
           }
 
           // As we are expecting to delete the entry the d2h copy might race
           // with another one that also tries to delete the entry. This happens
           // as the entry can be reused and the reuse might happen after the
           // copy-back was issued but before it completed. Since the reuse might
           // also copy-back a value we would race.
-          if (IsLast) {
-            if (TPR.Entry->addEventIfNecessary(Device, AsyncInfo) !=
+          if (TPR.isLast()) {
+            if (TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
                 OFFLOAD_SUCCESS)
               return OFFLOAD_FAIL;
           }
         }
       }
       if (DelEntry && FromMapper && I == 0) {
         DelEntry = false;
         FromMapperBase = HstPtrBegin;
       }
 
       // Add pointer to the buffer for post-synchronize processing.
       PostProcessingPtrs.emplace_back(HstPtrBegin, DataSize, ArgTypes[I],
-                                      DelEntry && !IsHostPtr, TPR);
+                                      DelEntry && !TPR.isHostPtr(), TPR);
     }
   }
 
   // TODO: We should not synchronize here but pass the AsyncInfo object to the
   //       allocate/deallocate device APIs.
   //
   // We need to synchronize before deallocating data.
   Ret = AsyncInfo.synchronize();
int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
 
   return Ret;
 }
 
 static int targetDataContiguous(ident_t *loc, DeviceTy &Device, void *ArgsBase,
                                 void *HstPtrBegin, int64_t ArgSize,
                                 int64_t ArgType, AsyncInfoTy &AsyncInfo) {
   TIMESCOPE_WITH_IDENT(loc);
-  bool IsLast, IsHostPtr;
-  TargetPointerResultTy TPR = Device.getTgtPtrBegin(
-      HstPtrBegin, ArgSize, IsLast, /*UpdateRefCount=*/false,
-      /*UseHoldRefCount=*/false, IsHostPtr, /*MustContain=*/true);
-  void *TgtPtrBegin = TPR.TargetPointer;
+  TargetPointerResultTy TPR =
+      Device.getTgtPtrBegin(HstPtrBegin, ArgSize, /*UpdateRefCount=*/false,
+                            /*UseHoldRefCount=*/false, /*MustContain=*/true);
+  void *TgtPtrBegin = TPR.getTargetPointer();
   if (!TgtPtrBegin) {
     DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
     if (ArgType & OMP_TGT_MAPTYPE_PRESENT) {
       MESSAGE("device mapping required by 'present' motion modifier does not "
               "exist for host address " DPxMOD " (%" PRId64 " bytes)",
               DPxPTR(HstPtrBegin), ArgSize);
       return OFFLOAD_FAIL;
     }
     return OFFLOAD_SUCCESS;
   }
 
-  if (IsHostPtr) {
+  if (TPR.isHostPtr()) {
     DP("hst data:" DPxMOD " unified and shared, becomes a noop\n",
        DPxPTR(HstPtrBegin));
     return OFFLOAD_SUCCESS;
   }
 
   if (ArgType & OMP_TGT_MAPTYPE_FROM) {
     DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
        ArgSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
if (!(ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {
         unsigned Idx = getParentIndex(ArgTypes[I]);
         int TgtIdx = TgtArgsPositions[Idx];
         assert(TgtIdx != -1 && "Base address must be translated already.");
         // The parent lambda must be processed already and it must be the last
         // in TgtArgs and TgtOffsets arrays.
         void *HstPtrVal = Args[I];
         void *HstPtrBegin = ArgBases[I];
         void *HstPtrBase = Args[Idx];
-        bool IsLast, IsHostPtr; // IsLast is unused.
         void *TgtPtrBase =
             (void *)((intptr_t)TgtArgs[TgtIdx] + TgtOffsets[TgtIdx]);
         DP("Parent lambda base " DPxMOD "\n", DPxPTR(TgtPtrBase));
         uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
         void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta);
         void *&PointerTgtPtrBegin = AsyncInfo.getVoidPtrLocation();
         TargetPointerResultTy TPR = Device.getTgtPtrBegin(
-            HstPtrVal, ArgSizes[I], IsLast, /*UpdateRefCount=*/false,
-            /*UseHoldRefCount=*/false, IsHostPtr);
-        PointerTgtPtrBegin = TPR.TargetPointer;
+            HstPtrVal, ArgSizes[I], /*UpdateRefCount=*/false,
+            /*UseHoldRefCount=*/false);
+        PointerTgtPtrBegin = TPR.getTargetPointer();
         if (!PointerTgtPtrBegin) {
           DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
              DPxPTR(HstPtrVal));
           continue;
         }
-        if (IsHostPtr) {
+        if (TPR.isHostPtr()) {
           DP("Unified memory is active, no need to map lambda captured"
              "variable (" DPxMOD ")\n",
              DPxPTR(HstPtrVal));
           continue;
         }
         DP("Update lambda reference (" DPxMOD ") -> [" DPxMOD "]\n",
            DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin));
         Ret = Device.submitData(TgtPtrBegin, &PointerTgtPtrBegin,
                                 sizeof(void *), AsyncInfo);
         if (Ret != OFFLOAD_SUCCESS) {
           REPORT("Copying data to device failed.\n");
           return OFFLOAD_FAIL;
         }
       }
       continue;
     }
     void *HstPtrBegin = Args[I];
     void *HstPtrBase = ArgBases[I];
     void *TgtPtrBegin;
     map_var_info_t HstPtrName = (!ArgNames) ? nullptr : ArgNames[I];
     ptrdiff_t TgtBaseOffset;
-    bool IsLast, IsHostPtr; // unused.
     TargetPointerResultTy TPR;
     if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {
       DP("Forwarding first-private value " DPxMOD " to the target construct\n",
          DPxPTR(HstPtrBase));
       TgtPtrBegin = HstPtrBase;
       TgtBaseOffset = 0;
     } else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE) {
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {
       if (Ret != OFFLOAD_SUCCESS) {
         REPORT("Failed to process %sprivate argument " DPxMOD "\n",
                (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin));
         return OFFLOAD_FAIL;
       }
     } else {
       if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)
         HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);
-      TPR = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
+      TPR = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I],
                                   /*UpdateRefCount=*/false,
-                                  /*UseHoldRefCount=*/false, IsHostPtr);
-      TgtPtrBegin = TPR.TargetPointer;
+                                  /*UseHoldRefCount=*/false);
+      TgtPtrBegin = TPR.getTargetPointer();
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
 #ifdef OMPTARGET_DEBUG
       void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);
       DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n",
          DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin));
 #endif
     }
     TgtArgsPositions[I] = TgtArgs.size();