[Openmp-commits] [openmp] 518a27e - [OpenMP] Fix ref count dec for implicit map of partial data

[Joel E. Denny via Openmp-commits]

Author: Joel E. Denny
Date: 2020-08-06T11:39:29-04:00
New Revision: 518a27e5591c211ceeef3091edc59012e6ace2b2

URL: https://github.com/llvm/llvm-project/commit/518a27e5591c211ceeef3091edc59012e6ace2b2
DIFF: https://github.com/llvm/llvm-project/commit/518a27e5591c211ceeef3091edc59012e6ace2b2.diff

LOG: [OpenMP] Fix ref count dec for implicit map of partial data

D85342 broke this case.  The new test case presents an example.

Reviewed By: grokos

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

Added: 
    openmp/libomptarget/test/mapping/target_implicit_partial_map.c

Modified: 
    openmp/libomptarget/src/omptarget.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 2b158ff34268..aa4ab75573b4 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -497,6 +497,7 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
     }
 
     bool IsLast, IsHostPtr;
+    bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;
     bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
                      (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ);
     bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
@@ -504,9 +505,8 @@ int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
     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,
-                                              /*MustContain=*/true);
+    void *TgtPtrBegin = Device.getTgtPtrBegin(
+        HstPtrBegin, DataSize, IsLast, UpdateRef, IsHostPtr, !IsImplicit);
     if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
       DP("Mapping does not exist (%s)\n",
          (HasPresentModifier ? "'present' map type modifier" : "ignored"));

diff  --git a/openmp/libomptarget/test/mapping/target_implicit_partial_map.c b/openmp/libomptarget/test/mapping/target_implicit_partial_map.c
new file mode 100644
index 000000000000..aa713ff0f8a0
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/target_implicit_partial_map.c
@@ -0,0 +1,39 @@
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+//
+// END.
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int arr[100];
+
+#pragma omp target data map(alloc: arr[50:2]) // partially mapped
+  {
+#pragma omp target // would implicitly map with full size but already present
+    {
+      arr[50] = 5;
+      arr[51] = 6;
+    } // must treat as present (dec ref count) even though full size not present
+  } // wouldn't delete if previous ref count dec didn't happen
+
+  // CHECK: still present: 0
+  fprintf(stderr, "still present: %d\n",
+          omp_target_is_present(&arr[50], omp_get_default_device()));
+
+  return 0;
+}