[Clang][OpenMP] Fix run time crash when use_device_addr is used.

It is data mapping ordering problem.

According omp spec
If one or more map clauses are present, the list item conversions that
are performed for any use_device_ptr or use_device_addr clause occur
after all variables are mapped on entry to the region according to those
map clauses.

The change is to put mapping data for use_device_addr at end of data
mapping array.

Differential Revision: https://reviews.llvm.org/D134556

1 files changed, 18 insertions, 0 deletions

diff --git a/openmp/libomptarget/test/mapping/target_use_device_addr.c b/openmp/libomptarget/test/mapping/target_use_device_addr.c
new file mode 100644
index 0000000..05a5aea
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/target_use_device_addr.c
@@ -0,0 +1,18 @@
+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+int main() {
+  short x[10];
+  short *xp = &x[0];
+
+  x[1] = 111;
+
+  printf("%d, %p\n", xp[1], &xp[1]);
+#pragma omp target data use_device_addr(xp [1:3]) map(tofrom : x)
+#pragma omp target is_device_ptr(xp)
+  { xp[1] = 222; }
+  // CHECK: 222
+  printf("%d, %p\n", xp[1], &xp[1]);
+}