diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp --- a/openmp/libomptarget/src/omptarget.cpp +++ b/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"); 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 --- /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 + +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]); +} +