[OPENMP50]Allow overlapping mapping in target constructs.
OpenMP 5.0 removed a lot of restriction for overlapped mapped items
comparing to OpenMP 4.5. Patch restricts the checks for overlapped data
mappings only for OpenMP 4.5 and less and reorders mapping of the
arguments so, that present and alloc mappings are processed first and
then all others.
Differential Revision: https://reviews.llvm.org/D86119
GitOrigin-RevId: 60d71a286b5a03653fc99cd09423d603feb897de
diff --git a/libomptarget/src/omptarget.cpp b/libomptarget/src/omptarget.cpp
index ddefdd8..d2d0e08 100644
--- a/libomptarget/src/omptarget.cpp
+++ b/libomptarget/src/omptarget.cpp
@@ -258,7 +258,7 @@
for (unsigned I = 0, E = MapperComponents.Components.size(); I < E; ++I) {
auto &C =
MapperComponents
- .Components[target_data_function == targetDataEnd ? I : E - I - 1];
+ .Components[target_data_function == targetDataEnd ? E - I - 1 : I];
MapperArgsBase[I] = C.Base;
MapperArgs[I] = C.Begin;
MapperArgSizes[I] = C.Size;
diff --git a/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp b/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
new file mode 100644
index 0000000..87697af
--- /dev/null
+++ b/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
@@ -0,0 +1,66 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+ int a;
+ double *b;
+} C;
+#pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+ int e;
+ C f;
+ int h;
+ short *g;
+} D;
+#pragma omp declare mapper(default \
+ : D r) map(from \
+ : r.e) map(mapper(id1), tofrom \
+ : r.f) map(tofrom \
+ : r.g [0:r.h])
+
+int main() {
+ constexpr int N = 10;
+ D s;
+ s.e = 111;
+ s.f.a = 222;
+ double x[2];
+ x[1] = 20;
+ short y[N];
+ y[1] = 30;
+ s.f.b = &x[0];
+ s.g = &y[0];
+ s.h = N;
+
+ D *sp = &s;
+ D **spp = &sp;
+
+ printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
+ spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
+ spp[0][0].g == &y[0] ? 1 : 0);
+ // CHECK: 111 222 20.00000 1 30 1
+
+ __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]),
+ p1 = reinterpret_cast<__intptr_t>(&y[0]);
+#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1)
+ {
+ printf("%d %d %d %d\n", spp[0][0].f.a,
+ spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0, spp[0][0].g[1],
+ spp[0][0].g == reinterpret_cast<void *>(p1) ? 1 : 0);
+ // CHECK: 222 0 30 0
+ spp[0][0].e = 333;
+ spp[0][0].f.a = 444;
+ spp[0][0].f.b[1] = 40;
+ spp[0][0].g[1] = 50;
+ }
+ printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
+ spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
+ spp[0][0].g == &y[0] ? 1 : 0);
+ // CHECK: 333 222 40.00000 1 50 1
+}
