OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.

Hi,

This patch adds the OpenMP runtime routine "omp_get_mapped_ptr" which was
introduced in OpenMP 5.1 (specification section 3.8.11):

"The omp_get_mapped_ptr routine returns the device pointer that is associated
with a host pointer for a given device."

"The device_num argument must be greater than or equal to zero and less than or
equal to the result of omp_get_num_devices()."

"A call to this routine for a pointer that is not NULL (or C_NULL_PTR, for
Fortran) and does not have an associated pointer on the given device results in
a NULL pointer."

"The routine returns NULL (or C_NULL_PTR, for Fortran) if unsuccessful.
Otherwise it returns the device pointer, which is ptr if device_num is the value
returned by omp_get_initial_device()."

Implementation and tests were added for C/C++ and Fortran.

There is a small inconvenience considering zero-length arrays as list items of
the "target map" construct: it seems that zero-length arrays are not associated
correctly there, such that omp_get_mapped_ptr returns NULL instead of the
associated device pointer - in contrast to the situation where a device pointer
is associated with the host pointer via omp_target_associate_ptr.
However, the result for omp_get_mapped_ptr is consistent with
omp_target_is_present (which returns 0, i.e. "not present") in this situation.

The patch was tested on x86_64-linux with nvptx and amdgcn offloading. All with
no regressions.

Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.

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.

On Fri, Mar 04, 2022 at 03:47:31PM +0100, Marcel Vollweiler wrote:
> 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.
> 
> diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
> index 2ac5809..00a4858 100644
> --- a/libgomp/libgomp.map
> +++ b/libgomp/libgomp.map
> @@ -224,6 +224,7 @@ OMP_5.1 {
>  	omp_set_teams_thread_limit_8_;
>  	omp_get_teams_thread_limit;
>  	omp_get_teams_thread_limit_;
> +	omp_get_mapped_ptr;
>  } OMP_5.0.2;

I think it is too late for this to be targetted for GCC 12, and
for GCC 13 it will need to go into OMP_5.1.1 symver.

>  GOMP_1.0 {
> diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
> index 161a423..c163b56 100644
> --- a/libgomp/libgomp.texi
> +++ b/libgomp/libgomp.texi
> @@ -314,7 +314,7 @@ The OpenMP 4.5 specification is fully supported.
>  @item @code{omp_target_is_accessible} runtime routine @tab N @tab
>  @item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async}
>        runtime routines @tab N @tab
> -@item @code{omp_get_mapped_ptr} runtime routine @tab N @tab
> +@item @code{omp_get_mapped_ptr} runtime routine @tab Y @tab
>  @item @code{omp_calloc}, @code{omp_realloc}, @code{omp_aligned_alloc} and
>        @code{omp_aligned_calloc} runtime routines @tab Y @tab
>  @item @code{omp_alloctrait_key_t} enum: @code{omp_atv_serialized} added,
> diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
> index 89c5d65..18d0152 100644
> --- a/libgomp/omp.h.in
> +++ b/libgomp/omp.h.in
> @@ -282,6 +282,7 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int,
>  extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__,
>  				     __SIZE_TYPE__, int) __GOMP_NOTHROW;
>  extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW;
> +extern void *omp_get_mapped_ptr (const void *, int) __GOMP_NOTHROW;
>  
>  extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW;
>  extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__)
> diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
> index daf40dc..506f15c 100644
> --- a/libgomp/omp_lib.f90.in
> +++ b/libgomp/omp_lib.f90.in
> @@ -835,6 +835,15 @@
>            end function omp_target_disassociate_ptr
>          end interface
>  
> +        interface
> +          function omp_get_mapped_ptr (ptr, device_num) bind(c)
> +            use, intrinsic :: iso_c_binding, only : c_ptr, c_int
> +            type(c_ptr) :: omp_get_mapped_ptr
> +            type(c_ptr), value :: ptr
> +            integer(c_int), value :: device_num
> +          end function omp_get_mapped_ptr
> +        end interface
> +
>  #if _OPENMP >= 201811
>  !GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested
>  #endif
> diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
> index ff857a4..0f48510 100644
> --- a/libgomp/omp_lib.h.in
> +++ b/libgomp/omp_lib.h.in
> @@ -416,3 +416,12 @@
>            integer(c_int), value :: device_num
>          end function omp_target_disassociate_ptr
>        end interface
> +
> +      interface
> +        function omp_get_mapped_ptr (ptr, device_num) bind(c)
> +          use, intrinsic :: iso_c_binding, only : c_ptr, c_int
> +          type(c_ptr) :: omp_get_mapped_ptr
> +          type(c_ptr), value :: ptr
> +          integer(c_int), value :: device_num
> +        end function omp_get_mapped_ptr
> +      end interface
> diff --git a/libgomp/target.c b/libgomp/target.c
> index 9017458..735d70b 100644
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -3665,6 +3665,49 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
>    return ret;
>  }
>  
> +void *
> +omp_get_mapped_ptr (const void *ptr, int device_num)
> +{
> +  if (device_num < 0 || device_num > omp_get_num_devices ())
> +    return NULL;
> +
> +  if (device_num == omp_get_initial_device ())
> +    return (void*)ptr;

Space before * and space after )

> +  struct gomp_device_descr *devicep = resolve_device (device_num);
> +  if (devicep == NULL)
> +    return NULL;
> +
> +  if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
> +      || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
> +    return (void*)ptr;

Likewise.
> +
> +  gomp_mutex_lock (&devicep->lock);
> +
> +  struct splay_tree_s *mem_map = &devicep->mem_map;
> +  struct splay_tree_key_s cur_node;
> +  void *ret = NULL;
> +  uintptr_t offset = 0;

offset should be moved to the only place that defines it.
> +
> +  cur_node.host_start = (uintptr_t) ptr;
> +  cur_node.host_end = cur_node.host_start;
> +  splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
> +
> +  if (n && n->host_start == cur_node.host_start)
> +    {
> +      ret = (void*) n->tgt->tgt_start + n->tgt_offset;
> +    }

Single statement body, so without {}s and reindented, space before *.
> +  else if (n)
> +    {
> +      offset = cur_node.host_start - n->host_start;
      uintptr_t offset = cur_node.host_start - n->host_start;

> +      ret = (void*) n->tgt->tgt_start + n->tgt_offset + offset;

Space before *.

Though, looking at this more, what is the point of the first if?
The second if would compute offset = 0...

Also, void * arithmetics is a GNU extension, maybe better use char *.

> +  if (omp_get_mapped_ptr (q, -1) != NULL)
> +    __builtin_abort ();

When you do include stdlib.h, what is the point of using __builtin_abort ?
Just use abort then.

	Jakub

Hi Jakub,

>> diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
>> index 2ac5809..00a4858 100644
>> --- a/libgomp/libgomp.map
>> +++ b/libgomp/libgomp.map
>> @@ -224,6 +224,7 @@ OMP_5.1 {
>>      omp_set_teams_thread_limit_8_;
>>      omp_get_teams_thread_limit;
>>      omp_get_teams_thread_limit_;
>> +    omp_get_mapped_ptr;
>>   } OMP_5.0.2;
>
> I think it is too late for this to be targetted for GCC 12, and
> for GCC 13 it will need to go into OMP_5.1.1 symver.

Agreed and changed accordingly.

>> +void *
>> +omp_get_mapped_ptr (const void *ptr, int device_num)
>> +{
>> +  if (device_num < 0 || device_num > omp_get_num_devices ())
>> +    return NULL;
>> +
>> +  if (device_num == omp_get_initial_device ())
>> +    return (void*)ptr;
>
> Space before * and space after )

Changed.

>> +  struct gomp_device_descr *devicep = resolve_device (device_num);
>> +  if (devicep == NULL)
>> +    return NULL;
>> +
>> +  if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
>> +      || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
>> +    return (void*)ptr;
>
> Likewise.

Changed.

>> +
>> +  gomp_mutex_lock (&devicep->lock);
>> +
>> +  struct splay_tree_s *mem_map = &devicep->mem_map;
>> +  struct splay_tree_key_s cur_node;
>> +  void *ret = NULL;
>> +  uintptr_t offset = 0;
>
> offset should be moved to the only place that defines it.

Changed.

>> +
>> +  cur_node.host_start = (uintptr_t) ptr;
>> +  cur_node.host_end = cur_node.host_start;
>> +  splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
>> +
>> +  if (n && n->host_start == cur_node.host_start)
>> +    {
>> +      ret = (void*) n->tgt->tgt_start + n->tgt_offset;
>> +    }
>
> Single statement body, so without {}s and reindented, space before *.
>> +  else if (n)
>> +    {
>> +      offset = cur_node.host_start - n->host_start;
>        uintptr_t offset = cur_node.host_start - n->host_start;
>
>> +      ret = (void*) n->tgt->tgt_start + n->tgt_offset + offset;
>
> Space before *.
>
> Though, looking at this more, what is the point of the first if?
> The second if would compute offset = 0...

Absolutely true :)
Changed.

>
> Also, void * arithmetics is a GNU extension, maybe better use char *.

I changed it to (enclosing parentheses):

        ret = (void *) (n->tgt->tgt_start + n->tgt_offset + offset);

i.e. pointer arithmetic is done on uintptr_t, but I am not completely sure if
that's sufficient in terms of compatibility. On the other hand,

        ret = (void *) ((char *) n->tgt->tgt_start + (char *) n->tgt_offset + (char *)
offset);

is perhaps overcomplicated if not really necessary. What do you think?

>> +  if (omp_get_mapped_ptr (q, -1) != NULL)
>> +    __builtin_abort ();
>
> When you do include stdlib.h, what is the point of using __builtin_abort ?
> Just use abort then.

Good point. Changed.

Thanks,

Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.

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.

diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 2ac5809..608a54c 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -226,6 +226,11 @@ OMP_5.1 {
 	omp_get_teams_thread_limit_;
 } OMP_5.0.2;
 
+OMP_5.1.1 {
+  global:
+	omp_get_mapped_ptr;
+} OMP_5.1;
+
 GOMP_1.0 {
   global:
 	GOMP_atomic_end;
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 161a423..c163b56 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -314,7 +314,7 @@ The OpenMP 4.5 specification is fully supported.
 @item @code{omp_target_is_accessible} runtime routine @tab N @tab
 @item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async}
       runtime routines @tab N @tab
-@item @code{omp_get_mapped_ptr} runtime routine @tab N @tab
+@item @code{omp_get_mapped_ptr} runtime routine @tab Y @tab
 @item @code{omp_calloc}, @code{omp_realloc}, @code{omp_aligned_alloc} and
       @code{omp_aligned_calloc} runtime routines @tab Y @tab
 @item @code{omp_alloctrait_key_t} enum: @code{omp_atv_serialized} added,
diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index 89c5d65..18d0152 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -282,6 +282,7 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int,
 extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__,
 				     __SIZE_TYPE__, int) __GOMP_NOTHROW;
 extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW;
+extern void *omp_get_mapped_ptr (const void *, int) __GOMP_NOTHROW;
 
 extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW;
 extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__)
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index daf40dc..506f15c 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -835,6 +835,15 @@
           end function omp_target_disassociate_ptr
         end interface
 
+        interface
+          function omp_get_mapped_ptr (ptr, device_num) bind(c)
+            use, intrinsic :: iso_c_binding, only : c_ptr, c_int
+            type(c_ptr) :: omp_get_mapped_ptr
+            type(c_ptr), value :: ptr
+            integer(c_int), value :: device_num
+          end function omp_get_mapped_ptr
+        end interface
+
 #if _OPENMP >= 201811
 !GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested
 #endif
diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index ff857a4..0f48510 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -416,3 +416,12 @@
           integer(c_int), value :: device_num
         end function omp_target_disassociate_ptr
       end interface
+
+      interface
+        function omp_get_mapped_ptr (ptr, device_num) bind(c)
+          use, intrinsic :: iso_c_binding, only : c_ptr, c_int
+          type(c_ptr) :: omp_get_mapped_ptr
+          type(c_ptr), value :: ptr
+          integer(c_int), value :: device_num
+        end function omp_get_mapped_ptr
+      end interface
diff --git a/libgomp/target.c b/libgomp/target.c
index 9017458..89c837b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -3665,6 +3665,44 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
   return ret;
 }
 
+void *
+omp_get_mapped_ptr (const void *ptr, int device_num)
+{
+  if (device_num < 0 || device_num > omp_get_num_devices ())
+    return NULL;
+
+  if (device_num == omp_get_initial_device ())
+    return (void *) ptr;
+
+  struct gomp_device_descr *devicep = resolve_device (device_num);
+  if (devicep == NULL)
+    return NULL;
+
+  if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
+      || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+    return (void *) ptr;
+
+  gomp_mutex_lock (&devicep->lock);
+
+  struct splay_tree_s *mem_map = &devicep->mem_map;
+  struct splay_tree_key_s cur_node;
+  void *ret = NULL;
+
+  cur_node.host_start = (uintptr_t) ptr;
+  cur_node.host_end = cur_node.host_start;
+  splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
+
+  if (n)
+    {
+      uintptr_t offset = cur_node.host_start - n->host_start;
+      ret = (void *) (n->tgt->tgt_start + n->tgt_offset + offset);
+    }
+
+  gomp_mutex_unlock (&devicep->lock);
+
+  return ret;
+}
+
 int
 omp_pause_resource (omp_pause_resource_t kind, int device_num)
 {
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;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90
new file mode 100644
index 0000000..de05179
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90
@@ -0,0 +1,43 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id
+  type(c_ptr) :: p
+  integer, target :: q
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+
+  if (d < 0 .or. d >= omp_get_num_devices ()) &
+    d = id
+
+  p = omp_target_alloc (c_sizeof (q), d)
+  if (.not. c_associated (p)) &
+    stop 0  ! okay
+
+  if (omp_target_associate_ptr (c_loc (q), p, c_sizeof (q), &
+                                0_c_size_t, d) == 0) then
+
+    if(c_associated (omp_get_mapped_ptr (c_loc (q), -1))) &
+      stop 1
+
+    if(c_associated (omp_get_mapped_ptr (c_loc (q), &
+                     omp_get_num_devices () + 1))) &
+      stop 2
+
+    if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), id), c_loc (q))) &
+      stop 3
+
+    if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), d), p)) &
+      stop 4
+
+    if (omp_target_disassociate_ptr (c_loc (q), d) /= 0) &
+      stop 5
+
+    if(c_associated (omp_get_mapped_ptr (c_loc (q), d))) &
+      stop 6
+  end if
+
+  call omp_target_free (p, d)
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90
new file mode 100644
index 0000000..66a0b88
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90
@@ -0,0 +1,175 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id, i, j
+  integer, target :: a, b(1:2), c(1:2), e(0:127)
+  type(c_ptr) :: p1, p2, p3, q, devptrs(0:63)
+
+  a = 42;
+  q = c_loc (a);
+  e = [(i, i = 0, 127)]
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+
+  if (d < 0 .or. d >= omp_get_num_devices ()) &
+    d = id
+
+  if (d /= id) then
+    !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
+      !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+      p1 = c_loc (a);
+      p2 = c_loc (b);
+      p3 = c_loc (c(2))
+      devptrs = [(c_loc (e(i)), i = 32, 95)]
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), p2) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
+          .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+        stop 0
+
+      do j = 0, 31
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 1
+      end do
+      do j = 32, 95
+        if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
+          stop 2
+      end do
+      do j = 96, 128
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 3
+      end do
+    !$omp end target data
+
+    if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
+        .or. c_associated (omp_get_mapped_ptr (q, d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+      stop 4
+      do j = 0, 127
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 5
+      end do
+
+    !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
+      !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+      p1 = c_loc (a);
+      p2 = c_loc (b);
+      p3 = c_loc (c(2))
+      devptrs = [(c_loc (e(i)), i = 32, 95)]
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
+          .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+        stop 6
+
+      do j = 0, 31
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 7
+      end do
+      do j = 32, 95
+        if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
+          stop 8
+      end do
+      do j = 96, 128
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 9
+      end do
+    !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
+
+    if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
+        .or. c_associated (omp_get_mapped_ptr (q, d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d))) &
+      stop 10
+    do j = 0, 127
+      if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+        stop 11
+    end do
+
+  else ! d == id
+
+    !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
+      !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+      p1 = c_loc (a);
+      p2 = c_loc (b);
+      p3 = c_loc (c(2))
+      devptrs = [(c_loc (e(i)), i = 32, 95)]
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+      stop 12
+
+      do j = 0, 127
+        if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+          stop 13
+      end do
+    !$omp end target data
+
+    if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+      stop 14
+    do j = 0, 127
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+        stop 15
+    end do
+
+    !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
+      !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+      p1 = c_loc (a);
+      p2 = c_loc (b);
+      p3 = c_loc (c(2))
+      devptrs = [(c_loc (e(i)), i = 32, 95)]
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+        stop 16
+
+      do j = 0, 127
+        if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+          stop 17
+      end do
+    !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
+
+    if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+      stop 18
+
+    do j = 0, 127
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+        stop 19
+    end do
+  end if
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90
new file mode 100644
index 0000000..8e7ccac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90
@@ -0,0 +1,48 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id
+  type(c_ptr) :: p, p1, p2
+  integer, target :: a(1:0), b(1:2)
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+
+  if (d < 0 .or. d >= omp_get_num_devices ()) &
+    d = id
+
+  p = omp_target_alloc (c_sizeof (c_int), d)
+  if (.not. c_associated (p)) &
+    stop 0  ! okay
+
+  if (omp_target_associate_ptr (c_loc (a), p, c_sizeof (c_int), &
+                                0_c_size_t, d) == 0) then
+
+  if(.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p)) &
+    stop 1
+
+  if (omp_target_disassociate_ptr (c_loc (a), d) /= 0) &
+    stop 2
+
+  if(c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
+    stop 3
+
+  !$omp target data map(alloc: a) device(d)
+    !$omp target map(from: p1) map(alloc: a) device(d)
+    p1 = c_loc (a);
+    !$omp end target
+    if (c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
+      stop 4
+  !$omp end target data
+
+  !$omp target data map(alloc: b(1:0)) device(d)
+    !$omp target map(from: p2) map(alloc: b(1:0)) device(d)
+    p2 = c_loc (b(1));
+    !$omp end target
+    if (c_associated (omp_get_mapped_ptr (c_loc (b(1)), d))) &
+      stop 5
+  !$omp end target data
+  end if
+  call omp_target_free (p, d)
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90
new file mode 100644
index 0000000..4300a55
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90
@@ -0,0 +1,84 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id
+  type(c_ptr) :: p1, p2
+
+  type t
+    integer :: m1, m2
+  end type t
+  type(t), target :: s
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+
+  if (d < 0 .or. d >= omp_get_num_devices ()) &
+    d = id
+
+  if (d /= id) then
+    !$omp target data map(alloc: s, s%m2) device(d)
+      !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+      p1 = c_loc (s);
+      p2 = c_loc (s%m2);
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
+        stop 0
+    !$omp end target data
+
+    if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
+      stop 1
+
+    !$omp target enter data map (alloc: s, s%m2) device (d)
+      !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+      p1 = c_loc (s);
+      p2 = c_loc (s%m2);
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
+        stop 2
+    !$omp target exit data map (delete: s, s%m2) device (d)
+
+    if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
+      stop 3
+
+  else ! d == id
+
+    !$omp target data map(alloc: s, s%m2) device(d)
+      !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+      p1 = c_loc (s);
+      p2 = c_loc (s%m2);
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+        stop 4
+    !$omp end target data
+
+    if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+      stop 5
+
+    !$omp target enter data map (alloc: s, s%m2) device (d)
+      !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+      p1 = c_loc (s);
+      p2 = c_loc (s%m2);
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+        stop 6
+
+    !$omp target exit data map (delete: s, s%m2) device (d)
+
+    if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+      stop 7
+  end if
+
+end program main

Hi Jakub,

This is an update to the patch from Tue Mar 8:

https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591343.html

I just added "get_mapped_ptr" to the "omp_runtime_apis" array in omp-low.cc and
replaced "omp_get_num_devices" by "gomp_get_num_devices" in target.c.

Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.

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.

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 77176ef..02a0f72 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -3962,6 +3962,7 @@ omp_runtime_api_call (const_tree fndecl)
       "target_is_present",
       "target_memcpy",
       "target_memcpy_rect",
+      "get_mapped_ptr",
       NULL,
       /* Now omp_* calls that are available as omp_* and omp_*_; however, the
 	 DECL_NAME is always omp_* without tailing underscore.  */
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 2ac5809..608a54c 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -226,6 +226,11 @@ OMP_5.1 {
 	omp_get_teams_thread_limit_;
 } OMP_5.0.2;
 
+OMP_5.1.1 {
+  global:
+	omp_get_mapped_ptr;
+} OMP_5.1;
+
 GOMP_1.0 {
   global:
 	GOMP_atomic_end;
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 161a423..c163b56 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -314,7 +314,7 @@ The OpenMP 4.5 specification is fully supported.
 @item @code{omp_target_is_accessible} runtime routine @tab N @tab
 @item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async}
       runtime routines @tab N @tab
-@item @code{omp_get_mapped_ptr} runtime routine @tab N @tab
+@item @code{omp_get_mapped_ptr} runtime routine @tab Y @tab
 @item @code{omp_calloc}, @code{omp_realloc}, @code{omp_aligned_alloc} and
       @code{omp_aligned_calloc} runtime routines @tab Y @tab
 @item @code{omp_alloctrait_key_t} enum: @code{omp_atv_serialized} added,
diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index 89c5d65..18d0152 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -282,6 +282,7 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int,
 extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__,
 				     __SIZE_TYPE__, int) __GOMP_NOTHROW;
 extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW;
+extern void *omp_get_mapped_ptr (const void *, int) __GOMP_NOTHROW;
 
 extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW;
 extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__)
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index daf40dc..506f15c 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -835,6 +835,15 @@
           end function omp_target_disassociate_ptr
         end interface
 
+        interface
+          function omp_get_mapped_ptr (ptr, device_num) bind(c)
+            use, intrinsic :: iso_c_binding, only : c_ptr, c_int
+            type(c_ptr) :: omp_get_mapped_ptr
+            type(c_ptr), value :: ptr
+            integer(c_int), value :: device_num
+          end function omp_get_mapped_ptr
+        end interface
+
 #if _OPENMP >= 201811
 !GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested
 #endif
diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index ff857a4..0f48510 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -416,3 +416,12 @@
           integer(c_int), value :: device_num
         end function omp_target_disassociate_ptr
       end interface
+
+      interface
+        function omp_get_mapped_ptr (ptr, device_num) bind(c)
+          use, intrinsic :: iso_c_binding, only : c_ptr, c_int
+          type(c_ptr) :: omp_get_mapped_ptr
+          type(c_ptr), value :: ptr
+          integer(c_int), value :: device_num
+        end function omp_get_mapped_ptr
+      end interface
diff --git a/libgomp/target.c b/libgomp/target.c
index 9017458..86930ea 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -3665,6 +3665,44 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
   return ret;
 }
 
+void *
+omp_get_mapped_ptr (const void *ptr, int device_num)
+{
+  if (device_num < 0 || device_num > gomp_get_num_devices ())
+    return NULL;
+
+  if (device_num == omp_get_initial_device ())
+    return (void *) ptr;
+
+  struct gomp_device_descr *devicep = resolve_device (device_num);
+  if (devicep == NULL)
+    return NULL;
+
+  if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
+      || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+    return (void *) ptr;
+
+  gomp_mutex_lock (&devicep->lock);
+
+  struct splay_tree_s *mem_map = &devicep->mem_map;
+  struct splay_tree_key_s cur_node;
+  void *ret = NULL;
+
+  cur_node.host_start = (uintptr_t) ptr;
+  cur_node.host_end = cur_node.host_start;
+  splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
+
+  if (n)
+    {
+      uintptr_t offset = cur_node.host_start - n->host_start;
+      ret = (void *) (n->tgt->tgt_start + n->tgt_offset + offset);
+    }
+
+  gomp_mutex_unlock (&devicep->lock);
+
+  return ret;
+}
+
 int
 omp_pause_resource (omp_pause_resource_t kind, int device_num)
 {
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;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90
new file mode 100644
index 0000000..de05179
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90
@@ -0,0 +1,43 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id
+  type(c_ptr) :: p
+  integer, target :: q
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+
+  if (d < 0 .or. d >= omp_get_num_devices ()) &
+    d = id
+
+  p = omp_target_alloc (c_sizeof (q), d)
+  if (.not. c_associated (p)) &
+    stop 0  ! okay
+
+  if (omp_target_associate_ptr (c_loc (q), p, c_sizeof (q), &
+                                0_c_size_t, d) == 0) then
+
+    if(c_associated (omp_get_mapped_ptr (c_loc (q), -1))) &
+      stop 1
+
+    if(c_associated (omp_get_mapped_ptr (c_loc (q), &
+                     omp_get_num_devices () + 1))) &
+      stop 2
+
+    if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), id), c_loc (q))) &
+      stop 3
+
+    if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), d), p)) &
+      stop 4
+
+    if (omp_target_disassociate_ptr (c_loc (q), d) /= 0) &
+      stop 5
+
+    if(c_associated (omp_get_mapped_ptr (c_loc (q), d))) &
+      stop 6
+  end if
+
+  call omp_target_free (p, d)
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90
new file mode 100644
index 0000000..66a0b88
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90
@@ -0,0 +1,175 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id, i, j
+  integer, target :: a, b(1:2), c(1:2), e(0:127)
+  type(c_ptr) :: p1, p2, p3, q, devptrs(0:63)
+
+  a = 42;
+  q = c_loc (a);
+  e = [(i, i = 0, 127)]
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+
+  if (d < 0 .or. d >= omp_get_num_devices ()) &
+    d = id
+
+  if (d /= id) then
+    !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
+      !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+      p1 = c_loc (a);
+      p2 = c_loc (b);
+      p3 = c_loc (c(2))
+      devptrs = [(c_loc (e(i)), i = 32, 95)]
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), p2) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
+          .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+        stop 0
+
+      do j = 0, 31
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 1
+      end do
+      do j = 32, 95
+        if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
+          stop 2
+      end do
+      do j = 96, 128
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 3
+      end do
+    !$omp end target data
+
+    if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
+        .or. c_associated (omp_get_mapped_ptr (q, d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+      stop 4
+      do j = 0, 127
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 5
+      end do
+
+    !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
+      !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+      p1 = c_loc (a);
+      p2 = c_loc (b);
+      p3 = c_loc (c(2))
+      devptrs = [(c_loc (e(i)), i = 32, 95)]
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
+          .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+        stop 6
+
+      do j = 0, 31
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 7
+      end do
+      do j = 32, 95
+        if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
+          stop 8
+      end do
+      do j = 96, 128
+        if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+          stop 9
+      end do
+    !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
+
+    if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
+        .or. c_associated (omp_get_mapped_ptr (q, d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d))) &
+      stop 10
+    do j = 0, 127
+      if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+        stop 11
+    end do
+
+  else ! d == id
+
+    !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
+      !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+      p1 = c_loc (a);
+      p2 = c_loc (b);
+      p3 = c_loc (c(2))
+      devptrs = [(c_loc (e(i)), i = 32, 95)]
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+      stop 12
+
+      do j = 0, 127
+        if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+          stop 13
+      end do
+    !$omp end target data
+
+    if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+      stop 14
+    do j = 0, 127
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+        stop 15
+    end do
+
+    !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
+      !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+      p1 = c_loc (a);
+      p2 = c_loc (b);
+      p3 = c_loc (c(2))
+      devptrs = [(c_loc (e(i)), i = 32, 95)]
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+        stop 16
+
+      do j = 0, 127
+        if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+          stop 17
+      end do
+    !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
+
+    if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+      stop 18
+
+    do j = 0, 127
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+        stop 19
+    end do
+  end if
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90
new file mode 100644
index 0000000..8e7ccac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90
@@ -0,0 +1,48 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id
+  type(c_ptr) :: p, p1, p2
+  integer, target :: a(1:0), b(1:2)
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+
+  if (d < 0 .or. d >= omp_get_num_devices ()) &
+    d = id
+
+  p = omp_target_alloc (c_sizeof (c_int), d)
+  if (.not. c_associated (p)) &
+    stop 0  ! okay
+
+  if (omp_target_associate_ptr (c_loc (a), p, c_sizeof (c_int), &
+                                0_c_size_t, d) == 0) then
+
+  if(.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p)) &
+    stop 1
+
+  if (omp_target_disassociate_ptr (c_loc (a), d) /= 0) &
+    stop 2
+
+  if(c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
+    stop 3
+
+  !$omp target data map(alloc: a) device(d)
+    !$omp target map(from: p1) map(alloc: a) device(d)
+    p1 = c_loc (a);
+    !$omp end target
+    if (c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
+      stop 4
+  !$omp end target data
+
+  !$omp target data map(alloc: b(1:0)) device(d)
+    !$omp target map(from: p2) map(alloc: b(1:0)) device(d)
+    p2 = c_loc (b(1));
+    !$omp end target
+    if (c_associated (omp_get_mapped_ptr (c_loc (b(1)), d))) &
+      stop 5
+  !$omp end target data
+  end if
+  call omp_target_free (p, d)
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90
new file mode 100644
index 0000000..4300a55
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90
@@ -0,0 +1,84 @@
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d, id
+  type(c_ptr) :: p1, p2
+
+  type t
+    integer :: m1, m2
+  end type t
+  type(t), target :: s
+
+  d = omp_get_default_device ()
+  id = omp_get_initial_device ()
+
+  if (d < 0 .or. d >= omp_get_num_devices ()) &
+    d = id
+
+  if (d /= id) then
+    !$omp target data map(alloc: s, s%m2) device(d)
+      !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+      p1 = c_loc (s);
+      p2 = c_loc (s%m2);
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
+        stop 0
+    !$omp end target data
+
+    if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
+      stop 1
+
+    !$omp target enter data map (alloc: s, s%m2) device (d)
+      !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+      p1 = c_loc (s);
+      p2 = c_loc (s%m2);
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
+          .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
+        stop 2
+    !$omp target exit data map (delete: s, s%m2) device (d)
+
+    if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
+        .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
+      stop 3
+
+  else ! d == id
+
+    !$omp target data map(alloc: s, s%m2) device(d)
+      !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+      p1 = c_loc (s);
+      p2 = c_loc (s%m2);
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+        stop 4
+    !$omp end target data
+
+    if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+      stop 5
+
+    !$omp target enter data map (alloc: s, s%m2) device (d)
+      !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+      p1 = c_loc (s);
+      p2 = c_loc (s%m2);
+      !$omp end target
+
+      if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+        stop 6
+
+    !$omp target exit data map (delete: s, s%m2) device (d)
+
+    if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+        .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+      stop 7
+  end if
+
+end program main

On Thu, Mar 10, 2022 at 05:01:35PM +0100, Marcel Vollweiler wrote:
> --- a/gcc/omp-low.cc
> +++ b/gcc/omp-low.cc
> @@ -3962,6 +3962,7 @@ omp_runtime_api_call (const_tree fndecl)
>        "target_is_present",
>        "target_memcpy",
>        "target_memcpy_rect",
> +      "get_mapped_ptr",
>        NULL,
>        /* Now omp_* calls that are available as omp_* and omp_*_; however, the
>  	 DECL_NAME is always omp_* without tailing underscore.  */

The entries in each NULL separated subsection are supposed to be sorted
alphabetically.
Other than that LGTM, but stage1 is still far...

	Jakub