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-50a.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

@@ -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);
 

libgomp/icv-device.c

@@ -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;
 }
 

libgomp/libgomp.texi

@@ -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.
 
 

libgomp/target.c

@@ -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;

libgomp/testsuite/libgomp.c/target-48.c

@@ -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;
+}

libgomp/testsuite/libgomp.c/target-49.c

@@ -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;
+}

libgomp/testsuite/libgomp.c/target-50.c

@@ -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;
+}

libgomp/testsuite/libgomp.c/target-50a.c

@@ -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;
+}

libgomp/testsuite/libgomp.c/target-51.c

@@ -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;
+}

libgomp/testsuite/libgomp.c/target-52.c

@@ -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;
+}

libgomp/testsuite/libgomp.c/target-53.c

@@ -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;
+}

libgomp/testsuite/libgomp.c/target-54.c

@@ -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;
+}