[Openmp-commits] [openmp] 5adb3a6 - [libomptarget] Fix copy-to motion for PTR_AND_OBJ entries where PTR is a struct member.

[George Rokos via Openmp-commits]

Author: George Rokos
Date: 2020-10-16T16:14:01-07:00
New Revision: 5adb3a6d86eecade2cb94b1a04d35e673d4e5866

URL: https://github.com/llvm/llvm-project/commit/5adb3a6d86eecade2cb94b1a04d35e673d4e5866
DIFF: https://github.com/llvm/llvm-project/commit/5adb3a6d86eecade2cb94b1a04d35e673d4e5866.diff

LOG: [libomptarget] Fix copy-to motion for PTR_AND_OBJ entries where PTR is a struct member.

This patch fixes a problem whereby the pointee object of a PTR_AND_OBJ entry with a `map(to)` motion clause can be overwritten on the device even if its reference counter is >=1.

Currently, we check the reference counter of the parent struct in order to determine whether the motion clause should be respected, but since the pointee object is not part of the struct, it's got its own reference counter which should be used to enqueue the copy or discard it.

The same behavior has already been implemented in targetDataEnd (omptarget.cpp:539-540), but we somehow missed doing the same in targetDataBegin.

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

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

Modified: 
    openmp/libomptarget/src/omptarget.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index e44b9aad6202..d6fb4b3a94d9 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -372,8 +372,11 @@ int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
           HasCloseModifier) {
         if (IsNew || (arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS)) {
           copy = true;
-        } else if (arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) {
+        } else if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
+                   !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
           // Copy data only if the "parent" struct has RefCount==1.
+          // If this is a PTR_AND_OBJ entry, the OBJ is not part of the struct,
+          // so exclude it from this check.
           int32_t parent_idx = getParentIndex(arg_types[i]);
           uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
           assert(parent_rc > 0 && "parent struct not found");

diff  --git a/openmp/libomptarget/test/mapping/ptr_and_obj_motion.c b/openmp/libomptarget/test/mapping/ptr_and_obj_motion.c
new file mode 100644
index 000000000000..b08734542662
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/ptr_and_obj_motion.c
@@ -0,0 +1,48 @@
+// 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
+// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+
+#include <stdio.h>
+
+typedef struct {
+  double *dataptr;
+  int dummy1;
+  int dummy2;
+} DV;
+
+void init(double vertexx[]) {
+  #pragma omp target map(vertexx[0:100])
+  {
+    printf("In init: %lf, expected 100.0\n", vertexx[77]);
+    vertexx[77] = 77.0;
+  }
+}
+
+void change(DV *dvptr) {
+  #pragma omp target map(dvptr->dataptr[0:100])
+  {
+    printf("In change: %lf, expected 77.0\n", dvptr->dataptr[77]);
+    dvptr->dataptr[77] += 1.0;
+  }
+}
+
+int main() {
+  double vertexx[100];
+  vertexx[77] = 100.0;
+
+  DV dv;
+  dv.dataptr = &vertexx[0];
+
+  #pragma omp target enter data map(to:vertexx[0:100])
+
+  init(vertexx);
+  change(&dv);
+
+  #pragma omp target exit data map(from:vertexx[0:100])
+
+  // CHECK: Final: 78.0
+  printf("Final: %lf\n", vertexx[77]);
+}
+