[OpenMP][OpenACC] Implement `ompx_hold` map type modifier extension in runtime (2/2)

[jdenny]

Summary

This patch implements OpenMP runtime support for an original OpenMP
extension we have developed to support OpenACC: the ompx_hold map
type modifier. The previous patch in this series, D106509, implements
Clang support and documents the new functionality in detail.

[jdenny]

Updated for changes in D106509:

• Renamed hold to ompx_hold.
• Added -fopenmp-extensions to tests. Without it, Clang doesn't recognize ompx_hold.

[jdenny]

Rebased.

[jdenny]

Rebased. Ping.

[grokos]

openmp/libomptarget/src/device.cpp

200–203

IsNew is already false at this point, no need to reassign.

343

Currently, we don't decrement the RefCount at this point. Is this a bug in the existing code?

[jdenny]

Thanks for reviewing.

openmp/libomptarget/src/device.cpp

200–203

I'll remove it.

343

Without this patch, the only way I know of that ForceDelete=true but IsLast=false is if the reference count is infinite. Decrementing the reference count would have no effect then (and neither did resetting), so I don't think there's an existing bug.

With this patch, a new possibility is that the hold reference count is non-zero. Because, as a result, IsLast=false, we're not to ready to unmap, so the final decrement of the dynamic reference counter after its reset needs to happen here.

[grokos]

LGTM.

openmp/libomptarget/src/device.cpp

343

Oh, of course, makes sense.

[jdenny]

Thanks!

I'll have to work on landing next week when I can watch the bots.

openmp/libomptarget/include/omptarget.h

enum tgt_map_type {
   // copy by value - not mapped
   OMP_TGT_MAPTYPE_LITERAL         = 0x100,
   // mapping is implicit
   OMP_TGT_MAPTYPE_IMPLICIT        = 0x200,
   // copy data to device
   OMP_TGT_MAPTYPE_CLOSE           = 0x400,
   // runtime error if not already allocated
   OMP_TGT_MAPTYPE_PRESENT         = 0x1000,
+  // use a separate reference counter so that the data cannot be unmapped within
+  // the structured region
+  // This is an OpenMP extension for the sake of OpenACC support.
+  OMP_TGT_MAPTYPE_OMPX_HOLD       = 0x2000,
   // descriptor for non-contiguous target-update
   OMP_TGT_MAPTYPE_NON_CONTIG      = 0x100000000000,
   // member of struct, member given by [16 MSBs] - 1
   OMP_TGT_MAPTYPE_MEMBER_OF       = 0xffff000000000000
 };
 
 enum OpenMPOffloadingDeclareTargetFlags {
   /// Mark the entry as having a 'link' attribute.

openmp/libomptarget/src/api.cpp

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;
   void *TgtPtr = Device.getTgtPtrBegin(const_cast<void *>(ptr), 0, IsLast,
-                                       false, IsHostPtr);
+                                       /*UpdateRefCount=*/false,
+                                       /*UseHoldRefCount=*/false, IsHostPtr);
   int rc = (TgtPtr != 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;
   DP("Call to omp_target_is_present returns %d\n", rc);

openmp/libomptarget/src/device.h

struct HostDataToTargetTy {
   uintptr_t HstPtrBase; // host info.
   uintptr_t HstPtrBegin;
   uintptr_t HstPtrEnd;       // non-inclusive.
   map_var_info_t HstPtrName; // Optional source name of mapped variable.
 
   uintptr_t TgtPtrBegin; // target info.
 
 private:
-  /// use mutable to allow modification via std::set iterator which is const.
-  mutable uint64_t RefCount;
+  /// The dynamic reference count is the standard reference count as of OpenMP
+  /// 4.5.  The hold reference count is an OpenMP extension for the sake of
+  /// OpenACC support.
+  ///
+  /// The 'ompx_hold' map type modifier is permitted only on "omp target" and
+  /// "omp target data", and "delete" is permitted only on "omp target exit
+  /// data" and associated runtime library routines.  As a result, we really
+  /// need to implement "reset" functionality only for the dynamic reference
+  /// counter.  Likewise, only the dynamic reference count can be infinite
+  /// because, for example, omp_target_associate_ptr and "omp declare target
+  /// link" operate only on it.  Nevertheless, it's actually easier to follow
+  /// the code (and requires less assertions for special cases) when we just
+  /// implement these features generally across both reference counters here.
+  /// Thus, it's the users of this class that impose those restrictions.
+  ///
+  /// Use mutable to allow modification via std::set iterator which is const.
+  ///@{
+  mutable uint64_t DynRefCount;
+  mutable uint64_t HoldRefCount;
+  ///@}
   static const uint64_t INFRefCount = ~(uint64_t)0;
+  static std::string refCountToStr(uint64_t RefCount) {
+    return RefCount == INFRefCount ? "INF" : std::to_string(RefCount);
+  }
   /// This mutex will be locked when data movement is issued. For targets that
   /// doesn't support async data movement, this mutex can guarantee that after
   /// it is released, memory region on the target is update to date. For targets
   /// that support async data movement, this can guarantee that data movement
   /// has been issued. This mutex *must* be locked right before releasing the
   /// mapping table lock.
   std::shared_ptr<std::mutex> UpdateMtx;
 
 public:
   HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB,
-                     map_var_info_t Name = nullptr, bool IsINF = false)
+                     bool UseHoldRefCount, map_var_info_t Name = nullptr,
+                     bool IsINF = false)
       : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), HstPtrName(Name),
-        TgtPtrBegin(TB), RefCount(IsINF ? INFRefCount : 1),
+        TgtPtrBegin(TB), DynRefCount(UseHoldRefCount ? 0
+                                     : IsINF         ? INFRefCount
+                                                     : 1),
+        HoldRefCount(!UseHoldRefCount ? 0
+                     : IsINF          ? INFRefCount
+                                      : 1),
         UpdateMtx(std::make_shared<std::mutex>()) {}
 
-  uint64_t getRefCount() const { return RefCount; }
+  /// Get the total reference count.  This is smarter than just getDynRefCount()
+  /// + getHoldRefCount() because it handles the case where at least one is
+  /// infinity and the other is non-zero.
+  uint64_t getTotalRefCount() const {
+    if (DynRefCount == INFRefCount || HoldRefCount == INFRefCount)
+      return INFRefCount;
+    return DynRefCount + HoldRefCount;
+  }
 
-  uint64_t resetRefCount() const {
-    if (RefCount != INFRefCount)
-      RefCount = 1;
+  /// Get the dynamic reference count.
+  uint64_t getDynRefCount() const { return DynRefCount; }
 
-    return RefCount;
-  }
+  /// Get the hold reference count.
+  uint64_t getHoldRefCount() const { return HoldRefCount; }
 
-  uint64_t incRefCount() const {
-    if (RefCount != INFRefCount) {
-      ++RefCount;
-      assert(RefCount < INFRefCount && "refcount overflow");
-    }
+  /// Reset the specified reference count unless it's infinity.  Reset to 1
+  /// (even if currently 0) so it can be followed by a decrement.
+  void resetRefCount(bool UseHoldRefCount) const {
+    uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
+    if (ThisRefCount != INFRefCount)
+      ThisRefCount = 1;
 
-    return RefCount;
+  /// Increment the specified reference count unless it's infinity.
+  void incRefCount(bool UseHoldRefCount) const {
+    uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
+    if (ThisRefCount != INFRefCount) {
+      ++ThisRefCount;
+      assert(ThisRefCount < INFRefCount && "refcount overflow");
     }
-
-  uint64_t decRefCount() const {
-    if (RefCount != INFRefCount) {
-      assert(RefCount > 0 && "refcount underflow");
-      --RefCount;
-    }
 
-    return RefCount;
+  /// Decrement the specified reference count unless it's infinity or zero, and
+  /// return the total reference count.
+  uint64_t decRefCount(bool UseHoldRefCount) const {
+    uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
+    uint64_t OtherRefCount = UseHoldRefCount ? DynRefCount : HoldRefCount;
+    if (ThisRefCount != INFRefCount) {
+      if (ThisRefCount > 0)
+        --ThisRefCount;
+      else
+        assert(OtherRefCount > 0 && "total refcount underflow");
+    }
+    return getTotalRefCount();
   }
 
-  bool isRefCountInf() const { return RefCount == INFRefCount; }
+  /// Is the dynamic (and thus the total) reference count infinite?
+  bool isDynRefCountInf() const { return DynRefCount == INFRefCount; }
 
-  std::string refCountToStr() const {
-    return isRefCountInf() ? "INF" : std::to_string(getRefCount());
-  }
+  /// Convert the dynamic reference count to a debug string.
+  std::string dynRefCountToStr() const { return refCountToStr(DynRefCount); }
 
-  /// Should one decrement of the reference count (after resetting it if
-  /// \c AfterReset) remove this mapping?
-  bool decShouldRemove(bool AfterReset = false) const {
+  /// Convert the hold reference count to a debug string.
+  std::string holdRefCountToStr() const { return refCountToStr(HoldRefCount); }
+
+  /// Should one decrement of the specified reference count (after resetting it
+  /// if \c AfterReset) remove this mapping?
+  bool decShouldRemove(bool UseHoldRefCount, bool AfterReset = false) const {
+    uint64_t ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
+    uint64_t OtherRefCount = UseHoldRefCount ? DynRefCount : HoldRefCount;
+    if (OtherRefCount > 0)
+      return false;
     if (AfterReset)
-      return !isRefCountInf();
-    return getRefCount() == 1;
+      return ThisRefCount != INFRefCount;
+    return ThisRefCount == 1;
   }
 
   void lock() const { UpdateMtx->lock(); }
 
   void unlock() const { UpdateMtx->unlock(); }
 };
 
 typedef uintptr_t HstPtrBeginTy;
struct DeviceTy {
   /// 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, MoveDataStateTy MoveData,
                    bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier,
-                   bool HasPresentModifier, AsyncInfoTy &AsyncInfo);
+                   bool HasPresentModifier, bool HasHoldModifier,
+                   AsyncInfoTy &AsyncInfo);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-                       bool UpdateRefCount, bool &IsHostPtr,
-                       bool MustContain = false, bool ForceDelete = false);
-  int deallocTgtPtr(void *TgtPtrBegin, int64_t Size,
-                    bool HasCloseModifier = false);
+                       bool UpdateRefCount, bool UseHoldRefCount,
+                       bool &IsHostPtr, bool MustContain = false,
+                       bool ForceDelete = false);
+  int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool HasCloseModifier,
+                    bool HasHoldModifier);
   int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
   int disassociatePtr(void *HstPtrBegin);
 
   // calls to RTL
   int32_t initOnce();
   __tgt_target_table *load_binary(void *Img);
 
   // device memory allocation/deallocation routines

openmp/libomptarget/src/device.cpp

int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
 
   // Mapping does not exist, allocate it with refCount=INF
   const HostDataToTargetTy &newEntry =
       *HostDataToTargetMap
            .emplace(
                /*HstPtrBase=*/(uintptr_t)HstPtrBegin,
                /*HstPtrBegin=*/(uintptr_t)HstPtrBegin,
                /*HstPtrEnd=*/(uintptr_t)HstPtrBegin + Size,
-               /*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin, /*Name=*/nullptr,
+               /*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin,
+               /*UseHoldRefCount=*/false, /*Name=*/nullptr,
                /*IsRefCountINF=*/true)
            .first;
   DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD
-     ", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD ", RefCount=%s\n",
+     ", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD ", DynRefCount=%s, "
+     "HoldRefCount=%s\n",
      DPxPTR(newEntry.HstPtrBase), DPxPTR(newEntry.HstPtrBegin),
      DPxPTR(newEntry.HstPtrEnd), DPxPTR(newEntry.TgtPtrBegin),
-     newEntry.refCountToStr().c_str());
+     newEntry.dynRefCountToStr().c_str(), newEntry.holdRefCountToStr().c_str());
   (void)newEntry;
 
   DataMapMtx.unlock();
 
   return OFFLOAD_SUCCESS;
 }
 
 int DeviceTy::disassociatePtr(void *HstPtrBegin) {
   DataMapMtx.lock();
 
   auto search = HostDataToTargetMap.find(HstPtrBeginTy{(uintptr_t)HstPtrBegin});
   if (search != HostDataToTargetMap.end()) {
     // Mapping exists
-    if (search->isRefCountInf()) {
+    if (search->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\n");
+    } else if (search->isDynRefCountInf()) {
       DP("Association found, removing it\n");
       HostDataToTargetMap.erase(search);
       DataMapMtx.unlock();
       return OFFLOAD_SUCCESS;
     } else {
       REPORT("Trying to disassociate a pointer which was not mapped via "
              "omp_target_associate_ptr\n");
     }
+  } else {
+    REPORT("Association not found\n");
   }
 
   // Mapping not found
   DataMapMtx.unlock();
-  REPORT("Association not found\n");
   return OFFLOAD_FAIL;
 }
 
 LookupResult DeviceTy::lookupMapping(void *HstPtrBegin, int64_t Size) {
   uintptr_t hp = (uintptr_t)HstPtrBegin;
   LookupResult lr;
 
   DP("Looking up mapping(HstPtrBegin=" DPxMOD ", Size=%" PRId64 ")...\n",
LookupResult DeviceTy::lookupMapping(void *HstPtrBegin, int64_t Size) {
   return lr;
 }
 
 TargetPointerResultTy
 DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
                            map_var_info_t HstPtrName, MoveDataStateTy MoveData,
                            bool IsImplicit, bool UpdateRefCount,
                            bool HasCloseModifier, bool HasPresentModifier,
-                           AsyncInfoTy &AsyncInfo) {
+                           bool HasHoldModifier, AsyncInfoTy &AsyncInfo) {
   void *TargetPointer = nullptr;
   bool IsHostPtr = false;
   bool IsNew = false;
 
   DataMapMtx.lock();
 
   LookupResult LR = lookupMapping(HstPtrBegin, Size);
   auto Entry = LR.Entry;
 
   // 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;
-    assert(HT.getRefCount() > 0 && "expected existing RefCount > 0");
-    if (UpdateRefCount)
+    IsNew = false;
+    const char *RefCountAction;
+    assert(HT.getTotalRefCount() > 0 && "expected existing RefCount > 0");
+    if (UpdateRefCount) {
       // After this, RefCount > 1.
-      HT.incRefCount();
-    else
+      HT.incRefCount(HasHoldModifier);
+      RefCountAction = " (incremented)";
+    } else {
       // It might have been allocated with the parent, but it's still new.
-      IsNew = HT.getRefCount() == 1;
+      IsNew = 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 ", RefCount=%s (%s), Name=%s\n",
+         ", Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n",
          (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(Ptr),
-         Size, HT.refCountToStr().c_str(),
-         UpdateRefCount ? "incremented" : "update suppressed",
+         Size, HT.dynRefCountToStr().c_str(), DynRefCountAction,
+         HT.holdRefCountToStr().c_str(), HoldRefCountAction,
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
     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)",
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;
     uintptr_t Ptr = (uintptr_t)allocData(Size, HstPtrBegin);
     Entry = HostDataToTargetMap
                 .emplace((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
-                         (uintptr_t)HstPtrBegin + Size, Ptr, HstPtrName)
+                         (uintptr_t)HstPtrBegin + Size, Ptr, HasHoldModifier,
+                         HstPtrName)
                 .first;
     INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
          "Creating new map entry with "
          "HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, "
-         "RefCount=%s, Name=%s\n",
-         DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size, Entry->refCountToStr().c_str(),
+         "DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
+         DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size,
+         Entry->dynRefCountToStr().c_str(), Entry->holdRefCountToStr().c_str(),
          (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
     TargetPointer = (void *)Ptr;
   }
 
   if (IsNew && MoveData == MoveDataStateTy::UNKNOWN)
     MoveData = MoveDataStateTy::REQUIRED;
 
   // If the target pointer is valid, and we need to transfer data, issue the
DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
 
   return {{IsNew, IsHostPtr}, Entry, TargetPointer};
 }
 
 // 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.
 void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-                               bool UpdateRefCount, bool &IsHostPtr,
-                               bool MustContain, bool ForceDelete) {
+                               bool UpdateRefCount, bool UseHoldRefCount,
+                               bool &IsHostPtr, bool MustContain,
+                               bool ForceDelete) {
   void *rc = NULL;
   IsHostPtr = false;
   IsLast = false;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
 
   if (lr.Flags.IsContained ||
       (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
     auto &HT = *lr.Entry;
-    // We do not decrement the reference count to zero here.  deallocTgtPtr does
-    // that atomically with removing the mapping.  Otherwise, before this thread
+    // We do not zero the total reference count here.  deallocTgtPtr does that
+    // atomically with removing the mapping.  Otherwise, before this thread
     // removed the mapping in deallocTgtPtr, another thread could retrieve the
     // mapping, increment and decrement back to zero, and then both threads
     // would try to remove the mapping, resulting in a double free.
-    IsLast = HT.decShouldRemove(ForceDelete);
+    IsLast = HT.decShouldRemove(UseHoldRefCount, ForceDelete);
     const char *RefCountAction;
     if (!UpdateRefCount) {
-      RefCountAction = "update suppressed";
+      RefCountAction = " (update suppressed)";
     } else if (ForceDelete) {
-      HT.resetRefCount();
-      assert(IsLast == HT.decShouldRemove() &&
+      HT.resetRefCount(UseHoldRefCount);
+      assert(IsLast == HT.decShouldRemove(UseHoldRefCount) &&
              "expected correct IsLast prediction for reset");
       if (IsLast)
-        RefCountAction = "reset, deferred final decrement";
-      else
-        RefCountAction = "reset";
+        RefCountAction = " (reset, deferred final decrement)";
+      else {
+        HT.decRefCount(UseHoldRefCount);
+        RefCountAction = " (reset)";
+      }
     } else if (IsLast) {
-      RefCountAction = "deferred final decrement";
+      RefCountAction = " (deferred final decrement)";
     } else {
-      RefCountAction = "decremented";
-      HT.decRefCount();
+      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 ", RefCount=%s (%s)\n",
-         DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.refCountToStr().c_str(),
-         RefCountAction);
+         "Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s\n",
+         DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.dynRefCountToStr().c_str(),
+         DynRefCountAction, HT.holdRefCountToStr().c_str(), HoldRefCountAction);
     rc = (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);
if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     uintptr_t tp = HT.TgtPtrBegin + (hp - HT.HstPtrBegin);
     return (void *)tp;
   }
 
   return NULL;
 }
 
 int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size,
-                            bool HasCloseModifier) {
+                            bool HasCloseModifier, bool HasHoldModifier) {
   if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
       !HasCloseModifier)
     return OFFLOAD_SUCCESS;
   // Check if the pointer is contained in any sub-nodes.
   int rc;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
-    if (HT.decRefCount() == 0) {
+    if (HT.decRefCount(HasHoldModifier) == 0) {
       DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n",
          DPxPTR(HT.TgtPtrBegin), Size);
       deleteData((void *)HT.TgtPtrBegin);
       INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
            "Removing map entry with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
            ", Size=%" PRId64 ", Name=%s\n",
            DPxPTR(HT.HstPtrBegin), DPxPTR(HT.TgtPtrBegin), Size,
            (HT.HstPtrName) ? getNameFromMapping(HT.HstPtrName).c_str()

openmp/libomptarget/src/omptarget.cpp

for (__tgt_offload_entry *CurrDeviceEntry = TargetTable->EntriesBegin,
         DP("Add mapping from host " DPxMOD " to device " DPxMOD " with size %zu"
            "\n",
            DPxPTR(CurrHostEntry->addr), DPxPTR(CurrDeviceEntry->addr),
            CurrDeviceEntry->size);
         Device.HostDataToTargetMap.emplace(
             (uintptr_t)CurrHostEntry->addr /*HstPtrBase*/,
             (uintptr_t)CurrHostEntry->addr /*HstPtrBegin*/,
             (uintptr_t)CurrHostEntry->addr + CurrHostEntry->size /*HstPtrEnd*/,
-            (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/, nullptr,
+            (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/,
+            false /*UseHoldRefCount*/, nullptr /*Name*/,
             true /*IsRefCountINF*/);
       }
     }
     Device.DataMapMtx.unlock();
   }
   PM->TrlTblMtx.unlock();
 
   if (rc != OFFLOAD_SUCCESS) {
for (int32_t i = 0; i < arg_num; ++i) {
     void *Pointer_HstPtrBegin, *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);
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       // 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 *), nullptr,
           MoveDataStateTy::NONE, IsImplicit, UpdateRef, HasCloseModifier,
-          HasPresentModifier, AsyncInfo);
+          HasPresentModifier, HasHoldModifier, AsyncInfo);
       PointerTgtPtrBegin = Pointer_TPR.TargetPointer;
       IsHostPtr = Pointer_TPR.Flags.IsHostPointer;
       if (!PointerTgtPtrBegin) {
         REPORT("Call to getOrAllocTgtPtr returned null pointer (%s).\n",
                HasPresentModifier ? "'present' map type modifier"
                                   : "device failure or illegal mapping");
         return OFFLOAD_FAIL;
       }
for (int32_t i = 0; i < arg_num; ++i) {
     const bool HasFlagTo = arg_types[i] & OMP_TGT_MAPTYPE_TO;
     const bool HasFlagAlways = arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS;
     if (HasFlagTo && (!UseUSM || HasCloseModifier))
       MoveData = HasFlagAlways ? MoveDataStateTy::REQUIRED
                                : MoveData = MoveDataStateTy::UNKNOWN;
 
     auto TPR = Device.getTargetPointer(
         HstPtrBegin, HstPtrBase, data_size, HstPtrName, MoveData, IsImplicit,
-        UpdateRef, HasCloseModifier, HasPresentModifier, AsyncInfo);
+        UpdateRef, HasCloseModifier, HasPresentModifier, HasHoldModifier,
+        AsyncInfo);
     void *TgtPtrBegin = TPR.TargetPointer;
     IsHostPtr = TPR.Flags.IsHostPointer;
     // 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 getOrAllocTgtPtr returned null pointer (%s).\n",
              HasPresentModifier ? "'present' map type modifier"
                                 : "device failure or illegal mapping");

 /// used to call the function \p DeviceTy::deallocTgtPtr.
 struct DeallocTgtPtrInfo {
   /// Host pointer used to look up into the map table
   void *HstPtrBegin;
   /// Size of the data
   int64_t DataSize;
   /// Whether it has \p close modifier
   bool HasCloseModifier;
+  /// Whether it has \p ompx_hold modifier
+  bool HasHoldModifier;
 
-  DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool HasCloseModifier)
-      : HstPtrBegin(HstPtr), DataSize(Size),
-        HasCloseModifier(HasCloseModifier) {}
+  DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool HasCloseModifier,
+                    bool HasHoldModifier)
+      : HstPtrBegin(HstPtr), DataSize(Size), HasCloseModifier(HasCloseModifier),
+        HasHoldModifier(HasHoldModifier) {}
 };
 } // namespace
 
 /// Internal function to undo the mapping and retrieve the data from the device.
 int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
                   void **ArgBases, void **Args, int64_t *ArgSizes,
                   int64_t *ArgTypes, map_var_info_t *ArgNames,
                   void **ArgMappers, AsyncInfoTy &AsyncInfo, bool FromMapper) {
for (int32_t I = ArgNum - 1; I >= 0; --I) {
     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 HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
     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
-    void *TgtPtrBegin =
-        Device.getTgtPtrBegin(HstPtrBegin, DataSize, IsLast, UpdateRef,
-                              IsHostPtr, !IsImplicit, ForceDelete);
+    void *TgtPtrBegin = Device.getTgtPtrBegin(
+        HstPtrBegin, DataSize, IsLast, UpdateRef, HasHoldModifier, IsHostPtr,
+        !IsImplicit, ForceDelete);
     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 ((ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
         } else {
           ++Itr;
         }
       }
       Device.ShadowMtx.unlock();
 
       // Add pointer to the buffer for later deallocation
       if (DelEntry)
-        DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier);
+        DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier,
+                                    HasHoldModifier);
     }
   }
 
   // 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();
   if (Ret != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   // Deallocate target pointer
   for (DeallocTgtPtrInfo &Info : DeallocTgtPtrs) {
     if (FromMapperBase && FromMapperBase == Info.HstPtrBegin)
       continue;
     Ret = Device.deallocTgtPtr(Info.HstPtrBegin, Info.DataSize,
-                               Info.HasCloseModifier);
+                               Info.HasCloseModifier, Info.HasHoldModifier);
     if (Ret != OFFLOAD_SUCCESS) {
       REPORT("Deallocating data from device failed.\n");
       return OFFLOAD_FAIL;
     }
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 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;
-  void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSize, IsLast, false,
-                                            IsHostPtr, /*MustContain=*/true);
+  void *TgtPtrBegin = Device.getTgtPtrBegin(
+      HstPtrBegin, ArgSize, IsLast, /*UpdateRefCount=*/false,
+      /*UseHoldRefCount=*/false, IsHostPtr, /*MustContain=*/true);
   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;
     }
if (!(ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {
         void *HstPtrBase = Args[Idx];
         bool IsLast, IsHostPtr; // 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();
-        PointerTgtPtrBegin = Device.getTgtPtrBegin(HstPtrVal, ArgSizes[I],
-                                                   IsLast, false, IsHostPtr);
+        PointerTgtPtrBegin = Device.getTgtPtrBegin(
+            HstPtrVal, ArgSizes[I], IsLast, /*UpdateRefCount=*/false,
+            /*UseHoldRefCount=*/false, IsHostPtr);
         if (!PointerTgtPtrBegin) {
           DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
              DPxPTR(HstPtrVal));
           continue;
         }
         if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
             TgtPtrBegin == HstPtrBegin) {
           DP("Unified memory is active, no need to map lambda captured"
if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {
         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);
       TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
-                                          false, IsHostPtr);
+                                          /*UpdateRefCount=*/false,
+                                          /*UseHoldRefCount=*/false, IsHostPtr);
       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();

openmp/libomptarget/src/private.h

static inline void dumpTargetPointerMappings(const ident_t *Loc,
                                              DeviceTy &Device) {
   if (Device.HostDataToTargetMap.empty())
     return;
 
   SourceInfo Kernel(Loc);
   INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
        "OpenMP Host-Device pointer mappings after block at %s:%d:%d:\n",
        Kernel.getFilename(), Kernel.getLine(), Kernel.getColumn());
-  INFO(OMP_INFOTYPE_ALL, Device.DeviceID, "%-18s %-18s %s %s %s\n", "Host Ptr",
-       "Target Ptr", "Size (B)", "RefCount", "Declaration");
+  INFO(OMP_INFOTYPE_ALL, Device.DeviceID, "%-18s %-18s %s %s %s %s\n",
+       "Host Ptr", "Target Ptr", "Size (B)", "DynRefCount", "HoldRefCount",
+       "Declaration");
   Device.DataMapMtx.lock();
   for (const auto &HostTargetMap : Device.HostDataToTargetMap) {
     SourceInfo Info(HostTargetMap.HstPtrName);
     INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
-         DPxMOD " " DPxMOD " %-8" PRIuPTR " %-8s %s at %s:%d:%d\n",
+         DPxMOD " " DPxMOD " %-8" PRIuPTR " %-11s %-12s %s at %s:%d:%d\n",
          DPxPTR(HostTargetMap.HstPtrBegin), DPxPTR(HostTargetMap.TgtPtrBegin),
          HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin,
-         HostTargetMap.refCountToStr().c_str(), Info.getName(),
+         HostTargetMap.dynRefCountToStr().c_str(),
+         HostTargetMap.holdRefCountToStr().c_str(), Info.getName(),
          Info.getFilename(), Info.getLine(), Info.getColumn());
   }
   Device.DataMapMtx.unlock();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// Print out the names and properties of the arguments to each kernel
 static inline void

openmp/libomptarget/test/mapping/ompx_hold/omp_target_disassociate_ptr.c

+// omp_target_disassociate_ptr should always fail if the hold reference count is
+// non-zero, regardless of the dynamic reference count.  When the latter is
+// finite, the implementation happens to choose to report the hold diagnostic.
+
+// RUN: %libomptarget-compile-generic -fopenmp-extensions
+// RUN: %not %libomptarget-run-generic 0   2>&1 | %fcheck-generic
+// RUN: %not %libomptarget-run-generic 1   2>&1 | %fcheck-generic
+// RUN: %not %libomptarget-run-generic inf 2>&1 | %fcheck-generic
+
+// RUN: %libomptarget-compile-generic -fopenmp-extensions -DHOLD_MORE
+// RUN: %not %libomptarget-run-generic 0   2>&1 | %fcheck-generic
+// RUN: %not %libomptarget-run-generic 1   2>&1 | %fcheck-generic
+// RUN: %not %libomptarget-run-generic inf 2>&1 | %fcheck-generic
+
+#include <omp.h>
+#include <stdio.h>
+#include <limits.h>
+#include <string.h>
+
+int main(int argc, char *argv[]) {
+  // Parse command line.
+  int DynRef;
+  if (argc != 2) {
+    fprintf(stderr, "bad arguments\n");
+    return 1;
+  }
+  if (0 == strcmp(argv[1], "inf"))
+    DynRef = INT_MAX;
+  else
+    DynRef = atoi(argv[1]);
+
+  // Allocate and set dynamic reference count as specified.
+  int DevNum = omp_get_default_device();
+  int X;
+  void *XDev = omp_target_alloc(sizeof X, DevNum);
+  if (!XDev) {
+    fprintf(stderr, "omp_target_alloc failed\n");
+    return 1;
+  }
+  if (DynRef == INT_MAX) {
+    if (omp_target_associate_ptr(&X, &XDev, sizeof X, 0, DevNum)) {
+      fprintf(stderr, "omp_target_associate_ptr failed\n");
+      return 1;
+    }
+  } else {
+    for (int I = 0; I < DynRef; ++I) {
+      #pragma omp target enter data map(alloc: X)
+    }
+  }
+
+  // Disassociate while hold reference count > 0.
+  int Status = 0;
+  #pragma omp target data map(ompx_hold,alloc: X)
+#if HOLD_MORE
+  #pragma omp target data map(ompx_hold,alloc: X)
+  #pragma omp target data map(ompx_hold,alloc: X)
+#endif
+  {
+    //      CHECK: Libomptarget error: Trying to disassociate a pointer with a
+    // CHECK-SAME: non-zero hold reference count
+    // CHECK-NEXT: omp_target_disassociate_ptr failed
+    if (omp_target_disassociate_ptr(&X, DevNum)) {
+      fprintf(stderr, "omp_target_disassociate_ptr failed\n");
+      Status = 1;
+    }
+  }
+  return Status;
+}

openmp/libomptarget/test/mapping/ompx_hold/struct.c

+// RUN: %libomptarget-compile-generic -fopenmp-extensions
+// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
+
+#include <omp.h>
+#include <stdio.h>
+
+#define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
+  printf("    presence of %s, %s, %s: %d, %d, %d\n",                           \
+         #Var1, #Var2, #Var3,                                                  \
+         omp_target_is_present(&(Var1), omp_get_default_device()),             \
+         omp_target_is_present(&(Var2), omp_get_default_device()),             \
+         omp_target_is_present(&(Var3), omp_get_default_device()))
+
+#define CHECK_VALUES(Var1, Var2)                                               \
+  printf("    values of %s, %s: %d, %d\n",                                     \
+         #Var1, #Var2, (Var1), (Var2))
+
+int main() {
+  struct S { int i; int j; } s;
+  // CHECK: presence of s, s.i, s.j: 0, 0, 0
+  CHECK_PRESENCE(s, s.i, s.j);
+
+  // =======================================================================
+  // Check that ompx_hold keeps entire struct present.
+
+  // -----------------------------------------------------------------------
+  // CHECK-LABEL: check:{{.*}}
+  printf("check: ompx_hold only on first member\n");
+  s.i = 20;
+  s.j = 30;
+  #pragma omp target data map(tofrom: s) map(ompx_hold,tofrom: s.i) \
+                                         map(tofrom: s.j)
+  {
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    CHECK_PRESENCE(s, s.i, s.j);
+    #pragma omp target map(tofrom: s)
+    {
+      s.i = 21;
+      s.j = 31;
+    }
+    #pragma omp target exit data map(delete: s, s.i)
+    // ompx_hold on s.i applies to all of s.
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    // CHECK-NEXT: values of s.i, s.j: 20, 30
+    CHECK_PRESENCE(s, s.i, s.j);
+    CHECK_VALUES(s.i, s.j);
+  }
+  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+  // CHECK-NEXT: values of s.i, s.j: 21, 31
+  CHECK_PRESENCE(s, s.i, s.j);
+  CHECK_VALUES(s.i, s.j);
+
+  // -----------------------------------------------------------------------
+  // CHECK-LABEL: check:{{.*}}
+  printf("check: ompx_hold only on last member\n");
+  s.i = 20;
+  s.j = 30;
+  #pragma omp target data map(tofrom: s) map(tofrom: s.i) \
+                                         map(ompx_hold,tofrom: s.j)
+  {
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    CHECK_PRESENCE(s, s.i, s.j);
+    #pragma omp target map(tofrom: s)
+    {
+      s.i = 21;
+      s.j = 31;
+    }
+    #pragma omp target exit data map(delete: s, s.i)
+    // ompx_hold on s.j applies to all of s.
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    // CHECK-NEXT: values of s.i, s.j: 20, 30
+    CHECK_PRESENCE(s, s.i, s.j);
+    CHECK_VALUES(s.i, s.j);
+  }
+  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+  // CHECK-NEXT: values of s.i, s.j: 21, 31
+  CHECK_PRESENCE(s, s.i, s.j);
+  CHECK_VALUES(s.i, s.j);
+
+  // -----------------------------------------------------------------------
+  // CHECK-LABEL: check:{{.*}}
+  printf("check: ompx_hold only on struct\n");
+  s.i = 20;
+  s.j = 30;
+  #pragma omp target data map(ompx_hold,tofrom: s) map(tofrom: s.i) \
+                                              map(tofrom: s.j)
+  {
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    CHECK_PRESENCE(s, s.i, s.j);
+    #pragma omp target map(tofrom: s)
+    {
+      s.i = 21;
+      s.j = 31;
+    }
+    #pragma omp target exit data map(delete: s, s.i)
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    // CHECK-NEXT: values of s.i, s.j: 20, 30
+    CHECK_PRESENCE(s, s.i, s.j);
+    CHECK_VALUES(s.i, s.j);
+  }
+  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+  // CHECK-NEXT: values of s.i, s.j: 21, 31
+  CHECK_PRESENCE(s, s.i, s.j);
+  CHECK_VALUES(s.i, s.j);
+
+  // =======================================================================
+  // Check that transfer to/from host checks reference count correctly.
+
+  // -----------------------------------------------------------------------
+  // CHECK-LABEL: check:{{.*}}
+  printf("check: parent DynRefCount=1 is not sufficient for transfer\n");
+  s.i = 20;
+  s.j = 30;
+  #pragma omp target data map(ompx_hold, tofrom: s)
+  #pragma omp target data map(ompx_hold, tofrom: s)
+  {
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    CHECK_PRESENCE(s, s.i, s.j);
+    #pragma omp target map(from: s.i, s.j)
+    {
+      s.i = 21;
+      s.j = 31;
+    } // No transfer here even though parent's DynRefCount=1.
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    // CHECK-NEXT: values of s.i, s.j: 20, 30
+    CHECK_PRESENCE(s, s.i, s.j);
+    CHECK_VALUES(s.i, s.j);
+    #pragma omp target map(to: s.i, s.j)
+    { // No transfer here even though parent's DynRefCount=1.
+      // CHECK-NEXT: values of s.i, s.j: 21, 31
+      CHECK_VALUES(s.i, s.j);
+    }
+  }
+  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+  // CHECK-NEXT: values of s.i, s.j: 21, 31
+  CHECK_PRESENCE(s, s.i, s.j);
+  CHECK_VALUES(s.i, s.j);
+
+  // -----------------------------------------------------------------------
+  // CHECK-LABEL: check:{{.*}}
+  printf("check: parent HoldRefCount=1 is not sufficient for transfer\n");
+  s.i = 20;
+  s.j = 30;
+  #pragma omp target data map(tofrom: s)
+  #pragma omp target data map(tofrom: s)
+  {
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    CHECK_PRESENCE(s, s.i, s.j);
+    #pragma omp target map(ompx_hold, from: s.i, s.j)
+    {
+      s.i = 21;
+      s.j = 31;
+    } // No transfer here even though parent's HoldRefCount=1.
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    // CHECK-NEXT: values of s.i, s.j: 20, 30
+    CHECK_PRESENCE(s, s.i, s.j);
+    CHECK_VALUES(s.i, s.j);
+    #pragma omp target map(ompx_hold, to: s.i, s.j)
+    { // No transfer here even though parent's HoldRefCount=1.
+      // CHECK-NEXT: values of s.i, s.j: 21, 31
+      CHECK_VALUES(s.i, s.j);
+    }
+  }
+  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+  // CHECK-NEXT: values of s.i, s.j: 21, 31
+  CHECK_PRESENCE(s, s.i, s.j);
+  CHECK_VALUES(s.i, s.j);
+
+  // -----------------------------------------------------------------------
+  // CHECK-LABEL: check:{{.*}}
+  //
+  // At the beginning of a region, if the parent's TotalRefCount=1, then the
+  // transfer should happen.
+  //
+  // At the end of a region, it also must be true that the reference count being
+  // decremented is the reference count that is 1.
+  printf("check: parent TotalRefCount=1 is not sufficient for transfer\n");
+  s.i = 20;
+  s.j = 30;
+  #pragma omp target data map(ompx_hold, tofrom: s)
+  {
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    CHECK_PRESENCE(s, s.i, s.j);
+    #pragma omp target map(ompx_hold, tofrom: s.i, s.j)
+    {
+      s.i = 21;
+      s.j = 31;
+    }
+    #pragma omp target exit data map(from: s.i, s.j)
+    // No transfer here even though parent's TotalRefCount=1.
+    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+    // CHECK-NEXT: values of s.i, s.j: 20, 30
+    CHECK_PRESENCE(s, s.i, s.j);
+    CHECK_VALUES(s.i, s.j);
+  }
+  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+  // CHECK-NEXT: values of s.i, s.j: 21, 31
+  CHECK_PRESENCE(s, s.i, s.j);
+  CHECK_VALUES(s.i, s.j);
+
+  return 0;
+}

openmp/libomptarget/test/mapping/ompx_hold/target-data.c

+// RUN: %libomptarget-compile-generic -fopenmp-extensions
+// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
+
+#include <omp.h>
+#include <stdio.h>
+
+#define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
+  printf("    presence of %s, %s, %s: %d, %d, %d\n",                           \
+         #Var1, #Var2, #Var3,                                                  \
+         omp_target_is_present(&Var1, omp_get_default_device()),               \
+         omp_target_is_present(&Var2, omp_get_default_device()),               \
+         omp_target_is_present(&Var3, omp_get_default_device()))
+
+int main() {
+  int m, r, d;
+  // CHECK: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // -----------------------------------------------------------------------
+  // CHECK-NEXT: check:{{.*}}
+  printf("check: dyn>0, hold=0, dec/reset dyn=0\n");
+
+  // CHECK-NEXT: structured{{.*}}
+  printf("  structured dec of dyn\n");
+  #pragma omp target data map(tofrom: m) map(alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(tofrom: m) map(alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // CHECK-NEXT: dynamic{{.*}}
+  printf("  dynamic dec/reset of dyn\n");
+  #pragma omp target data map(tofrom: m) map(alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(tofrom: m) map(alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target exit data map(from: m) map(release: r)
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+      // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+    // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // -----------------------------------------------------------------------
+  // CHECK: check:{{.*}}
+  printf("check: dyn=0, hold>0, dec/reset dyn=0, dec hold=0\n");
+
+  // Structured dec of dyn would require dyn>0.
+
+  // CHECK-NEXT: dynamic{{.*}}
+  printf("  dynamic dec/reset of dyn\n");
+  #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(ompx_hold, tofrom: m) \
+                            map(ompx_hold, alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target exit data map(from: m) map(release: r)
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // -----------------------------------------------------------------------
+  // CHECK: check:{{.*}}
+  printf("check: dyn>0, hold>0, dec/reset dyn=0, dec hold=0\n");
+
+  // CHECK-NEXT: structured{{.*}}
+  printf("  structured dec of dyn\n");
+  #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(ompx_hold, tofrom: m) \
+                            map(ompx_hold, alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target data map(tofrom: m) map(alloc: r, d)
+      {
+        // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+        CHECK_PRESENCE(m, r, d);
+        #pragma omp target data map(tofrom: m) map(alloc: r, d)
+        {
+          // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+          CHECK_PRESENCE(m, r, d);
+        }
+        // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+        CHECK_PRESENCE(m, r, d);
+      }
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // CHECK-NEXT: dynamic{{.*}}
+  printf("  dynamic dec/reset of dyn\n");
+  #pragma omp target enter data map(to: m) map(alloc: r, d)
+  // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+  CHECK_PRESENCE(m, r, d);
+  #pragma omp target enter data map(to: m) map(alloc: r, d)
+  // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+  CHECK_PRESENCE(m, r, d);
+  #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(ompx_hold, tofrom: m) \
+                            map(ompx_hold, alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target exit data map(from: m) map(release: r)
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // -----------------------------------------------------------------------
+  // CHECK: check:{{.*}}
+  printf("check: dyn>0, hold>0, dec hold=0, dec/reset dyn=0\n");
+
+  // CHECK-NEXT: structured{{.*}}
+  printf("  structured dec of dyn\n");
+  #pragma omp target data map(tofrom: m) map(alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(tofrom: m) map(alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target data map(ompx_hold, tofrom: m) \
+                              map(ompx_hold, alloc: r, d)
+      {
+        // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+        CHECK_PRESENCE(m, r, d);
+        #pragma omp target data map(ompx_hold, tofrom: m) \
+                                map(ompx_hold, alloc: r, d)
+        {
+          // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+          CHECK_PRESENCE(m, r, d);
+        }
+        // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+        CHECK_PRESENCE(m, r, d);
+      }
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // CHECK-NEXT: dynamic{{.*}}
+  printf("  dynamic dec/reset of dyn\n");
+  #pragma omp target enter data map(to: m) map(alloc: r, d)
+  // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+  CHECK_PRESENCE(m, r, d);
+  #pragma omp target enter data map(to: m) map(alloc: r, d)
+  // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+  CHECK_PRESENCE(m, r, d);
+  #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(ompx_hold, tofrom: m) \
+                            map(ompx_hold, alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+  CHECK_PRESENCE(m, r, d);
+  #pragma omp target exit data map(from: m) map(release: r)
+  // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+  CHECK_PRESENCE(m, r, d);
+  #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  return 0;
+}

openmp/libomptarget/test/mapping/ompx_hold/target.c

+// RUN: %libomptarget-compile-generic -fopenmp-extensions
+// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
+
+#include <omp.h>
+#include <stdio.h>
+
+#define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
+  printf("    presence of %s, %s, %s: %d, %d, %d\n",                           \
+         #Var1, #Var2, #Var3,                                                  \
+         omp_target_is_present(&Var1, omp_get_default_device()),               \
+         omp_target_is_present(&Var2, omp_get_default_device()),               \
+         omp_target_is_present(&Var3, omp_get_default_device()))
+
+int main() {
+  int m, r, d;
+  // CHECK: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // -----------------------------------------------------------------------
+  // CHECK-NEXT: check:{{.*}}
+  printf("check: dyn>0, hold=0, dec dyn=0\n");
+
+  // CHECK-NEXT: once
+  printf("  once\n");
+  #pragma omp target map(tofrom: m) map(alloc: r, d)
+  ;
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // CHECK-NEXT: twice
+  printf("  twice\n");
+  #pragma omp target data map(tofrom: m) map(alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target map(tofrom: m) map(alloc: r, d)
+    ;
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // -----------------------------------------------------------------------
+  // CHECK: check:{{.*}}
+  printf("check: dyn=0, hold>0, dec hold=0\n");
+
+  // CHECK-NEXT: once
+  printf("  once\n");
+  #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+  ;
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // CHECK-NEXT: twice
+  printf("  twice\n");
+  #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+    ;
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // -----------------------------------------------------------------------
+  // CHECK: check:{{.*}}
+  printf("check: dyn>0, hold>0, dec dyn=0, dec hold=0\n");
+
+  // CHECK-NEXT: once each
+  printf("  once each\n");
+  #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target map(tofrom: m) map(alloc: r, d)
+    ;
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // CHECK-NEXT: twice each
+  printf("  twice each\n");
+  #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(ompx_hold, tofrom: m) \
+                            map(ompx_hold, alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target data map(tofrom: m) map(alloc: r, d)
+      {
+        // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+        CHECK_PRESENCE(m, r, d);
+        #pragma omp target map(tofrom: m) map(alloc: r, d)
+        ;
+        // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+        CHECK_PRESENCE(m, r, d);
+      }
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // -----------------------------------------------------------------------
+  // CHECK: check:{{.*}}
+  printf("check: dyn>0, hold>0, dec hold=0, dec dyn=0\n");
+
+  // CHECK-NEXT: once each
+  printf("  once each\n");
+  #pragma omp target data map(tofrom: m) map(alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+    ;
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  // CHECK-NEXT: twice each
+  printf("  twice each\n");
+  #pragma omp target data map(tofrom: m) map(alloc: r, d)
+  {
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+    #pragma omp target data map(tofrom: m) map(alloc: r, d)
+    {
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+      #pragma omp target data map(ompx_hold, tofrom: m) \
+                              map(ompx_hold, alloc: r, d)
+      {
+        // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+        CHECK_PRESENCE(m, r, d);
+        #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+        ;
+        // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+        CHECK_PRESENCE(m, r, d);
+      }
+      // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+      CHECK_PRESENCE(m, r, d);
+    }
+    // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+    CHECK_PRESENCE(m, r, d);
+  }
+  // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+  CHECK_PRESENCE(m, r, d);
+
+  return 0;
+}

openmp/libomptarget/test/offloading/info.c

-// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -gline-tables-only && env LIBOMPTARGET_INFO=63 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda \
+// RUN:     -gline-tables-only -fopenmp-extensions
+// RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | \
+// RUN:   %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
 // REQUIRES: nvptx64-nvidia-cuda
 
 #include <stdio.h>
 #include <omp.h>
 
 #define N 64
 
 #pragma omp declare target
 int global;
 #pragma omp end declare target
 
 extern void __tgt_set_info_flag(unsigned);
 
 int main() {
   int A[N];
   int B[N];
   int C[N];
   int val = 1;
 
 // INFO: CUDA device 0 info: Device supports up to {{[0-9]+}} CUDA blocks and {{[0-9]+}} threads with a warp size of {{[0-9]+}}
 // INFO: Libomptarget device 0 info: Entering OpenMP data region at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
 // INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
 // INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
 // INFO: Libomptarget device 0 info: to(C[0:64])[256]
-// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=A[0:64]
-// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=B[0:64]
+// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=A[0:64]
+// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=0, HoldRefCount=1, Name=B[0:64]
 // INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64]
-// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=C[0:64]
+// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=C[0:64]
 // INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64]
 // INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
-// INFO: Libomptarget device 0 info: Host Ptr           Target Ptr         Size (B) RefCount Declaration
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: Host Ptr           Target Ptr         Size (B) DynRefCount HoldRefCount Declaration
+// INFO: Libomptarget device 0 info: {{.*}}             {{.*}}             256      1           0            C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: {{.*}}             {{.*}}             256      0           1            B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: {{.*}}             {{.*}}             256      1           0            A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
 // INFO: Libomptarget device 0 info: Entering OpenMP kernel at info.c:{{[0-9]+}}:{{[0-9]+}} with 1 arguments:
 // INFO: Libomptarget device 0 info: firstprivate(val)[4]
 // INFO: CUDA device 0 info: Launching kernel __omp_offloading_{{.*}}main{{.*}} with {{[0-9]+}} blocks and {{[0-9]+}} threads in Generic mode
 // INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
-// INFO: Libomptarget device 0 info: Host Ptr           Target Ptr         Size (B) RefCount Declaration
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: Host Ptr           Target Ptr         Size (B) DynRefCount HoldRefCount Declaration
+// INFO: Libomptarget device 0 info: {{.*}}             {{.*}}             256      1           0            C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: {{.*}}             {{.*}}             256      0           1            B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: {{.*}}             {{.*}}             256      1           0            A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
 // INFO: Libomptarget device 0 info: Exiting OpenMP data region at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
 // INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
 // INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
 // INFO: Libomptarget device 0 info: to(C[0:64])[256]
 // INFO: Libomptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=256, Name=B[0:64]
 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
 // INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:[[#%u,]]:[[#%u,]]:
-// INFO: Libomptarget device 0 info: Host Ptr  Target Ptr Size (B) RefCount Declaration
-// INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]]  4        INF      unknown at unknown:0:0
-#pragma omp target data map(alloc:A[0:N]) map(tofrom:B[0:N]) map(to:C[0:N])
+// INFO: Libomptarget device 0 info: Host Ptr  Target Ptr Size (B) DynRefCount HoldRefCount Declaration
+// INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]]  4        INF         0            unknown at unknown:0:0
+#pragma omp target data map(alloc:A[0:N]) map(ompx_hold,tofrom:B[0:N]) map(to:C[0:N])
 #pragma omp target firstprivate(val)
   { val = 1; }
 
   __tgt_set_info_flag(0x0);
 // INFO-NOT: Libomptarget device 0 info: {{.*}}
 #pragma omp target
   { }
 
   return 0;
 }