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

This revision was automatically updated to reflect the committed changes.
Closed by commit rG5adb3a6d86ee: [libomptarget] Fix copy-to motion for PTR_AND_OBJ entries where PTR is a struct… (authored by grokos).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D89597/new/

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.298780.patch
Type: text/x-patch
Size: 2284 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20201016/0b3eed24/attachment.bin>