@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.
@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
@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
@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
@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.
@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
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
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 ();
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;
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;
}
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;
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;
{
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);
--- /dev/null
+// { dg-do run }
+
+#define REQ_SELF_MAPS 1
+
+#pragma omp requires self_maps
+
+#include "omp_target_is_present.c"
--- /dev/null
+// { 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);
+ }
+}
projects / thirdparty / gcc.git / commitdiff
