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

[grokos]

Summary

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.

[JonChesterfield]

Yep, that seems clear cut. Thanks!

openmp/libomptarget/src/omptarget.cpp

for (int32_t i = 0; i < arg_num; ++i) {
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
       bool copy = false;
       if (!(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
           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");
           if (parent_rc == 1) {
             copy = true;
           }
         }
       }

openmp/libomptarget/test/mapping/ptr_and_obj_motion.c

+// 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]);
+}
+