[Openmp-commits] [PATCH] D72525: [LIBOMPTARGET] Do not increment/decrement the refcount for "declare target" objects

[George Rokos via Phabricator via Openmp-commits]

grokos created this revision.
grokos added reviewers: jdoerfert, ABataev, Hahnfeld.
grokos added a project: OpenMP.

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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D72525

Files:
  openmp/libomptarget/src/device.cpp
  openmp/libomptarget/test/mapping/delete_inf_refcount.c


Index: openmp/libomptarget/test/mapping/delete_inf_refcount.c
===================================================================
--- /dev/null
+++ 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;
+}
Index: openmp/libomptarget/src/device.cpp
===================================================================
--- openmp/libomptarget/src/device.cpp
+++ openmp/libomptarget/src/device.cpp
@@ -173,7 +173,7 @@
     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);
@@ -231,7 +231,7 @@
     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);
@@ -279,9 +279,9 @@
   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);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D72525.237380.patch
Type: text/x-patch
Size: 2382 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20200110/ea2fb2e0/attachment.bin>