libgomp: fix omp_target_is_present and omp_get_mapped_ptr

There were a few minor issues with the two routines, partially because of
not handling corner cases and partially some clarifications are only in
newer versions of the spec.

In particular, for omp_target_is_present
* NULL pointer aren't regarded as present
* For (unified-)shared memory, claiming that something has always corresponding
  storage is wrong - it mostly never has. (but it is omp_target_is_accessible).
* Even with shared memory, 'declare target' usually has device memory. For
  'link' it is made to point to the host, i.e. it is not mapped, all others
  are still mapped. (With 'requires self_mapping', 'enter' should also not be
  mapped (and turned internally to 'link'), only 'local' needs to be mapped.)

For omp_get_mapped_ptr
* For NULL we can return NULL early also for devices.
* For shared memory, we shouldn't touch link (it is not counting as mapped);
  hence return NULL for it.

The documentation was updated add some missing cross references as the more
useful ones were missing.  Additionally, the description for the two modified
routines has been updated.

libgomp/ChangeLog:

	* target.c (omp_target_is_present, omp_get_mapped_ptr): Update handling
	for nullptr and shared-memory devices.
	* libgomp.texi (omp_target_is_present, omp_get_mapped_ptr): Update
	description, add see-also @refs.
	(omp_target_is_accessible, omp_target_associate_ptr): Add see-also
	@refs.
	* testsuite/libgomp.c/omp_target_is_present.c: New test.
	* testsuite/libgomp.c/omp_target_is_present-2.c: New test.

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

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);

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"

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);
+    }
+}