Re: [Patch][v4] OpenMP: Move omp requires checks to libgomp

[Tobias Burnus <tobias@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 3986 bytes --]

Hi Jakub, hi all,

new version attached. It now checks during lto1 whether the values are
consistent – and fails with a hard error.

The actually used value (by libgomp) is stored as a scalar weak symbol –
while for checking, each translation unit stores the integer value for
lto (alongside the offload table). This is both used for checking and in
lto1 (device + host lto1), to restore the value of 'omp_requires_mask'
for further use. – Currently, it is only used on the host to make the
value available to libgomp. – However, a device lto1 could also use it.
(Usage: cf. Andrew's USM gcn patch.)

Unchanged from previous version, libgomp outputs an warning/note if a
device could be found but the requires prevented libgomp from using it.
This message is also shown with -foffload=disable but it is not shown
for OMP_TARGET_OFFLOAD=disable.

Other change is that API calls no longer count as relevant for 'omp
requires' – such that compilation units which only contain those will
not output anything (independent whether there is an 'omp requires' or not.)

On 09.06.22 16:19, Jakub Jelinek wrote:
> On Thu, Jun 09, 2022 at 02:46:34PM +0200, Tobias Burnus wrote:
>> On 09.06.22 13:40, Jakub Jelinek via Gcc-patches wrote:
> If it is from me, bet it was because of that (mis)understanding that
> device routines are device related runtime API calls.
> I'd suggest to only mark in the patch what is clear (which is device
> constructs) and defer the rest until it is clarified.
Done so.
>>> Shouldn't the vars in that section be const, so that it is a read-only
>>> section?
>>>
>>> Is unsigned_type_node what we want (say wouldn't be just unsigned_char_node
>>> be enough, currently we just need 3 bits).
>> Probably -that would be 8 bits, leaving 5 spare. I have not checked what
>> Andrew et al. do with the pinned-memory support by -f<some-flag>, but
>> that will likely use only 1 to 3 bits, if any.
> If it is SHF_MERGE, even 16-bit or 32-bit wouldn't be the end of the world,
> or if it is in LTO streamed out stuff, we can use a bitpack for it...

As the final binary will only contain a single variable, the size should
not matter much. I currently use 'unsigned' but it could surely be
shorter.  For the .o files, it also outputs the unsigned value for each
TU, but that's also small. I was thinking about adding more data (like
location data, be it location_t or __FILENAME__). However, it uses a
stripped-down stream writer - and to do so, location/string writing
requires a different object (and reading it, data_in). I did not regard
this as worthwhile and, thus, I only output the used requires clauses
and not where they were used.

> I think best would be a fatal error if people try to configure
> offloading targets for a compiler that doesn't support named sections,
> or perhaps that and presence of anything that should be offloaded.

I do not use any named section – but I could if it makes sense. In any
case, the question is whether the current weak symbol makes sense or
not. And whether there are problems in using weak symbols (in libgomp's
target.c + for non-ACCEL_COMPILER, but only when the symbol needs to be
written). I am also not sure about the best naming. – Thoughts?

Otherwise, tested with no offloading configured + with offloading to
nvptx (fully tested) and gcn (partially) [all x86_64-gnu-linux)

Tobias

PS: At some point, we need to think about handling calling from a
program's target region a declare-target device function which is inside
a shared library. I am sure, we currently do not handle it. – For that,
we then also have to think about how to handle the requires clauses.
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Attachment #2: omp-requires-v4.diff --]
[-- Type: text/x-patch, Size: 49650 bytes --]

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