[LIBOMPTARGET] Do not increment/decrement the refcount for "declare target" objects

[grokos]

Summary

The reference counter for global objects marked with declare target is INF. This patch prevents the runtime from incrementing /decrementing INF refcounts. Without it, the map(delete: global_object) directive actually deallocates the global on the device. With this patch, such a directive becomes a no-op.

[jdoerfert]

I'm generally fine with this. To avoid future problems we could make the RefCount private and introduce modifier methods that do the CONSIDERED_INF check always.

[grokos]

I like the abstraction approach more, too. Updated the diff.

[ABataev]

openmp/libomptarget/src/device.h

54

You don't need inline here, it is inlined automatically, I assume, since defined in class.

59

Maybe just use a bitfield/bool flag instead?

[ABataev]

openmp/libomptarget/src/device.h

54

Const function

77

Const.

[JonChesterfield]

Change looks good to me as is. Some notes / questions on the reference counting.

We may regret using long instead of int64_t, if this is used on x86-64 with 32 bit long. E.g. windows. Are we sure a 32 bit counter can't overflow in pathological cases?

Is the signed counter type a performance optimization to hit the UB on overflow? Expected it to be unsigned as the value is never negative (iiuc).

An alternative to the bool is to use ~0 to indicate infinite count. That works on the basis that we don't increment the recount if it's inf, so no overflow, and that recounts can't reach uint64_max on a nominally 64 bit address space.

The bool/long or sentinel value are probably subtle enough to warrant a class recount, which likely already exists in llvm somewhere.

[grokos]

That's a good observation, although I can't think of a normal-use scenario where the refcount could overflow. Of course it can overflow in pathological cases like:

for (huge_number_of_iterations) {
#pragma omp target enter data map (var)
}

but in such problematic code even a 64-bit counter could overflow. Actually the above example is taylor-made to intentionally make the refcount overflow.

[JonChesterfield]

In D72525#1817204, @grokos wrote:

but in such problematic code even a 64-bit counter could overflow. Actually the above example is taylor-made to intentionally make the refcount overflow.

Curiously enough, refcounts can't overflow when bit size matches or exceeds address space.

Each refcount increment indicates some live reference to the object which itself uses some space. So even if we were refcounting an object of size one, each time we copy that object (thus increment the refcount) we also need non-zero bytes of address space for the new reference to that object. That is, we run out of address space shortly before the refcount overflows, even in the most extreme case and disregarding the considerable time it takes to count to 2^48.

Unless it's a <=32 bit counter and we're using a 48 bit address space system, where repeatedly kicking the counter can overflow within a plausible length of time (iirc it's of the order of a few hours spinning on the counter).

[grokos]

Changed refcount type to uint64_t.

[jdoerfert]

I'm fine with this. If others are, please accept and don't wait for me.

openmp/libomptarget/src/device.h

33

Nit: public is not needed.

[JonChesterfield]

Thanks for the change to uint64_t

What do you think of computing INFRefCount as (RefCnt == INFRefCount), instead of using the separate bool? It's slightly less obvious what's going on than with the separate flag, but also means the two values can't accidentally drift out of sync.

[grokos]

Got rid of the extra bool, I agree it creates confusion...

[JonChesterfield]

Nice. Thanks!

openmp/libomptarget/src/device.h

 struct __tgt_bin_desc;
 struct __tgt_target_table;
 
 #define INF_REF_CNT (LONG_MAX>>1) // leave room for additions/subtractions
 #define CONSIDERED_INF(x) (x > (INF_REF_CNT>>1))
 
 /// Map between host data and target data.
 struct HostDataToTargetTy {
+private:
+  long RefCount;
+
+public:
   uintptr_t HstPtrBase; // host info.
   uintptr_t HstPtrBegin;
   uintptr_t HstPtrEnd; // non-inclusive.
 
   uintptr_t TgtPtrBegin; // target info.
 
-  long RefCount;
-
   HostDataToTargetTy()
       : HstPtrBase(0), HstPtrBegin(0), HstPtrEnd(0),
         TgtPtrBegin(0), RefCount(0) {}
   HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB)
       : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E),
         TgtPtrBegin(TB), RefCount(1) {}
   HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB,
       long RF)
       : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E),
         TgtPtrBegin(TB), RefCount(RF) {}
+
+  inline long getRefCount() {
+    return RefCount;
+  }
+
+  inline void setRefCount(long Value) {
+    if (!CONSIDERED_INF(RefCount))
+      RefCount = Value;
+  }
+
+  inline long incRefCount() {
+    if (!CONSIDERED_INF(RefCount))
+      ++RefCount;
+
+    return RefCount;
+  }
+
+  inline long decRefCount() {
+    if (!CONSIDERED_INF(RefCount))
+      --RefCount;
+
+    return RefCount;
+  }
+
+  inline bool isRefCountInf() {
+    return CONSIDERED_INF(RefCount);
+  }
 };
 
 typedef std::list<HostDataToTargetTy> HostDataToTargetListTy;
 
 struct LookupResult {
   struct {
     unsigned IsContained   : 1;
     unsigned ExtendsBefore : 1;

openmp/libomptarget/src/device.cpp

int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
   HostDataToTargetTy newEntry;
 
   // Set up missing fields
   newEntry.HstPtrBase = (uintptr_t) HstPtrBegin;
   newEntry.HstPtrBegin = (uintptr_t) HstPtrBegin;
   newEntry.HstPtrEnd = (uintptr_t) HstPtrBegin + Size;
   newEntry.TgtPtrBegin = (uintptr_t) TgtPtrBegin;
   // refCount must be infinite
-  newEntry.RefCount = INF_REF_CNT;
+  newEntry.setRefCount(INF_REF_CNT);
 
   DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD ", HstEnd="
       DPxMOD ", TgtBegin=" DPxMOD "\n", DPxPTR(newEntry.HstPtrBase),
       DPxPTR(newEntry.HstPtrBegin), DPxPTR(newEntry.HstPtrEnd),
       DPxPTR(newEntry.TgtPtrBegin));
   HostDataToTargetMap.push_front(newEntry);
 
   DataMapMtx.unlock();
 
   return OFFLOAD_SUCCESS;
 }
 
 int DeviceTy::disassociatePtr(void *HstPtrBegin) {
   DataMapMtx.lock();
 
   // Check if entry exists
   for (HostDataToTargetListTy::iterator ii = HostDataToTargetMap.begin();
       ii != HostDataToTargetMap.end(); ++ii) {
     if ((uintptr_t)HstPtrBegin == ii->HstPtrBegin) {
       // Mapping exists
-      if (CONSIDERED_INF(ii->RefCount)) {
+      if (ii->isRefCountInf()) {
         DP("Association found, removing it\n");
         HostDataToTargetMap.erase(ii);
         DataMapMtx.unlock();
         return OFFLOAD_SUCCESS;
       } else {
         DP("Trying to disassociate a pointer which was not mapped via "
             "omp_target_associate_ptr\n");
         break;

 long DeviceTy::getMapEntryRefCnt(void *HstPtrBegin) {
   uintptr_t hp = (uintptr_t)HstPtrBegin;
   long RefCnt = -1;
 
   DataMapMtx.lock();
   for (auto &HT : HostDataToTargetMap) {
     if (hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd) {
       DP("DeviceTy::getMapEntry: requested entry found\n");
-      RefCnt = HT.RefCount;
+      RefCnt = HT.getRefCount();
       break;
     }
   }
   DataMapMtx.unlock();
 
   if (RefCnt < 0) {
     DP("DeviceTy::getMapEntry: requested entry not found\n");
   }
void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
   // 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;
     IsNew = false;
 
     if (UpdateRefCount)
-      ++HT.RefCount;
+      HT.incRefCount();
 
     uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     DP("Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
         "Size=%ld,%s RefCount=%s\n", (IsImplicit ? " (implicit)" : ""),
         DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
         (UpdateRefCount ? " updated" : ""),
-        (CONSIDERED_INF(HT.RefCount)) ? "INF" :
-            std::to_string(HT.RefCount).c_str());
+        HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str());
     rc = (void *)tp;
   } else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) {
     // Explicit extension of mapped data - not allowed.
     DP("Explicit extension of mapping is not allowed.\n");
   } else if (Size) {
     // 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
void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
   void *rc = NULL;
   IsHostPtr = false;
   IsLast = false;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
 
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
-    IsLast = !(HT.RefCount > 1);
+    IsLast = HT.getRefCount() == 1;
 
-    if (HT.RefCount > 1 && UpdateRefCount)
-      --HT.RefCount;
+    if (!IsLast && UpdateRefCount)
+      HT.decRefCount();
 
     uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     DP("Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
         "Size=%ld,%s RefCount=%s\n", DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
         (UpdateRefCount ? " updated" : ""),
-        (CONSIDERED_INF(HT.RefCount)) ? "INF" :
-            std::to_string(HT.RefCount).c_str());
+        HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str());
     rc = (void *)tp;
   } else if (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=%ld RefCount=%s\n",
        DPxPTR((uintptr_t)HstPtrBegin), Size, (UpdateRefCount ? " updated" : ""));
     IsHostPtr = true;
if (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.RefCount = 1;
-    if (--HT.RefCount <= 0) {
-      assert(HT.RefCount == 0 && "did not expect a negative ref count");
+      HT.setRefCount(1);
+    if (HT.decRefCount() <= 0) {
+      assert(HT.getRefCount() == 0 && "did not expect a negative ref count");
       DP("Deleting tgt data " DPxMOD " of size %ld\n",
           DPxPTR(HT.TgtPtrBegin), Size);
       RTL->data_delete(RTLDeviceID, (void *)HT.TgtPtrBegin);
       DP("Removing%s mapping with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
           ", Size=%ld\n", (ForceDelete ? " (forced)" : ""),
           DPxPTR(HT.HstPtrBegin), DPxPTR(HT.TgtPtrBegin), Size);
       HostDataToTargetMap.erase(lr.Entry);
     }

openmp/libomptarget/test/mapping/delete_inf_refcount.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+#include <omp.h>
+
+#pragma omp declare target
+int isHost;
+#pragma omp end declare target
+
+int main(void) {
+  isHost = -1;
+
+#pragma omp target enter data map(to: isHost)
+
+#pragma omp target
+  { isHost = omp_is_initial_device(); }
+#pragma omp target update from(isHost)
+
+  if (isHost < 0) {
+    printf("Runtime error, isHost=%d\n", isHost);
+  }
+
+#pragma omp target exit data map(delete: isHost)
+
+  // CHECK: Target region executed on the device
+  printf("Target region executed on the %s\n", isHost ? "host" : "device");
+
+  return isHost;
+}