Now also with attached patch …
Tobias Burnus wrote:
This is a collateral patch from working on an update for the
implementation status, namely on adding the TR14 (OpenMP 6.1)
additions to
https://gcc.gnu.org/onlinedocs/libgomp/OpenMP-Implementation-Status.html
Actually, I missed this feature myself - most of the time
when calling API routines that take a device number as argument;
'0' feels wrong – and 'omp_get_default_device()' is somewhat
lengthy. In any case, this patch adds this extremely low
hanging fruit.
Comments before I commit it?
Tobias
PS: This only adds omp_default_device; the PR also mentions
some cleanup tasks, but that's deferred to later.
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.
gcc/fortran/intrinsic.texi | 1 +
gcc/fortran/openmp.cc | 6 ++-
libgomp/omp.h.in | 3 +-
libgomp/omp_lib.f90.in | 1 +
libgomp/omp_lib.h.in | 3 +-
libgomp/target.c | 48 +++++++++++++++--
.../libgomp.c-c++-common/omp-default-device.c | 59 +++++++++++++++++++++
libgomp/testsuite/libgomp.c/device_uid.c | 4 +-
libgomp/testsuite/libgomp.fortran/device_uid.f90 | 5 +-
.../libgomp.fortran/omp-default-device.f90 | 61 ++++++++++++++++++++++
10 files changed, 180 insertions(+), 11 deletions(-)
diff --git a/gcc/fortran/intrinsic.texi b/gcc/fortran/intrinsic.texi
index b2d1e455988..69cf10caf9d 100644
--- 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 770bc5b1200..9c9d0c2c534 100644
--- 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 60cb2b21be7..566a3c28b94 100644
--- 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 ce866c00121..74e0bfea344 100644
--- 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 9047095c5e0..9422515dc37 100644
--- 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 ac5b4b0b720..002a144b4ab 100644
--- 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
index 00000000000..5489f019b0b
--- /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 0412d06f615..83aba0f6a91 100644
--- 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 504f6caaf07..9bc02e4b8e6 100644
--- 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
index 00000000000..28e3496d547
--- /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
