[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 ↗

(On Diff #237562)

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

59 ↗

(On Diff #237562)

Maybe just use a bitfield/bool flag instead?

[ABataev]

openmp/libomptarget/src/device.h

54 ↗

(On Diff #237562)

Const function

77 ↗

(On Diff #237562)

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

30 ↗

(On Diff #237929)

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.cpp

void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
   // 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;
     IsNew = false;
 
-    if (UpdateRefCount)
+    if (!CONSIDERED_INF(HT.RefCount) && UpdateRefCount)
       ++HT.RefCount;
 
     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" :
void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
   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);
 
-    if (HT.RefCount > 1 && UpdateRefCount)
+    if (!CONSIDERED_INF(HT.RefCount) && HT.RefCount > 1 && UpdateRefCount)
       --HT.RefCount;
 
     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());
int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
   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)
+    if (ForceDelete && !CONSIDERED_INF(HT.RefCount))
       HT.RefCount = 1;
-    if (--HT.RefCount <= 0) {
+    if (!CONSIDERED_INF(HT.RefCount) && --HT.RefCount <= 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);
       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;
+}