[Openmp-commits] [PATCH] D89597: [libomptarget] Fix copy-to motion for PTR_AND_OBJ entries where PTR is a struct member

[George Rokos via Phabricator via Openmp-commits]

grokos created this revision.
grokos added reviewers: Hahnfeld, ABataev, jdoerfert, JonChesterfield.
grokos added a project: OpenMP.
grokos requested review of this revision.
Herald added a subscriber: sstefan1.

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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D89597

Files:
  openmp/libomptarget/src/omptarget.cpp
  openmp/libomptarget/test/mapping/ptr_and_obj_motion.c


Index: openmp/libomptarget/test/mapping/ptr_and_obj_motion.c
===================================================================
--- /dev/null
+++ 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]);
+}
+
Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -372,8 +372,11 @@
           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");


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D89597.298767.patch
Type: text/x-patch
Size: 2284 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20201016/f2df2e27/attachment.bin>