OpenMP: Move omp requires checks to libgomp

Handle reverse_offload, unified_address, and unified_shared_memory
requirements in libgomp by saving them alongside the offload table.
When the device lto1 runs, it extracts the data for mkoffload. The
latter than passes the value on to GOMP_offload_register_ver.

lto1 (either the host one, with -flto [+ ENABLE_OFFLOADING], or in the
offload-device lto1) also does the the consistency check is done,
erroring out when the 'omp requires' clause use is inconsistent.

For all in-principle supported devices, if a requirement cannot be fulfilled,
the device is excluded from the (supported) devices list. Currently, none of
those requirements are marked as supported for any of the non-host devices.

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_omp_target_data, c_parser_omp_target_update,
	c_parser_omp_target_enter_data, c_parser_omp_target_exit_data): Set
	OMP_REQUIRES_TARGET_USED.
	(c_parser_omp_requires): Remove sorry.

gcc/ChangeLog:

	* config/gcn/mkoffload.cc (process_asm): Write '#include <stdint.h>'.
	(process_obj): Pass omp_requires_mask to GOMP_offload_register_ver.
	(main): Ask lto1 to obtain omp_requires_mask and pass it on.
	* config/nvptx/mkoffload.cc (process, main): Likewise.
	* lto-cgraph.cc (omp_requires_to_name): New.
	(input_offload_tables): Save omp_requires_mask.
	(output_offload_tables): Read it, check for consistency,
	save value for mkoffload.
	* omp-low.cc (lower_omp_target): Force output_offloadtables
	call for OMP_REQUIRES_TARGET_USED.

gcc/cp/ChangeLog:

	* parser.cc (cp_parser_omp_target_data,
	cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data,
	cp_parser_omp_target_update): Set OMP_REQUIRES_TARGET_USED.
	(cp_parser_omp_requires): Remove sorry.

gcc/fortran/ChangeLog:

	* openmp.cc (gfc_match_omp_requires): Remove sorry.
	* parse.cc (decode_omp_directive): Don't regard 'declare target'
	as target usage for 'omp requires'; add more flags to
	omp_requires_mask.

include/ChangeLog:

	* gomp-constants.h (GOMP_VERSION): Bump to 2.
	(GOMP_REQUIRES_UNIFIED_ADDRESS, GOMP_REQUIRES_UNIFIED_SHARED_MEMORY,
	GOMP_REQUIRES_REVERSE_OFFLOAD, GOMP_REQUIRES_TARGET_USED):
	New defines.

libgomp/ChangeLog:

	* libgomp-plugin.h (GOMP_OFFLOAD_get_num_devices): Add
	omp_requires_mask arg.
	* plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices): Likewise;
	return -1 when device available but omp_requires_mask != 0.
	* plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_num_devices): Likewise.
	* oacc-host.c (host_get_num_devices, host_openacc_get_property):
	Update call.
	* oacc-init.c (resolve_device, acc_init_1, acc_shutdown_1,
	goacc_attach_host_thread_to_device, acc_get_num_devices,
	acc_set_device_num, get_property_any): Likewise.
	* target.c (omp_requires_mask): New global var.
	(gomp_requires_to_name): New.
	(GOMP_offload_register_ver): Handle passed omp_requires_mask.
	(gomp_target_init): Handle omp_requires_mask.
	* libgomp.texi (OpenMP 5.0): Update requires impl. status.
	(OpenMP 5.1): Add a missed item.
	(OpenMP 5.2): Mark linear-clause change as supported in C/C++.
	* testsuite/libgomp.c-c++-common/requires-1-aux.c: New test.
	* testsuite/libgomp.c-c++-common/requires-1.c: New test.
	* testsuite/libgomp.c-c++-common/requires-2-aux.c: New test.
	* testsuite/libgomp.c-c++-common/requires-2.c: New test.
	* testsuite/libgomp.c-c++-common/requires-3-aux.c: New test.
	* testsuite/libgomp.c-c++-common/requires-3.c: New test.
	* testsuite/libgomp.c-c++-common/requires-4-aux.c: New test.
	* testsuite/libgomp.c-c++-common/requires-4.c: New test.
	* testsuite/libgomp.c-c++-common/requires-5-aux.c: New test.
	* testsuite/libgomp.c-c++-common/requires-5.c: New test.
	* testsuite/libgomp.c-c++-common/requires-6.c: New test.
	* testsuite/libgomp.c-c++-common/requires-7-aux.c: New test.
	* testsuite/libgomp.c-c++-common/requires-7.c: New test.
	* testsuite/libgomp.fortran/requires-1-aux.f90: New test.
	* testsuite/libgomp.fortran/requires-1.f90: New test.

liboffloadmic/ChangeLog:

	* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_num_devices):
	Return -1 when device available but omp_requires_mask != 0.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/requires-4.c: Update dg-*.
	* c-c++-common/gomp/reverse-offload-1.c: Likewise.
	* c-c++-common/gomp/target-device-ancestor-2.c: Likewise.
	* c-c++-common/gomp/target-device-ancestor-3.c: Likewise.
	* c-c++-common/gomp/target-device-ancestor-4.c: Likewise.
	* c-c++-common/gomp/target-device-ancestor-5.c: Likewise.
	* gfortran.dg/gomp/target-device-ancestor-3.f90: Likewise.
	* gfortran.dg/gomp/target-device-ancestor-4.f90: Likewise.
	* gfortran.dg/gomp/target-device-ancestor-5.f90: Likewise.
	* gfortran.dg/gomp/target-device-ancestor-2.f90: Likewise. Move
	post-FE checks to ...
	* gfortran.dg/gomp/target-device-ancestor-2a.f90: ... this new file.
	* gfortran.dg/gomp/requires-8.f90: Update as we don't regard
	'declare target' for the 'requires' usage requirement.

Co-authored-by: Chung-Lin Tang <cltang@codesourcery.com>
Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>

22 files changed, 365 insertions, 20 deletions

diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 07ab700..ab3ed63 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -125,7 +125,7 @@ extern void GOMP_PLUGIN_fatal (const char *, ...)
 extern const char *GOMP_OFFLOAD_get_name (void);
 extern unsigned int GOMP_OFFLOAD_get_caps (void);
 extern int GOMP_OFFLOAD_get_type (void);
-extern int GOMP_OFFLOAD_get_num_devices (void);
+extern int GOMP_OFFLOAD_get_num_devices (unsigned int);
 extern bool GOMP_OFFLOAD_init_device (int);
 extern bool GOMP_OFFLOAD_fini_device (int);
 extern unsigned GOMP_OFFLOAD_version (void);
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index a75cd24..39426ff 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -189,8 +189,8 @@ The OpenMP 4.5 specification is fully supported.
       env variable @tab Y @tab
 @item Nested-parallel changes to @emph{max-active-levels-var} ICV @tab Y @tab
 @item @code{requires} directive @tab P
-      @tab Only fulfillable requirement are @code{atomic_default_mem_order}
-      and @code{dynamic_allocators}
+      @tab complete but no non-host devices provides @code{unified_address},
+      @code{unified_shared_memory} or @code{reverse_offload}
 @item @code{teams} construct outside an enclosing target region @tab Y @tab
 @item Non-rectangular loop nests @tab Y @tab
 @item @code{!=} as relational-op in canonical loop form for C/C++ @tab Y @tab
@@ -344,6 +344,8 @@ The OpenMP 4.5 specification is fully supported.
 @item @code{unconstrained} and @code{reproducible} modifiers on @code{order}
       clause @tab Y @tab
 @item Support @code{begin/end declare target} syntax in C/C++ @tab N @tab
+@item Pointer predetermined firstprivate getting initialized
+to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
 @end multitable
 
 
@@ -361,7 +363,7 @@ The OpenMP 4.5 specification is fully supported.
 @item Clauses on @code{end} directive can be on directive @tab N @tab
 @item Deprecation of no-argument @code{destroy} clause on @code{depobj}
       @tab N @tab
-@item @code{linear} clause syntax changes and @code{step} modifier @tab N @tab
+@item @code{linear} clause syntax changes and @code{step} modifier @tab P @tab only C/C++
 @item Deprecation of minus operator for reductions @tab N @tab
 @item Deprecation of separating @code{map} modifiers without comma @tab N @tab
 @item @code{declare mapper} with iterator and @code{present} modifiers
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index 5bb8899..eb11b9c 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -54,7 +54,7 @@ host_get_type (void)
 }
 
 static int
-host_get_num_devices (void)
+host_get_num_devices (unsigned int omp_requires_mask __attribute__((unused)))
 {
   return 1;
 }
@@ -229,7 +229,7 @@ host_openacc_get_property (int n, enum goacc_property prop)
 {
   union goacc_property_value nullval = { .val = 0 };
 
-  if (n >= host_get_num_devices ())
+  if (n >= host_get_num_devices (0))
     return nullval;
 
   switch (prop)
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index 1565aa0..42c3e74e 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -148,7 +148,7 @@ resolve_device (acc_device_t d, bool fail_is_error)
 	      if (dispatchers[d]
 		  && !strcasecmp (goacc_device_type,
 				  get_openacc_name (dispatchers[d]->name))
-		  && dispatchers[d]->get_num_devices_func () > 0)
+		  && dispatchers[d]->get_num_devices_func (0) > 0)
 		goto found;
 
 	    if (fail_is_error)
@@ -169,7 +169,7 @@ resolve_device (acc_device_t d, bool fail_is_error)
     case acc_device_not_host:
       /* Find the first available device after acc_device_not_host.  */
       while (known_device_type_p (++d))
-	if (dispatchers[d] && dispatchers[d]->get_num_devices_func () > 0)
+	if (dispatchers[d] && dispatchers[d]->get_num_devices_func (0) > 0)
 	  goto found;
       if (d_arg == acc_device_default)
 	{
@@ -302,7 +302,7 @@ acc_init_1 (acc_device_t d, acc_construct_t parent_construct, int implicit)
 
   base_dev = resolve_device (d, true);
 
-  ndevs = base_dev->get_num_devices_func ();
+  ndevs = base_dev->get_num_devices_func (0);
 
   if (ndevs <= 0 || goacc_device_num >= ndevs)
     acc_dev_num_out_of_range (d, goacc_device_num, ndevs);
@@ -351,7 +351,7 @@ acc_shutdown_1 (acc_device_t d)
   /* Get the base device for this device type.  */
   base_dev = resolve_device (d, true);
 
-  ndevs = base_dev->get_num_devices_func ();
+  ndevs = base_dev->get_num_devices_func (0);
 
   /* Unload all the devices of this type that have been opened.  */
   for (i = 0; i < ndevs; i++)
@@ -520,7 +520,7 @@ goacc_attach_host_thread_to_device (int ord)
       base_dev = cached_base_dev;
     }
   
-  num_devices = base_dev->get_num_devices_func ();
+  num_devices = base_dev->get_num_devices_func (0);
   if (num_devices <= 0 || ord >= num_devices)
     acc_dev_num_out_of_range (acc_device_type (base_dev->type), ord,
 			      num_devices);
@@ -599,7 +599,7 @@ acc_get_num_devices (acc_device_t d)
   if (!acc_dev)
     return 0;
 
-  n = acc_dev->get_num_devices_func ();
+  n = acc_dev->get_num_devices_func (0);
   if (n < 0)
     n = 0;
 
@@ -779,7 +779,7 @@ acc_set_device_num (int ord, acc_device_t d)
 
       cached_base_dev = base_dev = resolve_device (d, true);
 
-      num_devices = base_dev->get_num_devices_func ();
+      num_devices = base_dev->get_num_devices_func (0);
 
       if (num_devices <= 0 || ord >= num_devices)
         acc_dev_num_out_of_range (d, ord, num_devices);
@@ -814,7 +814,7 @@ get_property_any (int ord, acc_device_t d, acc_device_property_t prop)
 
   struct gomp_device_descr *dev = resolve_device (d, true);
 
-  int num_devices = dev->get_num_devices_func ();
+  int num_devices = dev->get_num_devices_func (0);
 
   if (num_devices <= 0 || ord >= num_devices)
     acc_dev_num_out_of_range (d, ord, num_devices);
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index 1c04368..ea327bf 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -3221,10 +3221,14 @@ GOMP_OFFLOAD_version (void)
 /* Return the number of GCN devices on the system.  */
 
 int
-GOMP_OFFLOAD_get_num_devices (void)
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
 {
   if (!init_hsa_context ())
     return 0;
+  /* Return -1 if no omp_requires_mask cannot be fulfilled but
+     devices were present.  */
+  if (hsa_context.agent_count > 0 && omp_requires_mask != 0)
+    return -1;
   return hsa_context.agent_count;
 }
 
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 387bcbb..bc63e274 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1175,9 +1175,14 @@ GOMP_OFFLOAD_get_type (void)
 }
 
 int
-GOMP_OFFLOAD_get_num_devices (void)
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
 {
-  return nvptx_get_num_devices ();
+  int num_devices = nvptx_get_num_devices ();
+  /* Return -1 if no omp_requires_mask cannot be fulfilled but
+     devices were present.  */
+  if (num_devices > 0 && omp_requires_mask != 0)
+    return -1;
+  return num_devices;
 }
 
 bool
diff --git a/libgomp/target.c b/libgomp/target.c
index c0844f2..4dac818 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -36,6 +36,7 @@
 # include <inttypes.h>  /* For PRIu64.  */
 #endif
 #include <string.h>
+#include <stdio.h>  /* For snprintf. */
 #include <assert.h>
 #include <errno.h>
 
@@ -98,6 +99,9 @@ static int num_devices;
 /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices.  */
 static int num_devices_openmp;
 
+/* OpenMP requires mask.  */
+static int omp_requires_mask;
+
 /* Similar to gomp_realloc, but release register_lock before gomp_fatal.  */
 
 static void *
@@ -2314,6 +2318,20 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep,
     }
 }
 
+static void
+gomp_requires_to_name (char *buf, size_t size, int requires_mask)
+{
+  char *end = buf + size, *p = buf;
+  if (requires_mask & GOMP_REQUIRES_UNIFIED_ADDRESS)
+    p += snprintf (p, end - p, "unified_address");
+  if (requires_mask & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY)
+    p += snprintf (p, end - p, "%sunified_shared_memory",
+		   (p == buf ? "" : ", "));
+  if (requires_mask & GOMP_REQUIRES_REVERSE_OFFLOAD)
+    p += snprintf (p, end - p, "%sreverse_offload",
+		   (p == buf ? "" : ", "));
+}
+
 /* This function should be called from every offload image while loading.
    It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
    the target, and TARGET_DATA needed by target plugin.  */
@@ -2323,13 +2341,43 @@ GOMP_offload_register_ver (unsigned version, const void *host_table,
 			   int target_type, const void *target_data)
 {
   int i;
+  int omp_req = 0;
 
   if (GOMP_VERSION_LIB (version) > GOMP_VERSION)
     gomp_fatal ("Library too old for offload (version %u < %u)",
 		GOMP_VERSION, GOMP_VERSION_LIB (version));
-  
+
+  if (GOMP_VERSION_LIB (version) > 1)
+    {
+      omp_req = (int) (size_t) ((void **) target_data)[0];
+      target_data = &((void **) target_data)[1];
+    }
+
   gomp_mutex_lock (&register_lock);
 
+  if (omp_req && omp_requires_mask && omp_requires_mask != omp_req)
+    {
+      char buf1[sizeof ("unified_address, unified_shared_memory, "
+			"reverse_offload")];
+      char buf2[sizeof ("unified_address, unified_shared_memory, "
+			"reverse_offload")];
+      gomp_requires_to_name (buf2, sizeof (buf2),
+			     omp_req != GOMP_REQUIRES_TARGET_USED
+			     ? omp_req : omp_requires_mask);
+      if (omp_req != GOMP_REQUIRES_TARGET_USED
+	  && omp_requires_mask != GOMP_REQUIRES_TARGET_USED)
+	{
+	  gomp_requires_to_name (buf1, sizeof (buf1), omp_requires_mask);
+	  gomp_fatal ("OpenMP 'requires' directive with non-identical clauses "
+		      "in multiple compilation units: '%s' vs. '%s'",
+		      buf1, buf2);
+	}
+      else
+	gomp_fatal ("OpenMP 'requires' directive with '%s' specified only in "
+		    "some compilation units", buf2);
+    }
+  omp_requires_mask = omp_req;
+
   /* Load image to all initialized devices.  */
   for (i = 0; i < num_devices; i++)
     {
@@ -4125,8 +4173,30 @@ gomp_target_init (void)
 
 	if (gomp_load_plugin_for_device (&current_device, plugin_name))
 	  {
-	    new_num_devs = current_device.get_num_devices_func ();
-	    if (new_num_devs >= 1)
+	    int omp_req = omp_requires_mask & ~GOMP_REQUIRES_TARGET_USED;
+	    new_num_devs = current_device.get_num_devices_func (omp_req);
+	    if (gomp_debug_var > 0 && new_num_devs < 0)
+	      {
+		bool found = false;
+		int type = current_device.get_type_func ();
+		for (int img = 0; img < num_offload_images; img++)
+		  if (type == offload_images[img].type)
+		    found = true;
+		if (found)
+		  {
+		    char buf[sizeof ("unified_address, unified_shared_memory, "
+				     "reverse_offload")];
+		    gomp_requires_to_name (buf, sizeof (buf), omp_req);
+		    char *name = (char *) malloc (cur_len + 1);
+		    memcpy (name, cur, cur_len);
+		    name[cur_len] = '\0';
+		    gomp_debug (1,
+				"%s devices present but 'omp requires %s' "
+			        "cannot be fulfilled", name, buf);
+		    free (name);
+		  }
+	      }
+	    else if (new_num_devs >= 1)
 	      {
 		/* Augment DEVICES and NUM_DEVICES.  */
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c
new file mode 100644
index 0000000..bdca662
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+  #pragma omp target
+  x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-1.c b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c
new file mode 100644
index 0000000..fedf977
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c
@@ -0,0 +1,24 @@
+/* { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } } */
+/* { dg-additional-sources requires-1-aux.c } */
+
+/* Check diagnostic by device-compiler's lto1.
+   Other file uses: 'requires unified_address'.  */
+
+#pragma omp requires unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+  #pragma omp target
+  for (int i = 0; i < 10; i++)
+    a[i] = 0;
+
+  foo ();
+  return 0;
+}
+
+/* { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 }  */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c
new file mode 100644
index 0000000..6175774
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c
@@ -0,0 +1,9 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+int x;
+
+void foo (void)
+{
+  #pragma omp target
+  x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-2.c b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c
new file mode 100644
index 0000000..be1830d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c
@@ -0,0 +1,25 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-options "-foffload=disable -flto" } */
+/* { dg-additional-sources requires-2-aux.c } */
+
+/* Check diagnostic by host's lto1.
+   Other file does not have any 'omp requires'. */
+
+#pragma omp requires unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+  #pragma omp target
+  for (int i = 0; i < 10; i++)
+    a[i] = 0;
+
+  foo ();
+  return 0;
+}
+
+/* { dg-error "OpenMP 'requires' directive with 'unified_shared_memory' specified only in some compilation units" "" { target *-*-* } 0 }  */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c
new file mode 100644
index 0000000..bdca662
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-3-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+  #pragma omp target
+  x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-3.c b/libgomp/testsuite/libgomp.c-c++-common/requires-3.c
new file mode 100644
index 0000000..4b07ffd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-3.c
@@ -0,0 +1,24 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-sources requires-3-aux.c } */
+
+/* Check diagnostic by device-compiler's lto1.
+   Other file uses: 'requires unified_address'.  */
+
+#pragma omp requires unified_address,unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+  #pragma omp target
+  for (int i = 0; i < 10; i++)
+    a[i] = 0;
+
+  foo ();
+  return 0;
+}
+
+/* { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_address, unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 }  */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c
new file mode 100644
index 0000000..b8b51ae
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4-aux.c
@@ -0,0 +1,13 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires reverse_offload
+
+/* Note: The file does not have neither of:
+   declare target directives, device constructs or device routines.  */
+
+int x;
+
+void foo (void)
+{
+  x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
new file mode 100644
index 0000000..128fdbb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
@@ -0,0 +1,23 @@
+/* { dg-do link { target offloading_enabled } } */
+/* { dg-additional-options "-flto" } */
+/* { dg-additional-sources requires-4-aux.c } */
+
+/* Check diagnostic by device-compiler's or host compiler's lto1.
+   Other file uses: 'requires reverse_offload', but that's inactive as
+   there are no declare target directives, device constructs nor device routines  */
+
+#pragma omp requires unified_address,unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+  #pragma omp target
+  for (int i = 0; i < 10; i++)
+    a[i] = 0;
+
+  foo ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c
new file mode 100644
index 0000000..d223749
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+int x;
+
+void foo (void)
+{
+  #pragma omp target
+  x = 1;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
new file mode 100644
index 0000000..c1e5540
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
@@ -0,0 +1,21 @@
+/* { dg-do run { target { offload_target_nvptx || offload_target_amdgcn } } } */
+/* { dg-additional-sources requires-5-aux.c } */
+
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+  #pragma omp target
+  for (int i = 0; i < 10; i++)
+    a[i] = 0;
+
+  foo ();
+  return 0;
+}
+
+/* (Only) if GOMP_DEBUG=1, should print at runtime the following:
+   "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled" */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
new file mode 100644
index 0000000..b00c745
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
@@ -0,0 +1,17 @@
+#pragma omp requires unified_shared_memory, unified_address, reverse_offload
+
+/* The requires line is not active as there is none of:
+     declare target directives, device constructs or device routines.
+   Thus, this code is expected to work everywhere.  */
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+  for (int i = 0; i < 10; i++)
+    a[i] = 0;
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c b/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c
new file mode 100644
index 0000000..0916db8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-7-aux.c
@@ -0,0 +1,11 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires unified_address
+
+int x;
+
+void foo (void)
+{
+  x = 1;
+  #pragma omp target enter data map(always,to: x)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-7.c b/libgomp/testsuite/libgomp.c-c++-common/requires-7.c
new file mode 100644
index 0000000..c94a4c1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-7.c
@@ -0,0 +1,24 @@
+/* { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } } */
+/* { dg-additional-sources requires-7-aux.c } */
+
+/* Check diagnostic by device-compiler's lto1.
+   Other file uses: 'requires unified_address'.  */
+
+#pragma omp requires unified_shared_memory
+
+int a[10];
+extern void foo (void);
+
+int
+main (void)
+{
+  #pragma omp target
+  for (int i = 0; i < 10; i++)
+    a[i] = 0;
+
+  foo ();
+  return 0;
+}
+
+/* { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 }  */
+/* { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" } */
diff --git a/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90 b/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90
new file mode 100644
index 0000000..a18caeb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/requires-1-aux.f90
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } }
+
+module m
+  integer x
+end module m
+
+subroutine foo
+  use m
+  implicit none
+  !$omp requires unified_address
+
+  x = 1
+  !$omp target enter data map(always,to: x)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/requires-1.f90 b/libgomp/testsuite/libgomp.fortran/requires-1.f90
new file mode 100644
index 0000000..33741af
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/requires-1.f90
@@ -0,0 +1,26 @@
+! { dg-do link { target { offload_target_nvptx || offload_target_amdgcn } } }
+! { dg-additional-sources requires-1-aux.f90 }
+
+! Check diagnostic by device-compiler's lto1.
+!   Other file uses: 'requires unified_address'.
+
+module m
+  integer :: a(10)
+  interface
+    subroutine foo
+    end
+  end interface
+end
+
+program main
+  !$omp requires unified_shared_memory
+
+  !$omp target
+    a = 0
+  !$omp end target
+
+  call foo ()
+end
+
+! { dg-error "OpenMP 'requires' directive with non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. 'unified_address'" "" { target *-*-* } 0 }
+! { dg-excess-errors "Ignore messages like: errors during merging of translation units|mkoffload returned 1 exit status" }