OpenMP: Add omp_default_device named constant [PR119677]

OpenMP TR 14 (OpenMP 6.1) adds omp_default_device < -1 as
named constant alongside omp_initial_device and omp_default_device.

GCC supports it already internally via GOMP_DEVICE_DEFAULT_OMP_61,
but this patch now adds the omp_default_device enum/PARAMETER to
omp.h / omp_lib.

Note that PR119677 requests some cleanups, which still have to be
done.

	PR libgomp/119677

gcc/fortran/ChangeLog:

	* intrinsic.texi (OpenMP Modules): Add omp_default_device.
	* openmp.cc (gfc_resolve_omp_context_selector): Accept
	omp_default_device as conforming device number.

libgomp/ChangeLog:

	* omp.h.in (omp_default_device): New enum value.
	* omp_lib.f90.in: New parameter.
	* omp_lib.h.in: Likewise
	* target.c (gomp_get_default_device): New. Split off from ...
	(resolve_device): ... here; call it.
	(omp_target_alloc, omp_target_free, omp_target_is_present,
	omp_target_memcpy_check, omp_target_memset, omp_target_memset_async,
	omp_target_associate_ptr, omp_get_mapped_ptr,
	omp_target_is_accessible, omp_pause_resource,
	omp_get_uid_from_device): Handle omp_default_device.
	* testsuite/libgomp.c/device_uid.c: Likewise.
	* testsuite/libgomp.fortran/device_uid.f90: Likewise.
	* testsuite/libgomp.c-c++-common/omp-default-device.c: New test.
	* testsuite/libgomp.fortran/omp-default-device.f90: New test.

diff --git a/gcc/fortran/intrinsic.texi b/gcc/fortran/intrinsic.texi
index b2d1e4559883e48b726a8ab0ae64150e6850ae03..69cf10caf9dbb5ba0d00e3e65aa6c93da2285522 100644 (file)
--- a/gcc/fortran/intrinsic.texi
+++ b/gcc/fortran/intrinsic.texi
@@ -16250,6 +16250,7 @@ The following scalar default-integer named constants:
 @table @asis
 @item @code{omp_initial_device}
 @item @code{omp_invalid_device}
+@item @code{omp_default_device}
 @end table
 
 

diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 770bc5b1200fff20428946fa224a32cd702a2e31..9c9d0c2c5340bb40702fc2f426b8cb56dd2dfe63 100644 (file)
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -12290,12 +12290,14 @@ gfc_resolve_omp_context_selector (gfc_omp_set_selector *oss,
                continue;
              }
            /* Device number must be conforming, which includes
-              omp_initial_device (-1) and omp_invalid_device (-4).  */
+              omp_initial_device (-1), omp_invalid_device (-4),
+              and omp_default_device (-5).  */
            if (property_kind == OMP_TRAIT_PROPERTY_DEV_NUM_EXPR
                && otp->expr->expr_type == EXPR_CONSTANT
                && mpz_sgn (otp->expr->value.integer) < 0
                && mpz_cmp_si (otp->expr->value.integer, -1) != 0
-               && mpz_cmp_si (otp->expr->value.integer, -4) != 0)
+               && mpz_cmp_si (otp->expr->value.integer, -4) != 0
+               && mpz_cmp_si (otp->expr->value.integer, -5) != 0)
              gfc_error ("property must be a conforming device number at %L",
                         &otp->expr->where);
            break;

diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index 60cb2b21be74f05e60a4c2e6c4eea3242d7c2acf..566a3c28b942f042018df60ca59cccf91871f2e7 100644 (file)
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -189,7 +189,8 @@ typedef enum omp_event_handle_t __GOMP_UINTPTR_T_ENUM
 enum
 {
   omp_initial_device = -1,
-  omp_invalid_device = -4
+  omp_invalid_device = -4,
+  omp_default_device = -5
 };
 
 typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM

diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index ce866c00121a20caf4754e71c519e35122c352cd..74e0bfea344c1df4634d472ba5a2d439cea913fd 100644 (file)
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -176,6 +176,7 @@
                  parameter :: omp_low_lat_mem_space = 4
         integer, parameter :: omp_initial_device = -1
         integer, parameter :: omp_invalid_device = -4
+        integer, parameter :: omp_default_device = -5
         integer (omp_interop_kind), &
                  parameter :: omp_interop_none = 0_omp_interop_kind
         integer (omp_interop_fr_kind), parameter :: omp_ifr_cuda = 1

diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index 9047095c5e0bd8902728c5211a59cff79270218a..9422515dc379ca8ff85db4f54897151fae705bf7 100644 (file)
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -182,9 +182,10 @@
       parameter (omp_const_mem_space = 2)
       parameter (omp_high_bw_mem_space = 3)
       parameter (omp_low_lat_mem_space = 4)
-      integer omp_initial_device, omp_invalid_device
+      integer omp_initial_device, omp_invalid_device, omp_default_device
       parameter (omp_initial_device = -1)
       parameter (omp_invalid_device = -4)
+      parameter (omp_default_device = -5)
       integer (omp_interop_kind) omp_interop_none
       parameter (omp_interop_none = 0_omp_interop_kind)
       integer (omp_interop_fr_kind) omp_ifr_cuda

diff --git a/libgomp/target.c b/libgomp/target.c
index ac5b4b0b7209a2e4b534bbc1ddd871f63db30414..002a144b4abb181b4e8508730e40aaf4f327f0fc 100644 (file)
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -139,6 +139,14 @@ gomp_get_num_devices (void)
   return num_devices_openmp;
 }
 
+static int
+gomp_get_default_device ()
+{
+  gomp_init_targets_once ();
+  struct gomp_task_icv *icv = gomp_icv (false);
+  return icv->default_device_var;
+}
+
 static struct gomp_device_descr *
 resolve_device (int device_id, bool remapped)
 {
@@ -148,11 +156,7 @@ resolve_device (int device_id, bool remapped)
 
   if ((remapped && device_id == GOMP_DEVICE_ICV)
       || device_id == GOMP_DEVICE_DEFAULT_OMP_61)
-    {
-      struct gomp_task_icv *icv = gomp_icv (false);
-      device_id = icv->default_device_var;
-      remapped = false;
-    }
+    device_id = gomp_get_default_device ();
 
   if (device_id < 0)
     {
@@ -4653,6 +4657,9 @@ GOMP_teams4 (unsigned int num_teams_low, unsigned int num_teams_high,
 void *
 omp_target_alloc (size_t size, int device_num)
 {
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
+
   if (device_num == omp_initial_device
       || device_num == gomp_get_num_devices ())
     return malloc (size);
@@ -4674,6 +4681,9 @@ omp_target_alloc (size_t size, int device_num)
 void
 omp_target_free (void *device_ptr, int device_num)
 {
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
+
   if (device_num == omp_initial_device
       || device_num == gomp_get_num_devices ())
     {
@@ -4811,6 +4821,9 @@ gomp_page_locked_host_free (void *ptr)
 int
 omp_target_is_present (const void *ptr, int device_num)
 {
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
+
   if (device_num == omp_initial_device
       || device_num == gomp_get_num_devices ())
     return 1;
@@ -4843,6 +4856,11 @@ omp_target_memcpy_check (int dst_device_num, int src_device_num,
                         struct gomp_device_descr **dst_devicep,
                         struct gomp_device_descr **src_devicep)
 {
+  if (dst_device_num == omp_default_device)
+    dst_device_num = gomp_get_default_device ();
+  if (src_device_num == omp_default_device)
+    src_device_num = gomp_get_default_device ();
+
   if (dst_device_num != gomp_get_num_devices ()
       /* Above gomp_get_num_devices has to be called unconditionally.  */
       && dst_device_num != omp_initial_device)
@@ -5323,6 +5341,9 @@ omp_target_memset_int (void *ptr, int val, size_t count,
 void*
 omp_target_memset (void *ptr, int val, size_t count, int device_num)
 {
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
+
   struct gomp_device_descr *devicep;
   if (device_num == omp_initial_device
       || device_num == gomp_get_num_devices ()
@@ -5359,6 +5380,9 @@ omp_target_memset_async (void *ptr, int val, size_t count, int device_num,
   unsigned flags = 0;
   int i;
 
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
+
   if (device_num == omp_initial_device
       || device_num == gomp_get_num_devices ()
       || (devicep = resolve_device (device_num, false)) == NULL
@@ -5387,6 +5411,9 @@ int
 omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
                          size_t size, size_t device_offset, int device_num)
 {
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
+
   if (device_num == omp_initial_device
       || device_num == gomp_get_num_devices ())
     return EINVAL;
@@ -5484,6 +5511,9 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
 void *
 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
       || device_num == omp_get_initial_device ())
     return (void *) ptr;
@@ -5520,6 +5550,9 @@ omp_get_mapped_ptr (const void *ptr, int device_num)
 int
 omp_target_is_accessible (const void *ptr, size_t size, int device_num)
 {
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
+
   if (device_num == omp_initial_device
       || device_num == gomp_get_num_devices ())
     return true;
@@ -5537,6 +5570,8 @@ int
 omp_pause_resource (omp_pause_resource_t kind, int device_num)
 {
   (void) kind;
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
   if (device_num == omp_initial_device
       || device_num == gomp_get_num_devices ())
     return gomp_pause_host ();
@@ -5847,6 +5882,9 @@ gomp_get_uid_for_device (struct gomp_device_descr *devicep, int device_num)
 const char *
 omp_get_uid_from_device (int device_num)
 {
+  if (device_num == omp_default_device)
+    device_num = gomp_get_default_device ();
+
   if (device_num < omp_initial_device || device_num > gomp_get_num_devices ())
     return NULL;
 

diff --git a/libgomp/testsuite/libgomp.c-c++-common/omp-default-device.c b/libgomp/testsuite/libgomp.c-c++-common/omp-default-device.c
new file mode 100644 (file)
index 0000000..5489f01
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/omp-default-device.c
@@ -0,0 +1,59 @@
+#include <omp.h>
+
+#if __cplusplus
+static_assert (omp_default_device < -1
+              && omp_default_device != omp_invalid_device, "");
+#else
+_Static_assert (omp_default_device < -1
+               && omp_default_device != omp_invalid_device, "");
+#endif
+
+static int
+is_same_dev (int d1, int d2)
+{
+  int num_dev = omp_get_num_devices ();
+  if (d1 == omp_initial_device)
+    d1 = num_dev;
+  if (d2 == omp_initial_device)
+    d2 = num_dev;
+  return (d1 == d2);
+}
+
+int
+main()
+{
+  int dev = -99;
+  int def_dev = omp_get_default_device ();
+  #pragma omp target map(from: dev) device(omp_default_device)
+    dev = omp_get_device_num ();
+
+  if (!is_same_dev (def_dev, dev))
+    __builtin_abort ();
+
+  for (def_dev = omp_initial_device; def_dev <= omp_get_num_devices ();
+       def_dev++)
+    {
+      const char* uid = omp_get_uid_from_device(def_dev);
+      omp_set_default_device (def_dev);
+      dev = -99;
+      #pragma omp target map(from: dev) device(omp_default_device)
+        dev = omp_get_device_num ();
+      if (!is_same_dev (def_dev, dev))
+        __builtin_abort ();
+
+      /* Shall not modify the ICV.  */
+      omp_set_default_device (omp_default_device);
+      if (def_dev != omp_get_default_device ())
+        __builtin_abort ();
+
+      /* Assume the ptr and no only the string is the same.  */
+      if (uid != omp_get_uid_from_device (omp_default_device))
+        __builtin_abort ();
+    }
+
+  omp_set_default_device (omp_invalid_device);
+  /* Shall not modify the ICV.  */
+  omp_set_default_device (omp_default_device);
+  if (omp_invalid_device != omp_get_default_device ())
+    __builtin_abort ();
+}

diff --git a/libgomp/testsuite/libgomp.c/device_uid.c b/libgomp/testsuite/libgomp.c/device_uid.c
index 0412d06f615598729763f55155890e9b1ec4f6f7..83aba0f6a91829eb25a4665f1d05f99f3d8bf56a 100644 (file)
--- a/libgomp/testsuite/libgomp.c/device_uid.c
+++ b/libgomp/testsuite/libgomp.c/device_uid.c
@@ -5,10 +5,12 @@
 int main()
 {
   const char **strs = (const char **) malloc (sizeof (char*) * (omp_get_num_devices () + 1));
-  for (int i = omp_invalid_device - 1; i <= omp_get_num_devices () + 1; i++)
+  for (int i = omp_default_device - 1; i <= omp_get_num_devices () + 1; i++)
     {
       const char *str = omp_get_uid_from_device (i);
       int dev = omp_get_device_from_uid (str);
+      if (i == omp_default_device)
+       i = omp_get_default_device ();
 // __builtin_printf("%i -> %s -> %d\n", i, str, dev);
       if (i < omp_initial_device || i > omp_get_num_devices ())
        {

diff --git a/libgomp/testsuite/libgomp.fortran/device_uid.f90 b/libgomp/testsuite/libgomp.fortran/device_uid.f90
index 504f6caaf07dc6d858caccb88f0004cfb57b945d..9bc02e4b8e618fd4f05d83d5bd57d8d2622bfba8 100644 (file)
--- a/libgomp/testsuite/libgomp.fortran/device_uid.f90
+++ b/libgomp/testsuite/libgomp.fortran/device_uid.f90
@@ -10,10 +10,13 @@ program main
 
   allocate(strs(0:omp_get_num_devices ()))
 
-  do i = omp_invalid_device - 1, omp_get_num_devices () + 1
+  do j = omp_default_device - 1, omp_get_num_devices () + 1
+    i = j
     str => omp_get_uid_from_device (i)
     dev = omp_get_device_from_uid (str)
 ! print *, i, str, dev
+    if (i == omp_default_device) &
+      i = omp_get_default_device ()
     if (i < omp_initial_device .or. i > omp_get_num_devices ()) then
       if (dev /= omp_invalid_device .or. associated(str)) &
         stop 1

diff --git a/libgomp/testsuite/libgomp.fortran/omp-default-device.f90 b/libgomp/testsuite/libgomp.fortran/omp-default-device.f90
new file mode 100644 (file)
index 0000000..28e3496
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/omp-default-device.f90
@@ -0,0 +1,61 @@
+program main
+  use omp_lib
+  implicit none (type, external)
+  integer :: dev, def_dev
+
+  if (omp_default_device >= -1 .or. omp_default_device == omp_invalid_device) &
+    error stop 1
+
+  dev = -99
+  def_dev = omp_get_default_device ()
+  !$omp target map(from: dev) device(omp_default_device)
+    dev = omp_get_device_num ()
+  !$omp end target
+
+  if (.not.is_same_dev (def_dev, dev)) &
+    error stop 2
+
+  do def_dev = omp_initial_device, omp_get_num_devices ()
+  block
+    character(:), pointer :: uid
+
+    uid => omp_get_uid_from_device(def_dev)
+    call omp_set_default_device (def_dev)
+    dev = -99
+    !$omp target map(from: dev) device(omp_default_device)
+      dev = omp_get_device_num ()
+    !$omp end target
+    if (.not.is_same_dev (def_dev, dev)) &
+      error stop 3
+
+    ! Shall not modify the ICV.  */
+    call omp_set_default_device (omp_default_device)
+    if (def_dev /= omp_get_default_device ()) &
+      error stop 4
+
+    ! Assume the ptr and no only the string is the same.  */
+    if (.not.associated(uid, omp_get_uid_from_device (omp_default_device))) &
+      error stop 5
+  end block
+  end do
+
+  call omp_set_default_device (omp_invalid_device)
+  ! Shall not modify the ICV.
+  call omp_set_default_device (omp_default_device)
+  if (omp_invalid_device /= omp_get_default_device ()) &
+    error stop 6
+
+contains
+
+  logical function is_same_dev (d1, d2)
+    integer, value :: d1, d2
+    integer :: num_dev
+
+    num_dev = omp_get_num_devices ()
+    if (d1 == omp_initial_device) &
+      d1 = num_dev
+    if (d2 == omp_initial_device) &
+      d2 = num_dev
+    is_same_dev = d1 == d2
+  end function is_same_dev
+end program