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 @@ -889,6 +889,7 @@ DP("Restoring original host pointer value " DPxMOD " for host pointer " DPxMOD "\n", DPxPTR(Itr->second.HstPtrVal), DPxPTR(ShadowHstPtrAddr)); + ++Itr; return OFFLOAD_SUCCESS; }; applyToShadowMapEntries(Device, CB, HstPtrBegin, ArgSize, TPR); @@ -911,6 +912,7 @@ sizeof(void *), AsyncInfo); if (Ret != OFFLOAD_SUCCESS) REPORT("Copying data to device failed.\n"); + ++Itr; return Ret; }; applyToShadowMapEntries(Device, CB, HstPtrBegin, ArgSize, TPR); diff --git a/openmp/libomptarget/test/offloading/bug53727.cpp b/openmp/libomptarget/test/offloading/bug53727.cpp new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/offloading/bug53727.cpp @@ -0,0 +1,57 @@ +// RUN: %libomptarget-compilexx-and-run-generic + +#include +#include + +constexpr const int N = 10; + +struct T { + int a; + int *p; +}; + +struct S { + int b; + T t; +}; + +int main(int argc, char *argv[]) { + S s; + s.t.p = new int[N]; + for (int i = 0; i < N; ++i) { + s.t.p[i] = i; + } + +#pragma omp target enter data map(to : s, s.t.p[:N]) + +#pragma omp target + { + for (int i = 0; i < N; ++i) { + s.t.p[i] += i; + } + } + +#pragma omp target update from(s.t.p[:N]) + + for (int i = 0; i < N; ++i) { + assert(s.t.p[i] == 2 * i); + s.t.p[i] += i; + } + +#pragma omp target update to(s.t.p[:N]) + +#pragma omp target + { + for (int i = 0; i < N; ++i) { + s.t.p[i] += i; + } + } + +#pragma omp target exit data map(from : s, s.t.p[:N]) + + for (int i = 0; i < N; ++i) { + assert(s.t.p[i] == 4 * i); + } + + return 0; +}