[OpenMP] Fix delete map type in ref count debug messages

[jdenny]

Summary

For example, without this patch:

$ cat test.c
int main() {
  int x;
  #pragma omp target enter data map(alloc: x)
  #pragma omp target enter data map(alloc: x)
  #pragma omp target enter data map(alloc: x)
  #pragma omp target exit data map(delete: x)
  ;
  return 0;
}
$ clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda test.c
$ LIBOMPTARGET_DEBUG=1 ./a.out |& grep 'Creating\|Mapping exists\|last'
Libomptarget --> Creating new map entry with HstPtrBegin=0x00007ffddf1eaea8, TgtPtrBegin=0x00000000013bb040, Size=4, RefCount=1, Name=unknown
Libomptarget --> Mapping exists with HstPtrBegin=0x00007ffddf1eaea8, TgtPtrBegin=0x00000000013bb040, Size=4, RefCount=2 (incremented), Name=unknown
Libomptarget --> Mapping exists with HstPtrBegin=0x00007ffddf1eaea8, TgtPtrBegin=0x00000000013bb040, Size=4, RefCount=3 (incremented), Name=unknown
Libomptarget --> Mapping exists with HstPtrBegin=0x00007ffddf1eaea8, TgtPtrBegin=0x00000000013bb040, Size=4, RefCount=2 (decremented)
Libomptarget --> There are 4 bytes allocated at target address 0x00000000013bb040 - is not last

RefCount is reported as decremented to 2, but it ought to be reset
because of the delete map type, and is not last is incorrect.

This patch migrates the reset of reference counts from
DeviceTy::deallocTgtPtr to DeviceTy::getTgtPtrBegin, which then
correctly reports the reset. Based on the IsLast result from
DeviceTy::getTgtPtrBegin, targetDataEnd then correctly reports `is
last` for any deletion. DeviceTy::deallocTgtPtr is responsible only
for the final reference count decrement and mapping removal.

An obscure side effect of this patch is that a delete map type when
the reference count is infinite yields DelEntry=IsLast=false in
targetDataEnd and so no longer results in a
DeviceTy::deallocTgtPtr call. Without this patch, that call is a
no-op anyway besides some unnecessary locking and mapping table
lookups.

[grokos]

openmp/libomptarget/src/device.cpp

361–362

It seems that after this patch ForceDelete is no longer necessary in the signature of deallocTgtPtr. There is only one instance where this variable is used and it's not a functional one; it is only used inside INFO to report whether the removal of a map entry is a forced one. I think that at this point such information is of no value since getTgtPtrBegin will already have printed out whether the entry is forcefully removed. Right?

[jdenny]

openmp/libomptarget/src/device.cpp

361–362

Good point. The "forcing" (the reset) doesn't happen here anymore. All that's left is a single decrement. I'll fix it. Thanks.

[jdenny]

Applied grokos's suggestion: Removed ForceDelete from deallocTgtPtr.

Rebased.

[grokos]

LGTM, thanks for the patch!

[jdenny]

Thanks for the review.

openmp/libomptarget/src/device.h

uint64_t decRefCount() const {
     return RefCount;
   }
 
   bool isRefCountInf() const { return RefCount == INFRefCount; }
 
   std::string refCountToStr() const {
     return isRefCountInf() ? "INF" : std::to_string(getRefCount());
   }
+
+  /// Should one decrement of the reference count (after resetting it if
+  /// \c AfterReset) remove this mapping?
+  bool decShouldRemove(bool AfterReset = false) const {
+    if (AfterReset)
+      return !isRefCountInf();
+    return getRefCount() == 1;
+  }
 };
 
 typedef uintptr_t HstPtrBeginTy;
 inline bool operator<(const HostDataToTargetTy &lhs, const HstPtrBeginTy &rhs) {
   return lhs.HstPtrBegin < rhs;
 }
 inline bool operator<(const HstPtrBeginTy &lhs, const HostDataToTargetTy &rhs) {
   return lhs < rhs.HstPtrBegin;
struct DeviceTy {
   LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
   void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
                          map_var_info_t HstPtrName, bool &IsNew,
                          bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount,
                          bool HasCloseModifier, bool HasPresentModifier);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
                        bool UpdateRefCount, bool &IsHostPtr,
-                       bool MustContain = false);
-  int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete,
+                       bool MustContain = false, bool ForceDelete = false);
+  int deallocTgtPtr(void *TgtPtrBegin, int64_t Size,
                     bool HasCloseModifier = false);
   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);
 

openmp/libomptarget/src/device.cpp

void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
   return rc;
 }
 
 // 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 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
     // 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.getRefCount() == 1;
+    IsLast = HT.decShouldRemove(ForceDelete);
     const char *RefCountAction;
-    if (!UpdateRefCount)
+    if (!UpdateRefCount) {
       RefCountAction = "update suppressed";
-    else if (IsLast)
+    } else if (ForceDelete) {
+      HT.resetRefCount();
+      assert(IsLast == HT.decShouldRemove() &&
+             "expected correct IsLast prediction for reset");
+      if (IsLast)
+        RefCountAction = "reset, deferred final decrement";
+      else
+        RefCountAction = "reset";
+    } else if (IsLast) {
       RefCountAction = "deferred final decrement";
-    else {
+    } else {
       RefCountAction = "decremented";
       HT.decRefCount();
     }
     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(),
if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
     uintptr_t tp = HT.TgtPtrBegin + (hp - HT.HstPtrBegin);
     return (void *)tp;
   }
 
   return NULL;
 }
 
-int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
+int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size,
                             bool HasCloseModifier) {
   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 (ForceDelete)
-      HT.resetRefCount();
     if (HT.decRefCount() == 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%s map entry with HstPtrBegin=" DPxMOD
-           ", TgtPtrBegin=" DPxMOD ", Size=%" PRId64 ", Name=%s\n",
-           (ForceDelete ? " (forced)" : ""), DPxPTR(HT.HstPtrBegin),
-           DPxPTR(HT.TgtPtrBegin), Size,
+           "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()
                            : "unknown");
       HostDataToTargetMap.erase(lr.Entry);
     }
     rc = OFFLOAD_SUCCESS;
   } else {
     REPORT("Section to delete (hst addr " DPxMOD ") does not exist in the"
            " allocated memory\n",

openmp/libomptarget/src/omptarget.cpp

 namespace {
 /// This structure contains information to deallocate a target pointer, aka.
 /// 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 is forced to be removed from the map table
-  bool ForceDelete;
   /// Whether it has \p close modifier
   bool HasCloseModifier;
 
-  DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool ForceDelete,
-                    bool HasCloseModifier)
-      : HstPtrBegin(HstPtr), DataSize(Size), ForceDelete(ForceDelete),
+  DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool HasCloseModifier)
+      : HstPtrBegin(HstPtr), DataSize(Size),
         HasCloseModifier(HasCloseModifier) {}
 };
 } // 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,
for (int32_t I = ArgNum - 1; I >= 0; --I) {
     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;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
-    void *TgtPtrBegin = Device.getTgtPtrBegin(
-        HstPtrBegin, DataSize, IsLast, UpdateRef, IsHostPtr, !IsImplicit);
+    void *TgtPtrBegin =
+        Device.getTgtPtrBegin(HstPtrBegin, DataSize, IsLast, UpdateRef,
+                              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
for (int32_t I = ArgNum - 1; I >= 0; --I) {
     // 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 || ForceDelete;
+    bool DelEntry = 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) {
         } else {
           ++Itr;
         }
       }
       Device.ShadowMtx.unlock();
 
       // Add pointer to the buffer for later deallocation
       if (DelEntry)
-        DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, ForceDelete,
-                                    HasCloseModifier);
+        DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier);
     }
   }
 
   // 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.ForceDelete, Info.HasCloseModifier);
+                               Info.HasCloseModifier);
     if (Ret != OFFLOAD_SUCCESS) {
       REPORT("Deallocating data from device failed.\n");
       return OFFLOAD_FAIL;
     }
   }
 
   return OFFLOAD_SUCCESS;
 }