OpenMP: Set default-device-var with OMP_TARGET_OFFLOAD=mandatory
OMP_TARGET_OFFLOAD=mandatory handling was before inconsistent. Hence, in
OpenMP 5.2 it was clarified/extended by having implications on the
default-device-var; additionally, omp_initial_device and omp_invalid_device
enum values/PARAMETERs were added; support for it was added
in r13-1066-g1158fe43407568 including aborting for omp_invalid_device and
non-conforming device numbers. Only the mandatory handling was missing.
Namely, while the default-device-var is usually initialized to value 0,
with 'mandatory' it must have the value 'omp_invalid_device' if and only if
zero non-host devices are available. (The OMP_DEFAULT_DEVICE env var
overrides this as it comes semantically after the initialization.)
To achieve this, default-device-var is now initialized to MIN_INT. If
there is no 'mandatory', it is set to 0 directly after env var parsing.
Otherwise, it is updated in gomp_target_init to either 0 or
omp_invalid_device. To ensure INT_MIN is never seen by the user, both
the omp_get_default_device API routine and omp_display_env (user call
and OMP_DISPLAY_ENV env var) call gomp_init_targets_once() in that case.
libgomp/ChangeLog:
* env.c (gomp_default_icv_values): Init default_device_var to
an nonconforming value - INT_MIN.
(initialize_env): After env-var parsing, set default_device_var to
device 0 unless OMP_TARGET_OFFLOAD=mandatory.
(omp_display_env): If default_device_var is INT_MIN, call
gomp_init_targets_once.
* icv-device.c (omp_get_default_device): Likewise.
* libgomp.texi (OMP_DEFAULT_DEVICE): Update init description.
(OpenMP 5.2 Impl. Status): Mark OMP_TARGET_OFFLOAD=mandatory as 'Y'.
* target.c (resolve_device): Improve error message device-num < 0
with 'mandatory' and no no-host devices available.
(gomp_target_init): Set default-device-var if INT_MIN.
* testsuite/libgomp.c/target-48.c: New test.
* testsuite/libgomp.c/target-49.c: New test.
* testsuite/libgomp.c/target-50.c: New test.
* testsuite/libgomp.c/target-51.c: New test.
* testsuite/libgomp.c/target-52.c: New test.
* testsuite/libgomp.c/target-53.c: New test.
* testsuite/libgomp.c/target-54.c: New test.
libgomp/env.c | 13 ++++++++--
libgomp/icv-device.c | 4 +++
libgomp/libgomp.texi | 4 ++-
libgomp/target.c | 15 ++++++++++-
libgomp/testsuite/libgomp.c/target-48.c | 31 +++++++++++++++++++++++
libgomp/testsuite/libgomp.c/target-49.c | 18 +++++++++++++
libgomp/testsuite/libgomp.c/target-50.c | 27 ++++++++++++++++++++
libgomp/testsuite/libgomp.c/target-50a.c | 43 ++++++++++++++++++++++++++++++++
libgomp/testsuite/libgomp.c/target-51.c | 24 ++++++++++++++++++
libgomp/testsuite/libgomp.c/target-52.c | 25 +++++++++++++++++++
libgomp/testsuite/libgomp.c/target-53.c | 22 ++++++++++++++++
libgomp/testsuite/libgomp.c/target-54.c | 20 +++++++++++++++
12 files changed, 242 insertions(+), 4 deletions(-)
@@ -62,13 +62,14 @@
#include "secure_getenv.h"
#include "environ.h"
-/* Default values of ICVs according to the OpenMP standard. */
+/* Default values of ICVs according to the OpenMP standard,
+ except for default-device-var. */
const struct gomp_default_icv gomp_default_icv_values = {
.nthreads_var = 1,
.thread_limit_var = UINT_MAX,
.run_sched_var = GFS_DYNAMIC,
.run_sched_chunk_size = 1,
- .default_device_var = 0,
+ .default_device_var = INT_MIN,
.max_active_levels_var = 1,
.bind_var = omp_proc_bind_false,
.nteams_var = 0,
@@ -1614,6 +1615,10 @@ omp_display_env (int verbose)
struct gomp_icv_list *none
= gomp_get_initial_icv_item (GOMP_DEVICE_NUM_FOR_NO_SUFFIX);
+ if (none->icvs.default_device_var == INT_MIN)
+ /* This implies OMP_TARGET_OFFLOAD=mandatory. */
+ gomp_init_targets_once ();
+
fputs ("\nOPENMP DISPLAY ENVIRONMENT BEGIN\n", stderr);
fputs (" _OPENMP = '201511'\n", stderr);
@@ -2213,6 +2218,10 @@ initialize_env (void)
gomp_global_icv.max_active_levels_var = gomp_supported_active_levels;
}
+ if (gomp_global_icv.default_device_var == INT_MIN
+ && gomp_target_offload_var != GOMP_TARGET_OFFLOAD_MANDATORY)
+ none->icvs.default_device_var = gomp_global_icv.default_device_var = 0;
+
/* Process GOMP_* variables and dependencies between parsed ICVs. */
parse_int_secure ("GOMP_DEBUG", &gomp_debug_var, true);
@@ -27,6 +27,7 @@
expected to replace. */
#include "libgomp.h"
+#include <limits.h>
void
omp_set_default_device (int device_num)
@@ -41,6 +42,9 @@ int
omp_get_default_device (void)
{
struct gomp_task_icv *icv = gomp_icv (false);
+ if (icv->default_device_var == INT_MIN)
+ /* This implies OMP_TARGET_OFFLOAD=mandatory. */
+ gomp_init_targets_once ();
return icv->default_device_var;
}
@@ -423,7 +423,7 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
@item Conforming device numbers and @code{omp_initial_device} and
@code{omp_invalid_device} enum/PARAMETER @tab Y @tab
@item Initial value of @emph{default-device-var} ICV with
- @code{OMP_TARGET_OFFLOAD=mandatory} @tab N @tab
+ @code{OMP_TARGET_OFFLOAD=mandatory} @tab Y @tab
@item @emph{interop_types} in any position of the modifier list for the @code{init} clause
of the @code{interop} construct @tab N @tab
@end multitable
@@ -2006,6 +2006,8 @@ Set to choose the device which is used in a @code{target} region, unless the
value is overridden by @code{omp_set_default_device} or by a @code{device}
clause. The value shall be the nonnegative device number. If no device with
the given device number exists, the code is executed on the host. If unset,
+@env{OMP_TARGET_OFFLOAD} is @code{mandatory} and no non-host devices are
+available, it is set to @code{omp_invalid_device}. Otherwise, if unset,
device number 0 will be used.
@@ -150,7 +150,11 @@ resolve_device (int device_id, bool remapped)
if (device_id == (remapped ? GOMP_DEVICE_HOST_FALLBACK
: omp_initial_device))
return NULL;
- if (device_id == omp_invalid_device)
+ if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY
+ && gomp_get_num_devices () == 0)
+ gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY but only the host "
+ "device is available");
+ else if (device_id == omp_invalid_device)
gomp_fatal ("omp_invalid_device encountered");
else if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY)
gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, "
@@ -5184,6 +5188,15 @@ gomp_target_init (void)
if (devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENACC_200)
goacc_register (&devs[i]);
}
+ if (gomp_global_icv.default_device_var == INT_MIN)
+ {
+ /* This implies OMP_TARGET_OFFLOAD=mandatory. */
+ struct gomp_icv_list *none;
+ none = gomp_get_initial_icv_item (GOMP_DEVICE_NUM_FOR_NO_SUFFIX);
+ gomp_global_icv.default_device_var = (num_devs_openmp
+ ? 0 : omp_invalid_device);
+ none->icvs.default_device_var = gomp_global_icv.default_device_var;
+ }
num_devices = num_devs;
num_devices_openmp = num_devs_openmp;
new file mode 100644
@@ -0,0 +1,31 @@
+/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices;
+ omp_invalid_device == -4 with GCC. */
+
+/* { dg-do run { target { ! offload_device } } } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
+
+/* { dg-output ".*OMP_DEFAULT_DEVICE = '-4'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */
+
+#include <omp.h>
+
+int
+main ()
+{
+ if (omp_get_default_device () != omp_invalid_device)
+ __builtin_abort ();
+
+ omp_set_default_device (omp_initial_device);
+
+ /* The spec is a bit unclear whether the line above sets the device number
+ (a) to -1 (= omp_initial_device) or
+ (b) to omp_get_initial_device() == omp_get_num_devices(). Therefore,
+ we accept either value. */
+
+ if (omp_get_default_device() != omp_get_initial_device()
+ && omp_get_default_device() != omp_initial_device)
+ __builtin_abort ();
+
+ omp_display_env (0);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,18 @@
+/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices,
+ which is enforced by using -foffload=disable. */
+
+/* { dg-do run } */
+/* { dg-additional-options "-foffload=disable" } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
+/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */
+
+/* See comment in target-50.c/target-50.c for why default-device-var can be '0'. */
+
+/* { dg-output ".*OMP_DEFAULT_DEVICE = '-4'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" { target { ! offload_device } } } */
+/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" { target offload_device } } */
+
+int
+main ()
+{
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,27 @@
+/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices;
+ here with using -foffload=disable.
+ As default-device-var is set to 0 (= host in this case), it should not fail. */
+
+/* Note that -foffload=disable will still find devices on the system and only
+ when trying to use them, it will fail as no binary data has been produced.
+ The "target offload_device" case is checked for in 'target-50a.c'. */
+
+/* { dg-do run { target { ! offload_device } } } */
+
+/* { dg-additional-options "-foffload=disable" } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
+/* { dg-set-target-env-var OMP_DEFAULT_DEVICE "0" } */
+/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */
+
+/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */
+
+int
+main ()
+{
+ int x;
+ #pragma omp target map(tofrom:x)
+ x = 5;
+ if (x != 5)
+ __builtin_abort ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,43 @@
+/* Check OMP_TARGET_OFFLOAD on systems with non-host devices but no executable
+ code due to -foffload=disable.
+
+ Note: While one might expect that -foffload=disable implies no non-host
+ devices, libgomp actually detects the devices and only fails when trying to
+ run as no executable code is availale for that device.
+ (Without MANDATORY it simply uses host fallback, which should usually be fine
+ but might have issues in corner cases.)
+
+ We have default-device-var = 0 (default but also explicitly set), which will
+ fail at runtime. For -foffload=disable without non-host devices, see
+ target-50.c testcase. */
+
+/* { dg-do run { target offload_device } } */
+
+/* { dg-additional-options "-foffload=disable" } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
+/* { dg-set-target-env-var OMP_DEFAULT_DEVICE "0" } */
+/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */
+
+/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int x;
+ /* We know that there are non-host devices. With GCC, we still find them as
+ available devices, hence, check for it. */
+ if (omp_get_num_devices() <= 0)
+ __builtin_abort ();
+
+ /* But due to -foffload=disable, there are no binary code for (default) device '0' */
+
+ /* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot be used for offloading.*" } */
+ /* { dg-shouldfail "OMP_TARGET_OFFLOAD=mandatory and no binary code for a non-host device" } */
+ #pragma omp target map(tofrom:x)
+ x = 5;
+ if (x != 5)
+ __builtin_abort ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,24 @@
+/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices,
+ which is enforced by using -foffload=disable. */
+
+/* { dg-do run } */
+/* { dg-additional-options "-foffload=disable" } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
+
+/* { dg-shouldfail "OMP_TARGET_OFFLOAD=mandatory and no available device" } */
+
+/* See comment in target-50.c/target-50.c for why the output differs. */
+
+/* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY but only the host device is available.*" { target { ! offload_device } } } */
+/* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY but device not found.*" { target offload_device } } */
+
+int
+main ()
+{
+ int x;
+ #pragma omp target map(tofrom:x)
+ x = 5;
+ if (x != 5)
+ __builtin_abort ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,25 @@
+/* Only run this with available non-host devices; in that case, GCC sets
+ the default-device-var to 0. */
+
+/* { dg-do run { target { offload_device } } } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
+/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */
+
+/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int x;
+ #pragma omp target map(tofrom:x)
+ x = 5 + omp_is_initial_device ();
+
+ if (x != 5)
+ __builtin_abort ();
+
+ if (0 != omp_get_default_device())
+ __builtin_abort ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,22 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "disabled" } */
+/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */
+
+/* { dg-output ".*OMP_DEFAULT_DEVICE = '\[0-9\]+'.*OMP_TARGET_OFFLOAD = 'DISABLED'.*" } */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int x;
+ #pragma omp target map(tofrom:x)
+ x = 5 + omp_is_initial_device ();
+
+ if (x != 5+1)
+ __builtin_abort ();
+
+ if (omp_get_default_device() != omp_get_initial_device())
+ __builtin_abort ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,20 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "default" } */
+/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */
+
+/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'DEFAULT'.*" } */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int x;
+ #pragma omp target map(tofrom:x)
+ x = 5 + omp_is_initial_device ();
+
+ if (x != 5 + (omp_get_default_device() == omp_get_initial_device()))
+ __builtin_abort ();
+
+ return 0;
+}