diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp index 14f5e7dc9d19..9b3533895f2a 100644 --- a/offload/libomptarget/OpenMP/Mapping.cpp +++ b/offload/libomptarget/OpenMP/Mapping.cpp @@ -326,6 +326,28 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer( // data transfer. if (LR.TPR.TargetPointer && !LR.TPR.Flags.IsHostPointer && HasFlagTo && (LR.TPR.Flags.IsNewEntry || HasFlagAlways) && Size != 0) { + + // If we have something like: + // #pragma omp target map(to: s.myarr[0:10]) map(to: s.myarr[0:10]) + // then we see two "new" mappings of the struct member s.myarr here -- + // and both have the "IsNewEntry" flag set, so trigger the copy to device + // below. But, the shadow pointer is only initialised on the target for + // the first copy, and the second copy clobbers it. So, this condition + // avoids the (second) copy here if we have already set shadow pointer info. + auto FailOnPtrFound = [HstPtrBegin, Size](ShadowPtrInfoTy &SP) { + if (SP.HstPtrAddr >= HstPtrBegin && + SP.HstPtrAddr < (void *)((char *)HstPtrBegin + Size)) + return OFFLOAD_FAIL; + return OFFLOAD_SUCCESS; + }; + if (LR.TPR.getEntry()->foreachShadowPointerInfo(FailOnPtrFound) == + OFFLOAD_FAIL) { + DP("Multiple new mappings of %" PRId64 " bytes detected (hst:" DPxMOD + ") -> (tgt:" DPxMOD ")\n", + Size, DPxPTR(HstPtrBegin), DPxPTR(LR.TPR.TargetPointer)); + return std::move(LR.TPR); + } + DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n", Size, DPxPTR(HstPtrBegin), DPxPTR(LR.TPR.TargetPointer)); diff --git a/offload/test/mapping/duplicate_mappings_1.cpp b/offload/test/mapping/duplicate_mappings_1.cpp new file mode 100644 index 000000000000..5bed3f161b4f --- /dev/null +++ b/offload/test/mapping/duplicate_mappings_1.cpp @@ -0,0 +1,27 @@ +// clang-format off +// RUN: %libomptarget-compilexx-generic -Wno-openmp-mapping && %libomptarget-run-generic + +// clang-format on + +#include + +struct Inner { + int *data; + Inner(int size) { data = new int[size](); } + ~Inner() { delete[] data; } +}; + +struct Outer { + Inner i; + Outer() : i(10) {} +}; + +int main() { + Outer o; +#pragma omp target map(tofrom : o.i.data[0 : 10]) map(tofrom : o.i.data[0 : 10]) + { + o.i.data[0] = 42; + } + assert(o.i.data[0] == 42); + return 0; +} diff --git a/offload/test/mapping/duplicate_mappings_2.cpp b/offload/test/mapping/duplicate_mappings_2.cpp new file mode 100644 index 000000000000..ca8112ee7254 --- /dev/null +++ b/offload/test/mapping/duplicate_mappings_2.cpp @@ -0,0 +1,29 @@ +// clang-format off +// RUN: %libomptarget-compilexx-generic -Wno-openmp-mapping && %libomptarget-run-generic + +#include + +// clang-format on + +struct Inner { + int *data; + Inner(int size) { data = new int[size](); } + ~Inner() { delete[] data; } +}; +#pragma omp declare mapper(Inner i) map(i, i.data[0 : 10]) + +struct Outer { + Inner i; + Outer() : i(10) {} +}; +#pragma omp declare mapper(Outer o) map(o, o.i) + +int main() { + Outer o; +#pragma omp target map(tofrom : o) + { + o.i.data[0] = 42; + } + assert(o.i.data[0] == 42); + return 0; +}