https://gcc.gnu.org/g:5da963d988e8ea3a13bd8dca87c6cec943af7b56
commit r16-5188-g5da963d988e8ea3a13bd8dca87c6cec943af7b56 Author: Tobias Burnus <[email protected]> Date: Wed Nov 12 10:18:18 2025 +0100 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: --- 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 b2d1e4559883..69cf10caf9db 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 770bc5b1200f..9c9d0c2c5340 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 60cb2b21be74..566a3c28b942 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 ce866c00121a..74e0bfea344c 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 9047095c5e0b..9422515dc379 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 ac5b4b0b7209..002a144b4abb 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 000000000000..5489f019b0b3 --- /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 0412d06f6155..83aba0f6a918 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 504f6caaf07d..9bc02e4b8e61 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 000000000000..28e3496d5470 --- /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
