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

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.

Differential Revision: https://reviews.llvm.org/D72525

openmp/libomptarget/src/device.cpp

@@ -44,16 +44,12 @@ int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
     }
   }
 
-  // Mapping does not exist, allocate it
+  // Mapping does not exist, allocate it with refCount=INF
-  HostDataToTargetTy newEntry;
+  HostDataToTargetTy newEntry((uintptr_t) HstPtrBegin /*HstPtrBase*/,
-
+                              (uintptr_t) HstPtrBegin /*HstPtrBegin*/,
-  // Set up missing fields
+                              (uintptr_t) HstPtrBegin + Size /*HstPtrEnd*/,
-  newEntry.HstPtrBase = (uintptr_t) HstPtrBegin;
+                              (uintptr_t) TgtPtrBegin /*TgtPtrBegin*/,
-  newEntry.HstPtrBegin = (uintptr_t) HstPtrBegin;
+                              true /*IsRefCountINF*/);
-  newEntry.HstPtrEnd = (uintptr_t) HstPtrBegin + Size;
-  newEntry.TgtPtrBegin = (uintptr_t) TgtPtrBegin;
-  // refCount must be infinite
-  newEntry.RefCount = INF_REF_CNT;
 
   DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD ", HstEnd="
       DPxMOD ", TgtBegin=" DPxMOD "\n", DPxPTR(newEntry.HstPtrBase),
@@ -74,7 +70,7 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) {
       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();
@@ -94,21 +90,21 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) {
 }
 
 // Get ref count of map entry containing HstPtrBegin
-long DeviceTy::getMapEntryRefCnt(void *HstPtrBegin) {
+uint64_t DeviceTy::getMapEntryRefCnt(void *HstPtrBegin) {
   uintptr_t hp = (uintptr_t)HstPtrBegin;
-  long RefCnt = -1;
+  uint64_t RefCnt = 0;
 
   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) {
+  if (RefCnt == 0) {
     DP("DeviceTy::getMapEntry: requested entry not found\n");
   }
 
@@ -174,15 +170,14 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
     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" :
+        HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str());
-            std::to_string(HT.RefCount).c_str());
     rc = (void *)tp;
   } else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) {
     // Explicit extension of mapped data - not allowed.
@@ -229,17 +224,16 @@ void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
 
   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)
+    if (!IsLast && UpdateRefCount)
-      --HT.RefCount;
+      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" :
+        HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str());
-            std::to_string(HT.RefCount).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
@@ -280,9 +274,8 @@ int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
     if (ForceDelete)
-      HT.RefCount = 1;
+      HT.resetRefCount();
-    if (--HT.RefCount <= 0) {
+    if (HT.decRefCount() == 0) {
-      assert(HT.RefCount == 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);

openmp/libomptarget/src/device.h

@@ -13,8 +13,8 @@
 #ifndef _OMPTARGET_DEVICE_H
 #define _OMPTARGET_DEVICE_H
 
+#include <cassert>
 #include <cstddef>
-#include <climits>
 #include <list>
 #include <map>
 #include <mutex>
@@ -25,9 +25,6 @@ struct RTLInfoTy;
 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 {
   uintptr_t HstPtrBase; // host info.
@@ -36,18 +33,48 @@ struct HostDataToTargetTy {
 
   uintptr_t TgtPtrBegin; // target info.
 
-  long RefCount;
+private:
+  uint64_t RefCount;
+  static const uint64_t INFRefCount = ~(uint64_t)0;
 
-  HostDataToTargetTy()
+public:
-      : 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)
+      bool IsINF = false)
       : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E),
-        TgtPtrBegin(TB), RefCount(RF) {}
+        TgtPtrBegin(TB), RefCount(IsINF ? INFRefCount : 1) {}
+
+  uint64_t getRefCount() const {
+    return RefCount;
+  }
+
+  uint64_t resetRefCount() {
+    if (RefCount != INFRefCount)
+      RefCount = 1;
+
+    return RefCount;
+  }
+
+  uint64_t incRefCount() {
+    if (RefCount != INFRefCount) {
+      ++RefCount;
+      assert(RefCount < INFRefCount && "refcount overflow");
+    }
+
+    return RefCount;
+  }
+
+  uint64_t decRefCount() {
+    if (RefCount != INFRefCount) {
+      assert(RefCount > 0 && "refcount underflow");
+      --RefCount;
+    }
+
+    return RefCount;
+  }
+
+  bool isRefCountInf() const {
+    return RefCount == INFRefCount;
+  }
 };
 
 typedef std::list<HostDataToTargetTy> HostDataToTargetListTy;
@@ -129,7 +156,7 @@ struct DeviceTy {
     return *this;
   }
 
-  long getMapEntryRefCnt(void *HstPtrBegin);
+  uint64_t getMapEntryRefCnt(void *HstPtrBegin);
   LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
   void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
       bool &IsNew, bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true,

openmp/libomptarget/src/omptarget.cpp

@@ -139,7 +139,7 @@ static int InitLibrary(DeviceTy& Device) {
             (uintptr_t)CurrHostEntry->addr /*HstPtrBegin*/,
             (uintptr_t)CurrHostEntry->addr + CurrHostEntry->size /*HstPtrEnd*/,
             (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/,
-            INF_REF_CNT /*RefCount*/));
+            true /*IsRefCountINF*/));
       }
     }
     Device.DataMapMtx.unlock();
@@ -301,7 +301,7 @@ int target_data_begin(DeviceTy &Device, int32_t arg_num,
         } else if (arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) {
           // Copy data only if the "parent" struct has RefCount==1.
           int32_t parent_idx = member_of(arg_types[i]);
-          long parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
+          uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
           assert(parent_rc > 0 && "parent struct not found");
           if (parent_rc == 1) {
             copy = true;
@@ -402,7 +402,7 @@ int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base,
               !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
             // Copy data only if the "parent" struct has RefCount==1.
             int32_t parent_idx = member_of(arg_types[i]);
-            long parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
+            uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
             assert(parent_rc > 0 && "parent struct not found");
             if (parent_rc == 1) {
               CopyMember = true;

openmp/libomptarget/test/mapping/delete_inf_refcount.c

@@ -0,0 +1,32 @@
+// 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;
+}