OpenMP: Move omp requires checks to libgomp

Handle reverse_offload, unified_address, and unified_shared_memory
requirements in libgomp by putting them into the '__offload_requires_mask'
weak variable. Additionally, store the value alongside the offload table
in lto - to permit checking the value for consistency in lto1. The value
is only stored when actually required due to 'omp (declare) target ...'.

In lto1 (either the host one, with -flto [+ ENABLE_OFFLOADING], or in the
offload-device lto1), the consistency check is done, erroring out when
an inconistency is found.

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_declaration_or_fndef): Set
        OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has
        "omp declare target" attribute.
        (c_parser_omp_target_data): Set OMP_REQUIRES_TARGET_USED in
        omp_requires_mask.
        (c_parser_omp_target_enter_data): Likewise.
        (c_parser_omp_target_exit_data): Likewise.
        (c_parser_omp_requires): Remove sorry.

gcc/cp/ChangeLog:

        * parser.cc (cp_parser_simple_declaration): Set
        OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has
        "omp declare target" attribute.
        (cp_parser_omp_target_data): Set OMP_REQUIRES_TARGET_USED in
        omp_requires_mask.
        (cp_parser_omp_target_enter_data): Likewise.
        (cp_parser_omp_target_exit_data): Likewise.
        (cp_parser_omp_requires): Remove sorry.

gcc/fortran/ChangeLog:

        * openmp.cc (gfc_match_omp_requires): Remove "not implemented yet".
        * parse.cc: Include "tree.h" and "omp-general.h".
        (gfc_parse_file): Add code to merge omp_requires to omp_requires_mask.

gcc/ChangeLog:

	* lto-cgraph.cc (output_offload_tables): Output omp_requires_mask,
	but only if OMP_REQUIRES_TARGET_USED.
	(omp_requires_to_name): New.
	(input_offload_tables): Read omp_requires_mask and check whether
	all compilation units use the same value.
	* omp-offload.cc (omp_finish_file): Output omp_requires_mask as
	weak symbol '__offload_requires_mask'.

include/ChangeLog:

	* gomp-constants.h (GOMP_REQUIRES_UNIFIED_ADDRESS,
	GOMP_REQUIRES_UNIFIED_SHARED_MEMORY,
	GOMP_REQUIRES_REVERSE_OFFLOAD): New.

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 (__offload_requires_mask): Declare extern weak var.
	(gomp_requires_to_name): New.
	(gomp_target_init): Pass __offload_requires_mask to
	get_num_devices_func, warn if devices present which do not fulfill
	requirements.
	* 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.

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-2.f90: Likewise. Move post-FE
        checks to ...
        * gfortran.dg/gomp/target-device-ancestor-2a.f90: ... this new file.

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

 gcc/c/c-parser.cc                                  | 21 +++++-
 gcc/cp/parser.cc                                   | 20 +++++-
 gcc/fortran/openmp.cc                              |  4 --
 gcc/fortran/parse.cc                               | 21 ++++++
 gcc/lto-cgraph.cc                                  | 53 +++++++++++++-
 gcc/omp-offload.cc                                 | 21 ++++++
 gcc/testsuite/c-c++-common/gomp/requires-4.c       |  2 -
 .../c-c++-common/gomp/reverse-offload-1.c          |  2 +-
 .../c-c++-common/gomp/target-device-ancestor-2.c   | 10 ++-
 .../c-c++-common/gomp/target-device-ancestor-3.c   |  2 +-
 .../c-c++-common/gomp/target-device-ancestor-4.c   |  4 +-
 .../c-c++-common/gomp/target-device-ancestor-5.c   |  2 +-
 .../gfortran.dg/gomp/target-device-ancestor-2.f90  | 70 ++-----------------
 .../gfortran.dg/gomp/target-device-ancestor-2a.f90 | 80 ++++++++++++++++++++++
 .../gfortran.dg/gomp/target-device-ancestor-3.f90  |  6 +-
 .../gfortran.dg/gomp/target-device-ancestor-4.f90  |  6 +-
 .../gfortran.dg/gomp/target-device-ancestor-5.f90  |  8 +--
 include/gomp-constants.h                           |  6 ++
 libgomp/libgomp-plugin.h                           |  2 +-
 libgomp/oacc-host.c                                |  4 +-
 libgomp/oacc-init.c                                | 16 ++---
 libgomp/plugin/plugin-gcn.c                        |  6 +-
 libgomp/plugin/plugin-nvptx.c                      |  9 ++-
 libgomp/target.c                                   | 40 ++++++++++-
 .../libgomp.c-c++-common/requires-1-aux.c          | 11 +++
 .../testsuite/libgomp.c-c++-common/requires-1.c    | 24 +++++++
 .../libgomp.c-c++-common/requires-2-aux.c          |  9 +++
 .../testsuite/libgomp.c-c++-common/requires-2.c    | 25 +++++++
 .../libgomp.c-c++-common/requires-3-aux.c          | 11 +++
 .../testsuite/libgomp.c-c++-common/requires-3.c    | 24 +++++++
 .../libgomp.c-c++-common/requires-4-aux.c          | 13 ++++
 .../testsuite/libgomp.c-c++-common/requires-4.c    | 23 +++++++
 .../libgomp.c-c++-common/requires-5-aux.c          | 11 +++
 .../testsuite/libgomp.c-c++-common/requires-5.c    | 20 ++++++
 .../testsuite/libgomp.c-c++-common/requires-6.c    | 17 +++++
 liboffloadmic/plugin/libgomp-plugin-intelmic.cpp   |  6 +-
 36 files changed, 493 insertions(+), 116 deletions(-)

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 1704a52be12..4748ce04737 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -2488,6 +2488,12 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
 	  break;
 	}
 
+      if (flag_openmp
+	  && lookup_attribute ("omp declare target",
+			       DECL_ATTRIBUTES (current_function_decl)))
+	omp_requires_mask
+	  = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
       if (DECL_DECLARED_INLINE_P (current_function_decl))
         tv = TV_PARSE_INLINE;
       else
@@ -20915,6 +20921,10 @@ c_parser_omp_teams (location_t loc, c_parser *parser,
 static tree
 c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
 {
+  if (flag_openmp)
+    omp_requires_mask
+      = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				"#pragma omp target data");
@@ -21057,6 +21067,10 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
       return true;
     }
 
+  if (flag_openmp)
+    omp_requires_mask
+      = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
 				"#pragma omp target enter data");
@@ -21143,6 +21157,10 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
       return true;
     }
 
+  if (flag_openmp)
+    omp_requires_mask
+      = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
 				"#pragma omp target exit data");
@@ -22763,9 +22781,6 @@ c_parser_omp_requires (c_parser *parser)
 	      c_parser_skip_to_pragma_eol (parser, false);
 	      return;
 	    }
-	  if (p && this_req != OMP_REQUIRES_DYNAMIC_ALLOCATORS)
-	    sorry_at (cloc, "%qs clause on %<requires%> directive not "
-			    "supported yet", p);
 	  if (p)
 	    c_parser_consume_token (parser);
 	  if (this_req)
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index da2f370cdca..6e26d123370 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -15389,6 +15389,11 @@ cp_parser_simple_declaration (cp_parser* parser,
 	  /* Otherwise, we're done with the list of declarators.  */
 	  else
 	    {
+	      if (flag_openmp && lookup_attribute ("omp declare target",
+						   DECL_ATTRIBUTES (decl)))
+		omp_requires_mask
+		  = (enum omp_requires) (omp_requires_mask
+					 | OMP_REQUIRES_TARGET_USED);
 	      pop_deferring_access_checks ();
 	      cp_finalize_omp_declare_simd (parser, &odsd);
 	      return;
@@ -44287,6 +44292,10 @@ cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok,
 static tree
 cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
 {
+  if (flag_openmp)
+    omp_requires_mask
+      = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				 "#pragma omp target data", pragma_tok);
@@ -44390,6 +44399,10 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
       return true;
     }
 
+  if (flag_openmp)
+    omp_requires_mask
+      = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
 				 "#pragma omp target enter data", pragma_tok);
@@ -44481,6 +44494,10 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
       return true;
     }
 
+  if (flag_openmp)
+    omp_requires_mask
+      = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
 				 "#pragma omp target exit data", pragma_tok);
@@ -46861,9 +46878,6 @@ cp_parser_omp_requires (cp_parser *parser, cp_token *pragma_tok)
 	      cp_parser_skip_to_pragma_eol (parser, pragma_tok);
 	      return false;
 	    }
-	  if (p && this_req != OMP_REQUIRES_DYNAMIC_ALLOCATORS)
-	    sorry_at (cloc, "%qs clause on %<requires%> directive not "
-			    "supported yet", p);
 	  if (p)
 	    cp_lexer_consume_token (parser->lexer);
 	  if (this_req)
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index aeb8a43e12e..a68711081e2 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -5488,10 +5488,6 @@ gfc_match_omp_requires (void)
       else
 	goto error;
 
-      if (requires_clause & ~(OMP_REQ_ATOMIC_MEM_ORDER_MASK
-			      | OMP_REQ_DYNAMIC_ALLOCATORS))
-	gfc_error_now ("Sorry, %qs clause at %L on REQUIRES directive is not "
-		       "yet supported", clause, &old_loc);
       if (!gfc_omp_requires_add_clause (requires_clause, clause, &old_loc, NULL))
 	goto error;
       requires_clauses |= requires_clause;
diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index 7356d1b5a3a..b142e169a5c 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -6908,6 +6908,27 @@ done:
       break;
     }
 
+  if (omp_requires & OMP_REQ_TARGET_MASK)
+    {
+      omp_requires_mask = (enum omp_requires) (omp_requires_mask
+					       | OMP_REQUIRES_TARGET_USED);
+      if (omp_requires & OMP_REQ_REVERSE_OFFLOAD)
+	omp_requires_mask
+	  = (enum omp_requires) (omp_requires_mask
+				 | OMP_REQUIRES_REVERSE_OFFLOAD);
+      if (omp_requires & OMP_REQ_UNIFIED_ADDRESS)
+	omp_requires_mask
+	  = (enum omp_requires) (omp_requires_mask
+				 | OMP_REQUIRES_UNIFIED_ADDRESS);
+      if (omp_requires & OMP_REQ_UNIFIED_SHARED_MEMORY)
+	omp_requires_mask
+	  = (enum omp_requires) (omp_requires_mask
+				 | OMP_REQUIRES_UNIFIED_SHARED_MEMORY);
+    }
+
+  if (omp_requires & OMP_REQ_DYNAMIC_ALLOCATORS)
+    omp_requires_mask
+	= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_DYNAMIC_ALLOCATORS);
   /* Do the parse tree dump.  */
   gfc_current_ns = flag_dump_fortran_original ? gfc_global_ns_list : NULL;
 
diff --git a/gcc/lto-cgraph.cc b/gcc/lto-cgraph.cc
index 237743ef0ba..1b67c4916c4 100644
--- a/gcc/lto-cgraph.cc
+++ b/gcc/lto-cgraph.cc
@@ -37,6 +37,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "pass_manager.h"
 #include "ipa-utils.h"
 #include "omp-offload.h"
+#include "omp-general.h"
 #include "stringpool.h"
 #include "attribs.h"
 #include "alloc-pool.h"
@@ -1068,12 +1069,28 @@ read_string (class lto_input_block *ib)
 void
 output_offload_tables (void)
 {
-  if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars))
+  bool output_requires = (flag_openmp
+			  && (omp_requires_mask & OMP_REQUIRES_TARGET_USED) != 0);
+  if (vec_safe_is_empty (offload_funcs) && vec_safe_is_empty (offload_vars)
+      && !output_requires)
     return;
 
   struct lto_simple_output_block *ob
     = lto_create_simple_output_block (LTO_section_offload_table);
 
+  if (output_requires)
+    {
+      HOST_WIDE_INT val = ((HOST_WIDE_INT) omp_requires_mask
+			   & (OMP_REQUIRES_UNIFIED_ADDRESS
+			      | OMP_REQUIRES_UNIFIED_SHARED_MEMORY
+			      | OMP_REQUIRES_REVERSE_OFFLOAD
+			      | OMP_REQUIRES_TARGET_USED));
+      /* (Mis)use LTO_symtab_edge for this variable.  */
+      streamer_write_enum (ob->main_stream, LTO_symtab_tags,
+			   LTO_symtab_last_tag, LTO_symtab_edge);
+      streamer_write_hwi_stream (ob->main_stream, val);
+    }
+
   for (unsigned i = 0; i < vec_safe_length (offload_funcs); i++)
     {
       symtab_node *node = symtab_node::get ((*offload_funcs)[i]);
@@ -1764,6 +1781,20 @@ input_symtab (void)
     }
 }
 
+static void
+omp_requires_to_name (char *buf, size_t size, unsigned 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 ? "" : ", "));
+}
+
 /* Input function/variable tables that will allow libgomp to look up offload
    target code, and store them into OFFLOAD_FUNCS and OFFLOAD_VARS.  */
 
@@ -1774,6 +1805,8 @@ input_offload_tables (bool do_force_output)
   struct lto_file_decl_data *file_data;
   unsigned int j = 0;
 
+  omp_requires_mask = (omp_requires) 0;
+
   while ((file_data = file_data_vec[j++]))
     {
       const char *data;
@@ -1811,6 +1844,24 @@ input_offload_tables (bool do_force_output)
 	      if (do_force_output)
 		varpool_node::get (var_decl)->force_output = 1;
 	    }
+	  else if (tag == LTO_symtab_edge)
+	    {
+	      static bool error_emitted = false;
+	      HOST_WIDE_INT val = streamer_read_hwi (ib);
+
+	      if (omp_requires_mask == 0)
+		omp_requires_mask = (omp_requires) val;
+	      else if (omp_requires_mask != val && !error_emitted)
+		{
+		  char buf[64], buf2[64];
+		  omp_requires_to_name (buf, sizeof (buf), omp_requires_mask);
+		  omp_requires_to_name (buf2, sizeof (buf2), val);
+		  error ("OpenMP %<requires%> directive with non-identical "
+			 "clauses in multiple compilation units: %qs vs. %qs",
+			 buf, buf2);
+		  error_emitted = true;
+		}
+	    }
 	  else
 	    fatal_error (input_location,
 			 "invalid offload table in %s", file_data->file_name);
diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index 3a89119371c..68e4f6e0993 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -55,6 +55,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "context.h"
 #include "convert.h"
 #include "opts.h"
+#include "varasm.h"
 
 /* Describe the OpenACC looping structure of a function.  The entire
    function is held in a 'NULL' loop.  */
@@ -398,6 +399,26 @@ omp_finish_file (void)
   unsigned num_funcs = vec_safe_length (offload_funcs);
   unsigned num_vars = vec_safe_length (offload_vars);
 
+#ifndef ACCEL_COMPILER
+  if (flag_openmp && (omp_requires_mask & OMP_REQUIRES_TARGET_USED) != 0)
+    {
+      tree var = build_decl (UNKNOWN_LOCATION, VAR_DECL,
+				 get_identifier ("__offload_requires_mask"),
+				 unsigned_type_node);
+      TREE_PUBLIC (var) = 1;
+      TREE_STATIC (var) = 1;
+      TREE_READONLY (var) = 1;
+      DECL_INITIAL (var)
+	= build_int_cst (unsigned_type_node,
+			 ((unsigned int) omp_requires_mask
+			  & (OMP_REQUIRES_UNIFIED_ADDRESS
+			     | OMP_REQUIRES_UNIFIED_SHARED_MEMORY
+			     | OMP_REQUIRES_REVERSE_OFFLOAD)));
+      declare_weak (var);
+      varpool_node::finalize_decl (var);
+    }
+#endif
+
   if (num_funcs == 0 && num_vars == 0)
     return;
 
diff --git a/gcc/testsuite/c-c++-common/gomp/requires-4.c b/gcc/testsuite/c-c++-common/gomp/requires-4.c
index 88ba7746cf8..8f45d83ea6e 100644
--- a/gcc/testsuite/c-c++-common/gomp/requires-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/requires-4.c
@@ -9,5 +9,3 @@ foo (void)
 #pragma omp requires unified_shared_memory	/* { dg-error "'unified_shared_memory' clause used lexically after first target construct or offloading API" } */
 #pragma omp requires unified_address	/* { dg-error "'unified_address' clause used lexically after first target construct or offloading API" } */
 #pragma omp requires reverse_offload	/* { dg-error "'reverse_offload' clause used lexically after first target construct or offloading API" } */
-
-/* { dg-prune-output "not supported yet" } */
diff --git a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
index 9a3fa5230f8..3452156f948 100644
--- a/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/reverse-offload-1.c
@@ -43,7 +43,7 @@ tg_fn (int *x, int *y)
   x2 = x2 + 2 + called_in_target1 ();
   y2 = y2 + 7;
 
-  #pragma omp target device(ancestor : 1) map(tofrom: x2)
+  #pragma omp target device(ancestor : 1) map(tofrom: x2)  /* { dg-message "sorry, unimplemented: 'ancestor' not yet supported" } */
     check_offload(&x2, &y2);
 
   if (x2 != 2+2+3+42 || y2 != 3 + 7)
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
index cf05c505004..b16e701bd5a 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-2.c
@@ -1,13 +1,11 @@
 /* { dg-do compile } */
 
-#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
 
 void
 foo (int n)
 {
-  /* The following test is marked with 'xfail' because a previous 'sorry' from
-     'reverse_offload' suppresses the 'sorry' for 'ancestor'.  */
-  #pragma omp target device (ancestor: 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+  #pragma omp target device (ancestor: 1)
   ;
 
 
@@ -19,9 +17,9 @@ foo (int n)
   #pragma omp target device (ancestor : 42) /* { dg-error "the 'device' clause expression must evaluate to '1'" } */
   ;
 
-  #pragma omp target device (ancestor : n) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+  #pragma omp target device (ancestor : n)
   ;
-  #pragma omp target device (ancestor : n + 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+  #pragma omp target device (ancestor : n + 1)
   ;
 
 
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c
index ea6e5a0cf6c..d16590107d2 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-3.c
@@ -11,7 +11,7 @@ int bar (void);
 
 /* { dg-do compile } */
 
-#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
 
 void
 foo (void)
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c
index b4b5620bbc0..241234f8daf 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-4.c
@@ -4,12 +4,12 @@
   /* Test to ensure that device-modifier 'ancestor' is parsed correctly in
      device clauses. */
 
-#pragma omp requires reverse_offload /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
 
 void
 foo (void)
 {
-  #pragma omp target device (ancestor: 1) /* { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } } */
+  #pragma omp target device (ancestor: 1) /* { dg-message "sorry, unimplemented: 'ancestor' not yet supported" } */
   ;
 
 }
diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
index b6ff84bcdab..b1520ff0636 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-ancestor-5.c
@@ -1,4 +1,4 @@
-#pragma omp requires reverse_offload  /* { dg-message "sorry, unimplemented: 'reverse_offload' clause on 'requires' directive not supported yet" } */
+#pragma omp requires reverse_offload
 
 void
 foo ()
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90
index 117a1d000a5..230c690d84c 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2.f90
@@ -4,19 +4,16 @@ implicit none
 
 integer :: a, b, c
 
-!$omp requires reverse_offload  ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+!$omp requires reverse_offload
 
 
-! The following test case is marked with 'xfail' because a previous 'sorry' from
-! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
-
-!$omp target device (ancestor: 1)  ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp target device (ancestor: 1)
 !$omp end target
 
-!$omp target device (ancestor : a)  ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp target device (ancestor : a)
 !$omp end target
 
-!$omp target device (ancestor : a + 1)  ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp target device (ancestor : a + 1)
 !$omp end target
 
 
@@ -32,61 +29,4 @@ integer :: a, b, c
 !$omp target device (42)
 !$omp end target
 
-
-! Ensure that no OpenMP constructs appear inside target regions with 'ancestor'.
-! The following test case is marked with 'xfail' because a previous 'sorry' from
-! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
-
-!$omp target device (ancestor: 1)
-  !$omp teams  ! { dg-error "" "OpenMP constructs are not allowed in target region with 'ancestor'" { xfail *-*-* } }
-  !$omp end teams
-!$omp end target
-
-!$omp target device (device_num: 1)
-  !$omp teams
-  !$omp end teams
-!$omp end target
-
-!$omp target device (1)
-  !$omp teams
-  !$omp end teams
-!$omp end target
-
-
-! Ensure that with 'ancestor' only the 'device', 'firstprivate', 'private',
-! 'defaultmap', and 'map' clauses appear on the construct.
-! The following test case is marked with 'xfail' because a previous 'sorry' from
-! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
-
-!$omp target nowait device (ancestor: 1)  ! { dg-error "" "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" { xfail *-*-* } }
-!$omp end target
-
-!$omp target device (ancestor: 1) nowait  ! { dg-error "" "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" { xfail *-*-* } }
-!$omp end target
-
-!$omp target nowait device (device_num: 1)
-!$omp end target
-
-!$omp target nowait device (1)
-!$omp end target
-
-!$omp target device (ancestor: 1) firstprivate (a) private (b) defaultmap (none) map (c)
-!$omp end target
-
-
-! Ensure that 'ancestor' is only used with 'target' constructs (not with
-! 'target data', 'target update' etc.).
-! The following test case is marked with 'xfail' because a previous 'sorry' from
-! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
-
-!$omp target data map (a) device (ancestor: 1)  ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
-!$omp end target data
-
-!$omp target enter data map (to: a) device (ancestor: 1)  ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
-!$omp target exit data map (from: a) device (ancestor: 1)  ! { dg-error "" "'device' clause with 'ancestor' is only allowed on 'target' construct" { xfail *-*-* } }
-
-!$omp target update to (a) device (ancestor: 1)  ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" "" { xfail *-*-* } }
-! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" "" { xfail *-*-* } .-1 }
-
-
-end
\ No newline at end of file
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90
new file mode 100644
index 00000000000..feb76fe2144
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-2a.f90
@@ -0,0 +1,80 @@
+! { dg-do compile }
+
+implicit none
+
+integer :: a, b, c
+
+!$omp requires reverse_offload
+
+!$omp target device (ancestor: 1)
+!$omp end target
+
+!$omp target device (ancestor : a)
+!$omp end target
+
+!$omp target device (ancestor : a + 1)
+!$omp end target
+
+
+!$omp target device (device_num:42)
+!$omp end target
+
+!$omp target device (42)
+!$omp end target
+
+
+! Ensure that no OpenMP constructs appear inside target regions with 'ancestor'.
+
+!$omp target device (ancestor: 1)
+  !$omp teams  ! { dg-error "OpenMP constructs are not allowed in target region with 'ancestor'" }
+  !$omp end teams
+!$omp end target
+
+!$omp target device (device_num: 1)
+  !$omp teams
+  !$omp end teams
+!$omp end target
+
+!$omp target device (1)
+  !$omp teams
+  !$omp end teams
+!$omp end target
+
+
+! Ensure that with 'ancestor' only the 'device', 'firstprivate', 'private',
+! 'defaultmap', and 'map' clauses appear on the construct.
+
+!$omp target nowait device (ancestor: 1)  ! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" }
+!$omp end target
+
+!$omp target device (ancestor: 1) nowait  ! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" }
+!$omp end target
+
+!$omp target nowait device (device_num: 1)
+!$omp end target
+
+!$omp target nowait device (1)
+!$omp end target
+
+!$omp target device (ancestor: 1) firstprivate (a) private (b) defaultmap (none) map (c)
+!$omp end target
+
+
+! Ensure that 'ancestor' is only used with 'target' constructs (not with
+! 'target data', 'target update' etc.).
+! The following test case is marked with 'xfail' because a previous 'sorry' from
+! 'reverse_offload' suppresses the 'sorry' for 'ancestor'.
+
+!$omp target data map (a) device (ancestor: 1)  ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+!$omp end target data
+
+!$omp target enter data map (to: a) device (ancestor: 1)  ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+!$omp target exit data map (from: a) device (ancestor: 1)  ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+
+!$omp target update to (a) device (ancestor: 1)  ! { dg-error "'device' clause with 'ancestor' is only allowed on 'target' construct" }
+
+!$omp target device (ancestor: 1) if(.false.)
+! { dg-error "with 'ancestor', only the 'device', 'firstprivate', 'private', 'defaultmap', and 'map' clauses may appear on the construct" "" { target *-*-* } .-1 }
+!$omp end target
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90
index f1145bde2ec..e8975e6a08b 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-3.f90
@@ -16,10 +16,10 @@ subroutine f1 ()
   implicit none
   integer :: n
 
-  !$omp requires reverse_offload  ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+  !$omp requires reverse_offload
 
   !$omp target device (ancestor : 1)
-    n = omp_get_thread_num ()  ! { dg-error "" "OpenMP runtime API call 'omp_get_thread_num' in a region with 'device\\(ancestor\\)' clause" { xfail *-*-* } }
+    n = omp_get_thread_num ()  ! { dg-error "OpenMP runtime API call 'omp_get_thread_num' in a region with 'device\\(ancestor\\)' clause" }
   !$omp end target
 
   !$omp target device (device_num : 1)
@@ -30,4 +30,4 @@ subroutine f1 ()
     n = omp_get_thread_num ()
   !$omp end target
 
-end
\ No newline at end of file
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90
index 63872fa51fb..ab56e2d1d52 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-4.f90
@@ -4,11 +4,11 @@
 ! Test to ensure that device-modifier 'ancestor' is parsed correctly in
 ! device clauses.
 
-!$omp requires reverse_offload  ! { dg-error "Sorry, 'reverse_offload' clause at \\(1\\) on REQUIRES directive is not yet supported" }
+!$omp requires reverse_offload
 
-!$omp target device (ancestor : 1)  ! { dg-message "" "sorry, unimplemented: 'ancestor' not yet supported" { xfail *-*-* } }
+!$omp target device (ancestor : 1)  ! { dg-message "sorry, unimplemented: 'ancestor' not yet supported" }
 !$omp end target
 
 end
 
-! TODO: dg-final { scan-tree-dump-times "pragma omp target \[^\n\r)]*device\\(ancestor:1\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "pragma omp target \[^\n\r)]*device\\(ancestor:1\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
index 06a11eb5092..ca8d4b282a0 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-ancestor-5.f90
@@ -6,7 +6,7 @@
 !
 
 module m
-  !$omp requires reverse_offload  ! { dg-error "REQUIRES directive is not yet supported" }
+  !$omp requires reverse_offload
 contains
   subroutine foo()
     !$omp target device(ancestor:1)
@@ -17,7 +17,7 @@ contains
     block
       block
         block
-          !$omp target device(ancestor:1)
+          !$omp target device(ancestor:1)  ! { dg-message "sorry, unimplemented: 'ancestor' not yet supported" }
           !$omp end target
         end block
       end block
@@ -26,7 +26,7 @@ contains
 end module m
 
 subroutine foo()
-  !$omp requires reverse_offload  ! { dg-error "REQUIRES directive is not yet supported" }
+  !$omp requires reverse_offload
   block
     block
       block
@@ -49,7 +49,7 @@ contains
 end subroutine foo
 
 program main
-  !$omp requires reverse_offload  ! { dg-error "REQUIRES directive is not yet supported" }
+  !$omp requires reverse_offload
 contains
   subroutine foo()
     !$omp target device(ancestor:1)
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index e4dd8ef3e1d..24804aa551f 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -341,6 +341,12 @@ enum gomp_map_kind
 #define GOMP_DEPEND_MUTEXINOUTSET	4
 #define GOMP_DEPEND_INOUTSET		5
 
+/* Flag values for requires-directive features, must match corresponding
+   OMP_REQUIRES_* values in gcc/omp-general.h.  */
+#define GOMP_REQUIRES_UNIFIED_ADDRESS       0x10
+#define GOMP_REQUIRES_UNIFIED_SHARED_MEMORY 0x20
+#define GOMP_REQUIRES_REVERSE_OFFLOAD       0x80
+
 /* HSA specific data structures.  */
 
 /* Identifiers of device-specific target arguments.  */
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 07ab700b80c..ab3ed638475 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/oacc-host.c b/libgomp/oacc-host.c
index 5bb889926d3..eb11b9cf16a 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 1565aa0f290..42c3e74e6ba 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 1c0436842da..ea327bf2ca0 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 387bcbbc52a..bc63e274cdf 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 c0844f2265a..12e6df28198 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,12 @@ static int num_devices;
 /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices.  */
 static int num_devices_openmp;
 
+/* Mask of requires directive clause values.  Offload plugins are queried
+   with this mask to see if all required features are supported.  */
+__attribute__((weak))
+extern unsigned int __offload_requires_mask;
+
+
 /* Similar to gomp_realloc, but release register_lock before gomp_fatal.  */
 
 static void *
@@ -4078,6 +4085,20 @@ gomp_target_fini (void)
     }
 }
 
+static void
+gomp_requires_to_name (char *buf, size_t size, unsigned 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 initializes the runtime for offloading.
    It parses the list of offload plugins, and tries to load these.
    On return, the variables NUM_DEVICES and NUM_DEVICES_OPENMP
@@ -4125,8 +4146,23 @@ 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 requires_mask = 0;
+	    if (&__offload_requires_mask != NULL)
+	      requires_mask = __offload_requires_mask;
+	    new_num_devs
+	      = current_device.get_num_devices_func (requires_mask);
+	    if (new_num_devs < 0)
+	      {
+		char buf[64];
+		gomp_requires_to_name (buf, sizeof (buf), requires_mask);
+		char *name = (char *) malloc (cur_len + 1);
+		memcpy (name, cur, cur_len);
+		name[cur_len] = '\0';
+		GOMP_PLUGIN_error ("note: %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 00000000000..bdca662e42f
--- /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 00000000000..fedf9779769
--- /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 00000000000..617577448ed
--- /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 00000000000..ac7f3ef512c
--- /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 non-identical clauses in multiple compilation units: 'unified_shared_memory' vs. ''" "" { 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 00000000000..bdca662e42f
--- /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 00000000000..4b07ffdd09b
--- /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 00000000000..b8b51ae8ca7
--- /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 00000000000..128fdbb8463
--- /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 00000000000..d223749f0a1
--- /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 00000000000..3d15bde21f0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
@@ -0,0 +1,20 @@
+/* { 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;
+}
+
+/* { dg-output "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 00000000000..b00c7459bbc
--- /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/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
index d1678d0514e..33bae0650b4 100644
--- a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
+++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp
@@ -168,8 +168,12 @@ GOMP_OFFLOAD_get_type (void)
 }
 
 extern "C" int
-GOMP_OFFLOAD_get_num_devices (void)
+GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask)
 {
+  /* Return -1 if no omp_requires_mask cannot be fulfilled but
+     devices were present.  */
+  if (num_devices > 0 && omp_requires_mask != 0)
+    return -1;
   TRACE ("(): return %d", num_devices);
   return num_devices;
 }