OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.

This patch adds the OpenMP runtime routine "omp_get_mapped_ptr" which was
introduced in OpenMP 5.1.

gcc/ChangeLog:

	* omp-low.cc (omp_runtime_api_call): Added get_mapped_ptr to
	omp_runtime_apis array.

libgomp/ChangeLog:

	* libgomp.map: Added omp_get_mapped_ptr.
	* libgomp.texi: Tagged omp_get_mapped_ptr as supported.
	* omp.h.in: Added omp_get_mapped_ptr.
	* omp_lib.f90.in: Added interface for omp_get_mapped_ptr.
	* omp_lib.h.in: Likewise.
	* target.c (omp_get_mapped_ptr): Added implementation of
	omp_get_mapped_ptr.
	* testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: New test.
	* testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c: New test.
	* testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c: New test.
	* testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c: New test.
	* testsuite/libgomp.fortran/get-mapped-ptr-1.f90: New test.
	* testsuite/libgomp.fortran/get-mapped-ptr-2.f90: New test.
	* testsuite/libgomp.fortran/get-mapped-ptr-3.f90: New test.
	* testsuite/libgomp.fortran/get-mapped-ptr-4.f90: New test.

4 files changed, 247 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c
new file mode 100644
index 0000000..97a60ca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c
@@ -0,0 +1,41 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  void *p , *q;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  p = omp_target_alloc (sizeof (int), d);
+  if (p == NULL)
+    return 0;
+
+  if (omp_target_associate_ptr (q, p, sizeof (int), 0, d) != 0)
+    return 0;
+
+  if (omp_get_mapped_ptr (q, -1) != NULL)
+    abort ();
+
+  if (omp_get_mapped_ptr (q, omp_get_num_devices () + 1) != NULL)
+    abort ();
+
+  if (omp_get_mapped_ptr (q, id) != q)
+    abort ();
+
+  if (omp_get_mapped_ptr (q, d) != p)
+    abort ();
+
+  if (omp_target_disassociate_ptr (q, d) != 0)
+    abort ();
+
+  if (omp_get_mapped_ptr (q, d) != NULL)
+    abort ();
+
+  omp_target_free (p, d);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c
new file mode 100644
index 0000000..194dade
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c
@@ -0,0 +1,106 @@
+#include <omp.h>
+#include <stdlib.h>
+#include <stdint.h>
+
+int
+main ()
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int a = 42;
+  int b[] = { 24, 42 };
+  int c[] = { 47, 11 };
+  int e[128];
+  int *q = &a;
+  void *p1 = NULL, *p2 = NULL, *p3 = NULL;
+  void *devptrs[128];
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  for (int i = 0; i < 128; i++)
+    e[i] = i;
+
+  #pragma omp target data map(alloc: a, b, c[1], e[32:64]) device(d)
+  {
+    #pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d)
+    {
+      p1 = &a;
+      p2 = &b;
+      p3 = &c[1];
+      for (int i = 32; i < 96; i++)
+	devptrs[i] = &e[i];
+    }
+
+    if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1)
+	|| omp_get_mapped_ptr (q, d) != (d == id ? q : p1)
+	|| omp_get_mapped_ptr (b, d) != (d == id ? b : p2)
+	|| omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2)
+	|| omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3)
+	|| omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
+      abort ();
+
+    for (int i = 0; i < 32; i++)
+      if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+	abort ();
+    for (int i = 32; i < 96; i++)
+      if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i]))
+	abort ();
+    for (int i = 96; i < 128; i++)
+      if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+	abort ();
+  }
+
+  if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL)
+      || omp_get_mapped_ptr (q, d) != (d == id ? q : NULL)
+      || omp_get_mapped_ptr (b, d) != (d == id ? b : NULL)
+      || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL)
+      || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL)
+      || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
+    abort ();
+  for (int i = 0; i < 128; i++)
+    if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+      abort ();
+
+  #pragma omp target enter data map (alloc: a, b, c[1], e[32:64]) device (d)
+  #pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d)
+  {
+    p1 = &a;
+    p2 = &b;
+    p3 = &c[1];
+    for (int i = 32; i < 96; i++)
+      devptrs[i] = &e[i];
+  }
+
+  if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1)
+      || omp_get_mapped_ptr (q, d) != (d == id ? q : p1)
+      || omp_get_mapped_ptr (b, d) != (d == id ? b : p2)
+      || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2)
+      || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3)
+      || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
+    abort ();
+  for (int i = 0; i < 32; i++)
+    if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+      abort ();
+  for (int i = 32; i < 96; i++)
+    if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i]))
+      abort ();
+  for (int i = 96; i < 128; i++)
+    if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+      abort ();
+
+  #pragma omp target exit data map (delete: a, b, c[1], e[32:64]) device (d)
+
+  if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL)
+      || omp_get_mapped_ptr (q, d) != (d == id ? q : NULL)
+      || omp_get_mapped_ptr (b, d) != (d == id ? b : NULL)
+      || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL)
+      || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL)
+      || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
+    abort ();
+  for (int i = 0; i < 128; i++)
+    if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+      abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c
new file mode 100644
index 0000000..747ef75
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c
@@ -0,0 +1,51 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  int a[0];
+  int b[] = { 24, 42 };
+  void *p1 = NULL, *p2 = NULL;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  void *p = omp_target_alloc (sizeof (int), d);
+  if (p == NULL)
+    return 0;
+
+  if (omp_target_associate_ptr (a, p, sizeof (int), 0, d) != 0)
+    return 0;
+
+  if (omp_get_mapped_ptr (a, d) != (d == id ? a : p))
+    abort ();
+
+  if (omp_target_disassociate_ptr (a, d) != 0)
+    abort ();
+
+  if (omp_get_mapped_ptr (a, d) != (d == id ? a : NULL))
+    abort ();
+
+  #pragma omp target data map(alloc: a, b[1:0]) device(d)
+  {
+    #pragma omp target map(from: p1, p2) map(alloc: a, b[1:0]) device(d)
+    {
+      p1 = &a;
+      p2 = &b[1];
+    }
+
+    /* This is probably expected to be p1/p2 instead of NULL. Zero-length arrays
+       as list items of the map clause are currently not inserted into the mem
+       map ?! However by returning NULL, omp_get_mapped_ptr is consistent with
+       omp_target_is_present.  */
+    if (omp_get_mapped_ptr (a, d) != NULL
+	|| omp_get_mapped_ptr (&b[1], d) != NULL)
+      abort ();
+  }
+
+  omp_target_free (p, d);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c
new file mode 100644
index 0000000..6f4bd62
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c
@@ -0,0 +1,49 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int d = omp_get_default_device ();
+  int id = omp_get_initial_device ();
+  struct s_t { int m1; char m2; } s;
+  void *p1 = NULL, *p2 = NULL;
+
+  if (d < 0 || d >= omp_get_num_devices ())
+    d = id;
+
+  #pragma omp target data map(alloc: s, s.m2) device(d)
+  {
+    #pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d)
+    {
+      p1 = &s;
+      p2 = &s.m2;
+    }
+    if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1)
+	|| omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2))
+      abort ();
+  }
+
+  if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL)
+      || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL))
+    abort ();
+
+  #pragma omp target enter data map(alloc: s, s.m2) device (d)
+  #pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d)
+  {
+    p1 = &s;
+    p2 = &s.m2;
+  }
+
+  if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1)
+      || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2))
+    abort ();
+
+  #pragma omp target exit data map (delete: s, s.m2) device (d)
+
+  if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL)
+      || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL))
+    abort ();
+
+  return 0;
+}