diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 258ea8a7619..b4442069f66 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -2122,18 +2122,9 @@ is not supported.
 @item @emph{Description}:
 This routine tests whether storage, identified by the host pointer @var{ptr}
 is mapped to the device specified by @var{device_num}.  If so, it returns
-a nonzero value and otherwise zero.
-
-In GCC, this includes self mapping such that @code{omp_target_is_present}
-returns @emph{true} when @var{device_num} specifies the host or when the host
-and the device share memory.  If @var{ptr} is a null pointer, @var{true} is
-returned and if @var{device_num} is an invalid device number, @var{false} is
-returned.
-
-If those conditions do not apply, @emph{true} is returned if the association has
-been established by an explicit or implicit @code{map} clause, the
-@code{declare target} directive or a call to the @code{omp_target_associate_ptr}
-routine.
+a nonzero value and otherwise zero.  In particular, it always returns zero
+for the null pointer and for invalid device numbers; for the host device,
+a nonzero value is returned for all non-null pointers.
 
 Running this routine in a @code{target} region except on the initial device
 is not supported.
@@ -2154,7 +2145,7 @@ is not supported.
 @end multitable
 
 @item @emph{See also}:
-@ref{omp_target_associate_ptr}
+@ref{omp_get_mapped_ptr}, @ref{omp_target_associate_ptr}, @ref{omp_target_is_accessible}
 
 @item @emph{Reference}:
 @uref{https://www.openmp.org, OpenMP specification v5.1}, Section 3.8.3
@@ -2204,7 +2195,7 @@ is not supported.
 @end multitable
 
 @item @emph{See also}:
-@ref{omp_target_associate_ptr}
+@ref{omp_target_is_present}
 
 @item @emph{Reference}:
 @uref{https://www.openmp.org, OpenMP specification v5.1}, Section 3.8.4
@@ -2599,7 +2590,8 @@ is not supported.
 
 @item @emph{See also}:
 @ref{omp_target_disassociate_ptr}, @ref{omp_target_is_present},
-@ref{omp_target_alloc}
+@ref{omp_get_mapped_ptr}, @ref{omp_target_alloc}
+
 
 @item @emph{Reference}:
 @uref{https://www.openmp.org, OpenMP specification v5.1}, Section 3.8.9
@@ -2657,13 +2649,20 @@ is not supported.
 @subsection @code{omp_get_mapped_ptr} -- Return device pointer to a host pointer
 @table @asis
 @item @emph{Description}:
-If the device number is refers to the initial device or to a device with
-memory accessible from the host (shared memory), the @code{omp_get_mapped_ptr}
-routines returns the value of the passed @var{ptr}.  Otherwise, if associated
+If the device number refers to the initial device, @code{omp_get_mapped_ptr}
+returns the value of the passed @var{ptr}.  Otherwise, if associated
 storage to the passed host pointer @var{ptr} exists on device associated with
 @var{device_num}, it returns that pointer. In all other cases and in cases of
 an error, a null pointer is returned.
 
+If the device number is not the initial device and the pointer points to a
+variable that is specified in a @code{declare target} directive: When
+requiring @code{unified_shared_memory} or @code{self_maps}, a null pointer is
+returned if the variable appears in a @code{link} or @code{enter} clause.
+Otherwise, the corresponding device memory is returned; with the @code{link}
+clause, GCC returns the address of the pointer-typed link variable on the device,
+not to the data that is mapped to that variable.
+
 The association of storage location is established either via an explicit or
 implicit @code{map} clause, the @code{declare target} directive or the
 @code{omp_target_associate_ptr} routine.
@@ -2685,7 +2684,7 @@ is not supported.
 @end multitable
 
 @item @emph{See also}:
-@ref{omp_target_associate_ptr}
+@ref{omp_target_is_present}, @ref{omp_target_associate_ptr}
 
 @item @emph{Reference}:
 @uref{https://www.openmp.org, OpenMP specification v5.1}, Section 3.8.11
@@ -4096,7 +4095,7 @@ The value can either be a predefined allocator or a predefined memory space
 or a predefined memory space followed by a colon and a comma-separated list
 of memory trait and value pairs, separated by @code{=}.
 
-See @ref{Memory allocation} for a list of supported prefedined allocators,
+See @ref{Memory allocation} for a list of supported predefined allocators,
 memory spaces, and traits.
 
 Note: The corresponding device environment variables are currently not
diff --git a/libgomp/target.c b/libgomp/target.c
index c106c8cafa3..d562b0493ea 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -4874,9 +4874,24 @@ gomp_page_locked_host_free (void *ptr)
 		device->name);
 }
 
+/* Check whether corresponding storage exists on the device.
+   - NULL pointer or invalid device: return 0
+   - host device: return 1
+   - Has corresponding storage: return 1
+   - Otherwise: return 0
+
+   Note that for GOMP_OFFLOAD_CAP_SHARED_MEM self mapping is used and
+   omp_target_associate_ptr is disabled; the only corresponding storage
+   exists then for declare_target with other clauses than an explicit or
+   implicit 'link' clause.
+   However, the link cause with shared memory does not count as mapped.  */
+
 int
 omp_target_is_present (const void *ptr, int device_num)
 {
+  if (ptr == NULL)
+    return 0;
+
   if (device_num == omp_default_device)
     device_num = gomp_get_default_device ();
 
@@ -4888,13 +4903,8 @@ omp_target_is_present (const void *ptr, int device_num)
   if (devicep == NULL)
     return 0;
 
-  if (ptr == NULL)
-    return 1;
-
-  if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
-      || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
-    return 1;
-
+  bool is_shared = (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
+		    || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM);
   gomp_mutex_lock (&devicep->lock);
   struct splay_tree_s *mem_map = &devicep->mem_map;
   struct splay_tree_key_s cur_node;
@@ -4902,7 +4912,7 @@ omp_target_is_present (const void *ptr, int device_num)
   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);
-  int ret = n != NULL;
+  int ret = n != NULL && (!is_shared || n->refcount != REFCOUNT_LINK);
   gomp_mutex_unlock (&devicep->lock);
   return ret;
 }
@@ -5570,7 +5580,8 @@ omp_get_mapped_ptr (const void *ptr, int device_num)
   if (device_num == omp_default_device)
     device_num = gomp_get_default_device ();
 
-  if (device_num == omp_initial_device
+  if (ptr == NULL
+      || device_num == omp_initial_device
       || device_num == omp_get_initial_device ())
     return (void *) ptr;
 
@@ -5578,10 +5589,8 @@ omp_get_mapped_ptr (const void *ptr, int 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;
-
+  bool is_shared = (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
+		    || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM);
   gomp_mutex_lock (&devicep->lock);
 
   struct splay_tree_s *mem_map = &devicep->mem_map;
@@ -5596,6 +5605,8 @@ omp_get_mapped_ptr (const void *ptr, int device_num)
     {
       uintptr_t offset = cur_node.host_start - n->host_start;
       ret = (void *) (n->tgt->tgt_start + n->tgt_offset + offset);
+      if (is_shared && n->refcount == REFCOUNT_LINK)
+	ret = NULL;
     }
 
   gomp_mutex_unlock (&devicep->lock);
diff --git a/libgomp/testsuite/libgomp.c/omp_target_is_present-2.c b/libgomp/testsuite/libgomp.c/omp_target_is_present-2.c
new file mode 100644
index 00000000000..ee4d8215f33
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_target_is_present-2.c
@@ -0,0 +1,7 @@
+// { dg-do run }
+
+#define REQ_SELF_MAPS 1
+
+#pragma omp requires self_maps
+
+#include "omp_target_is_present.c"
diff --git a/libgomp/testsuite/libgomp.c/omp_target_is_present.c b/libgomp/testsuite/libgomp.c/omp_target_is_present.c
new file mode 100644
index 00000000000..b7638e712ba
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_target_is_present.c
@@ -0,0 +1,206 @@
+// { dg-do run }
+
+// Check mainly omp_target_is_present - but also some related functions
+
+/* omp_target_is_present is only 1 if device == host or when there is corresponding
+   storage on the device, which implies ptr != omp_get_mapped_ptr (ptr, dev).
+
+   Note that a NULL ptr is regarded as not being present.  */
+
+#include <omp.h>
+
+#ifndef REQ_SELF_MAPS
+  #define REQ_SELF_MAPS 0
+#endif
+
+// FIXME: change enter to link clause for gLink, once implemented
+
+int gEnter = 3, gLink = 4, gLocal = 5;
+#pragma omp declare target enter(gEnter) link(gLink) enter(gLocal)
+
+void check_routines (int dev)
+{
+  int A = 1, B = 2;
+
+  int dev2 = dev;
+  if (dev2 == omp_default_device)
+    dev2 = omp_get_default_device ();
+
+  bool initial_dev = dev2 == omp_initial_device || dev2 == omp_get_num_devices();
+  bool self_mapping = false;
+  bool invalid_dev = dev == omp_invalid_device;
+  if (!invalid_dev && !initial_dev)
+    {
+      #pragma omp target map(to: self_mapping) device(dev)
+	self_mapping = true;
+      if (REQ_SELF_MAPS && !self_mapping)
+	__builtin_abort ();
+    }
+
+  if (omp_target_is_present (nullptr, dev) != 0)
+    __builtin_abort ();
+
+  if (omp_target_is_accessible (nullptr, 0, dev) != 0)
+    __builtin_abort ();
+
+
+  if (invalid_dev)
+    return; // Will otherwise fail with: libgomp: omp_invalid_device encountered
+
+
+  if (omp_target_is_present (&A, dev) != initial_dev)
+    __builtin_abort ();
+
+  // For link, it points to the pointer var - FIXME: update for self_maps implying 'link'
+  if (omp_target_is_present (&gEnter, dev) != !invalid_dev)
+    __builtin_abort ();
+
+  if (omp_target_is_present (&gLink, dev) != (!invalid_dev && (initial_dev || !REQ_SELF_MAPS)))
+    __builtin_abort ();
+
+  if (omp_target_is_present (&gLocal, dev) != !invalid_dev)
+    __builtin_abort ();
+
+  int *ptr = (int*) 0xDEEDBEEF;
+  if (!invalid_dev)
+    {
+      #pragma omp target enter data map(to: A) device(dev)
+      #pragma omp target enter data map(to: gEnter) device(dev)
+      #pragma omp target enter data map(to: gLink) device(dev)
+      #pragma omp target enter data map(to: gLocal) device(dev)
+
+      ptr = omp_target_alloc (sizeof (int), dev);
+      if (ptr == nullptr || !omp_target_is_accessible (ptr, sizeof (int), dev))
+	__builtin_abort ();
+    }
+
+  // Invalid
+  if ((initial_dev || invalid_dev) && omp_target_associate_ptr (ptr, ptr, sizeof (int), 0, dev) == 0)
+    __builtin_abort ();
+  if ((initial_dev || invalid_dev) && omp_target_associate_ptr (((char*)ptr) + 2, ptr, sizeof (int)-2, 2, dev) == 0)
+    __builtin_abort ();
+
+  // Should yield 0/success except for self mapping, host or invalid device
+  // use !! to convert the result to 0 or 1, as errors can also be, e.g. EINVAL
+  if (!!omp_target_associate_ptr (&B, ptr, sizeof (int), 0, dev)
+      != (self_mapping || initial_dev || invalid_dev))
+    __builtin_abort ();
+
+  // Try again, should still work as it is the same pointer
+  if (!!omp_target_associate_ptr (&B, ptr, sizeof (int), 0, dev)
+      != (self_mapping || initial_dev || invalid_dev))
+    __builtin_abort ();
+
+  if (!!omp_target_is_present (&A, dev)
+      != (initial_dev || (!self_mapping && !invalid_dev)))
+    __builtin_abort ();
+
+  if (!!omp_target_is_present (&B, dev)
+      != (initial_dev || (!self_mapping && !invalid_dev)))
+    __builtin_abort ();
+
+  if (!!omp_target_is_present (&gEnter, dev)
+      != (initial_dev || (/* !self_mapping && */ !invalid_dev)))
+    __builtin_abort ();
+
+  if (!!omp_target_is_present (&gLink, dev)
+      != (initial_dev || (!self_mapping && !invalid_dev)))
+    __builtin_abort ();
+
+  if (!!omp_target_is_present (&gLocal, dev) != !invalid_dev)
+    __builtin_abort ();
+
+  int *ptr2 = omp_get_mapped_ptr (&A, dev);
+  if (initial_dev)
+    {
+      if (ptr2 != &A)
+	__builtin_abort ();
+    }
+  else if (invalid_dev || self_mapping)
+    {
+      if (ptr2 != nullptr)
+	__builtin_abort ();
+    }
+  else if (ptr2 == &A || ptr2 == nullptr)
+    __builtin_abort ();
+
+  ptr2 = omp_get_mapped_ptr (&B, dev);
+  if (initial_dev)
+    {
+      if (ptr2 != &B)
+	__builtin_abort ();
+    }
+  else if (invalid_dev || self_mapping)
+    {
+      if (ptr2 != nullptr)
+	__builtin_abort ();
+    }
+  else if (ptr2 != ptr)
+    __builtin_abort ();
+
+  ptr2 = omp_get_mapped_ptr (&gEnter, dev);
+  if (initial_dev)
+    {
+      if (ptr2 != &gEnter)
+	__builtin_abort ();
+    }
+  else if (invalid_dev /* FIXME: || self_mapping */)
+    {
+      if (ptr2 != nullptr)
+	__builtin_abort ();
+    }
+  else if (ptr2 == &gEnter || ptr2 == nullptr)
+    __builtin_abort ();
+
+  ptr2 = omp_get_mapped_ptr (&gLink, dev);
+  if (initial_dev)
+    {
+      if (ptr2 != &gLink)
+	__builtin_abort ();
+    }
+  else if (invalid_dev || self_mapping)
+    {
+      if (ptr2 != nullptr)
+	__builtin_abort ();
+    }
+  else if (ptr2 == ptr || ptr2 == nullptr)
+    __builtin_abort ();
+
+  ptr2 = omp_get_mapped_ptr (&gLocal, dev);
+  if (initial_dev)
+    {
+      if (ptr2 != &gLocal)
+	__builtin_abort ();
+    }
+  else if (invalid_dev)
+    {
+      if (ptr2 != nullptr)
+	__builtin_abort ();
+    }
+  else if (ptr2 == &gLocal || ptr2 == nullptr)
+    __builtin_abort ();
+
+  if (!invalid_dev)
+    {
+      omp_target_free (ptr, dev);
+      #pragma omp target exit data map(release: A) device(dev)
+      #pragma omp target exit data map(release: gLink) device(dev)
+      #pragma omp target exit data map(release: gEnter) device(dev)
+      #pragma omp target exit data map(release: gLocal) device(dev)
+    }
+}
+
+int main()
+{
+  for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+    check_routines (dev);
+
+  check_routines (omp_invalid_device);
+  check_routines (omp_default_device);
+
+  for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)
+    {
+      omp_set_default_device (dev);
+      check_routines (omp_default_device);
+    }
+}