OpenMP: Move omp requires checks to libgomp

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

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

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

gcc/c/ChangeLog:

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

gcc/ChangeLog:

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

gcc/cp/ChangeLog:

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

gcc/fortran/ChangeLog:

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

include/ChangeLog:

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

libgomp/ChangeLog:

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

liboffloadmic/ChangeLog:

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

gcc/testsuite/ChangeLog:

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

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

20 files changed, 338 insertions, 110 deletions

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 97e3b23..9c02141 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -20915,6 +20915,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");
@@ -21010,6 +21014,10 @@ c_parser_omp_target_update (location_t loc, c_parser *parser,
       return false;
     }
 
+  if (flag_openmp)
+    omp_requires_mask
+      = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
+
   tree stmt = make_node (OMP_TARGET_UPDATE);
   TREE_TYPE (stmt) = void_type_node;
   OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses;
@@ -21057,6 +21065,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");
@@ -21151,6 +21163,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");
@@ -22779,9 +22795,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/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index ed93ae8..b8b3fec 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -611,6 +611,7 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
   struct regcount *regcounts = XOBFINISH (&regcounts_os, struct regcount *);
 
   fprintf (cfile, "#include <stdlib.h>\n");
+  fprintf (cfile, "#include <stdint.h>\n");
   fprintf (cfile, "#include <stdbool.h>\n\n");
 
   fprintf (cfile, "static const int gcn_num_vars = %d;\n\n", var_count);
@@ -664,7 +665,7 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
 /* Embed an object file into a C source file.  */
 
 static void
-process_obj (FILE *in, FILE *cfile)
+process_obj (FILE *in, FILE *cfile, uint32_t omp_requires)
 {
   size_t len = 0;
   const char *input = read_file (in, &len);
@@ -692,16 +693,18 @@ process_obj (FILE *in, FILE *cfile)
 
   fprintf (cfile,
 	   "static const struct gcn_image_desc {\n"
+	   "  uintptr_t omp_requires_mask;\n"
 	   "  const struct gcn_image *gcn_image;\n"
 	   "  unsigned kernel_count;\n"
 	   "  const struct hsa_kernel_description *kernel_infos;\n"
 	   "  unsigned global_variable_count;\n"
 	   "} target_data = {\n"
+	   "  %d,\n"
 	   "  &gcn_image,\n"
 	   "  sizeof (gcn_kernels) / sizeof (gcn_kernels[0]),\n"
 	   "  gcn_kernels,\n"
 	   "  gcn_num_vars\n"
-	   "};\n\n");
+	   "};\n\n", omp_requires);
 
   fprintf (cfile,
 	   "#ifdef __cplusplus\n"
@@ -1077,9 +1080,27 @@ main (int argc, char **argv)
       unsetenv ("COMPILER_PATH");
       unsetenv ("LIBRARY_PATH");
 
+      char *omp_requires_file;
+      if (save_temps)
+	omp_requires_file = concat (dumppfx, ".mkoffload.omp_requires", NULL);
+      else
+	omp_requires_file = make_temp_file (".mkoffload.omp_requires");
+
       /* Run the compiler pass.  */
+      xputenv (concat ("GCC_OFFLOAD_OMP_REQUIRES_FILE=", omp_requires_file, NULL));
       fork_execute (cc_argv[0], CONST_CAST (char **, cc_argv), true, ".gcc_args");
       obstack_free (&cc_argv_obstack, NULL);
+      unsetenv("GCC_OFFLOAD_OMP_REQUIRES_FILE");
+
+      in = fopen (omp_requires_file, "rb");
+      if (!in)
+	fatal_error (input_location, "cannot open omp_requires file %qs",
+		     omp_requires_file);
+      uint32_t omp_requires;
+      if (fread (&omp_requires, sizeof (omp_requires), 1, in) != 1)
+	fatal_error (input_location, "cannot read omp_requires file %qs",
+		     omp_requires_file);
+      fclose (in);
 
       in = fopen (gcn_s1_name, "r");
       if (!in)
@@ -1102,7 +1123,7 @@ main (int argc, char **argv)
       if (!in)
 	fatal_error (input_location, "cannot open intermediate gcn obj file");
 
-      process_obj (in, cfile);
+      process_obj (in, cfile, omp_requires);
 
       fclose (in);
 
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index b28c1a3..d8c81eb 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -231,7 +231,7 @@ access_check (const char *name, int mode)
 }
 
 static void
-process (FILE *in, FILE *out)
+process (FILE *in, FILE *out, uint32_t omp_requires)
 {
   size_t len = 0;
   const char *input = read_file (in, &len);
@@ -240,6 +240,8 @@ process (FILE *in, FILE *out)
   unsigned obj_count = 0;
   unsigned ix;
 
+  fprintf (out, "#include <stdint.h>\n\n");
+
   /* Dump out char arrays for each PTX object file.  These are
      terminated by a NUL.  */
   for (size_t i = 0; i != len;)
@@ -309,6 +311,7 @@ process (FILE *in, FILE *out)
 
   fprintf (out,
 	   "static const struct nvptx_tdata {\n"
+	   "  uintptr_t omp_requires_mask;\n"
 	   "  const struct ptx_obj *ptx_objs;\n"
 	   "  unsigned ptx_num;\n"
 	   "  const char *const *var_names;\n"
@@ -316,12 +319,12 @@ process (FILE *in, FILE *out)
 	   "  const struct nvptx_fn *fn_names;\n"
 	   "  unsigned fn_num;\n"
 	   "} target_data = {\n"
-	   "  ptx_objs, sizeof (ptx_objs) / sizeof (ptx_objs[0]),\n"
+	   "  %d, ptx_objs, sizeof (ptx_objs) / sizeof (ptx_objs[0]),\n"
 	   "  var_mappings,"
 	   "  sizeof (var_mappings) / sizeof (var_mappings[0]),\n"
 	   "  func_mappings,"
 	   "  sizeof (func_mappings) / sizeof (func_mappings[0])\n"
-	   "};\n\n");
+	   "};\n\n", omp_requires);
 
   fprintf (out, "#ifdef __cplusplus\n"
 	   "extern \"C\" {\n"
@@ -583,19 +586,37 @@ main (int argc, char **argv)
       unsetenv ("COMPILER_PATH");
       unsetenv ("LIBRARY_PATH");
 
+      char *omp_requires_file;
+      if (save_temps)
+	omp_requires_file = concat (dumppfx, ".mkoffload.omp_requires", NULL);
+      else
+	omp_requires_file = make_temp_file (".mkoffload.omp_requires");
+
+      xputenv (concat ("GCC_OFFLOAD_OMP_REQUIRES_FILE=", omp_requires_file, NULL));
       fork_execute (new_argv[0], CONST_CAST (char **, new_argv), true,
 		    ".gcc_args");
       obstack_free (&argv_obstack, NULL);
+      unsetenv("GCC_OFFLOAD_OMP_REQUIRES_FILE");
 
       xputenv (concat ("GCC_EXEC_PREFIX=", execpath, NULL));
       xputenv (concat ("COMPILER_PATH=", cpath, NULL));
       xputenv (concat ("LIBRARY_PATH=", lpath, NULL));
 
+      in = fopen (omp_requires_file, "rb");
+      if (!in)
+	fatal_error (input_location, "cannot open omp_requires file %qs",
+		     omp_requires_file);
+      uint32_t omp_requires;
+      if (fread (&omp_requires, sizeof (omp_requires), 1, in) != 1)
+	fatal_error (input_location, "cannot read omp_requires file %qs",
+		     omp_requires_file);
+      fclose (in);
+
       in = fopen (ptx_name, "r");
       if (!in)
 	fatal_error (input_location, "cannot open intermediate ptx file");
 
-      process (in, out);
+      process (in, out, omp_requires);
       fclose (in);
     }
 
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 6b3763b..df657a3 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -44329,6 +44329,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);
@@ -44432,6 +44436,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);
@@ -44531,6 +44539,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);
@@ -44625,6 +44637,10 @@ cp_parser_omp_target_update (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 stmt = make_node (OMP_TARGET_UPDATE);
   TREE_TYPE (stmt) = void_type_node;
   OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses;
@@ -46919,9 +46935,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 93e40f2..51b429a 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 7356d1b..0b4c596 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -1168,7 +1168,8 @@ decode_omp_directive (void)
     }
   switch (ret)
     {
-    case ST_OMP_DECLARE_TARGET:
+    /* Set omp_target_seen; exclude ST_OMP_DECLARE_TARGET.
+       FIXME: Get clarification, cf. OpenMP Spec Issue #3240.  */
     case ST_OMP_TARGET:
     case ST_OMP_TARGET_DATA:
     case ST_OMP_TARGET_ENTER_DATA:
@@ -6879,11 +6880,14 @@ done:
 
   /* Fixup for external procedures and resolve 'omp requires'.  */
   int omp_requires;
+  bool omp_target_seen;
   omp_requires = 0;
+  omp_target_seen = false;
   for (gfc_current_ns = gfc_global_ns_list; gfc_current_ns;
        gfc_current_ns = gfc_current_ns->sibling)
     {
       omp_requires |= gfc_current_ns->omp_requires;
+      omp_target_seen |= gfc_current_ns->omp_target_seen;
       gfc_check_externals (gfc_current_ns);
     }
   for (gfc_current_ns = gfc_global_ns_list; gfc_current_ns;
@@ -6908,6 +6912,22 @@ done:
       break;
     }
 
+  if (omp_target_seen)
+    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 237743e..4862965 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,7 +1069,10 @@ 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
@@ -1098,6 +1102,19 @@ output_offload_tables (void)
 			       (*offload_vars)[i]);
     }
 
+  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);
+    }
+
   streamer_write_uhwi_stream (ob->main_stream, 0);
   lto_destroy_simple_output_block (ob);
 
@@ -1764,6 +1781,20 @@ input_symtab (void)
     }
 }
 
+static void
+omp_requires_to_name (char *buf, size_t size, HOST_WIDE_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.  */
 
@@ -1773,6 +1804,10 @@ input_offload_tables (bool do_force_output)
   struct lto_file_decl_data **file_data_vec = lto_get_file_decl_data ();
   struct lto_file_decl_data *file_data;
   unsigned int j = 0;
+  const char *requires_fn = NULL;
+  tree requires_decl = NULL_TREE;
+
+  omp_requires_mask = (omp_requires) 0;
 
   while ((file_data = file_data_vec[j++]))
     {
@@ -1784,6 +1819,7 @@ input_offload_tables (bool do_force_output)
       if (!ib)
 	continue;
 
+      tree tmp_decl = NULL_TREE;
       enum LTO_symtab_tags tag
 	= streamer_read_enum (ib, LTO_symtab_tags, LTO_symtab_last_tag);
       while (tag)
@@ -1799,6 +1835,7 @@ input_offload_tables (bool do_force_output)
 		 LTO mode.  */
 	      if (do_force_output)
 		cgraph_node::get (fn_decl)->mark_force_output ();
+	      tmp_decl = fn_decl;
 	    }
 	  else if (tag == LTO_symtab_variable)
 	    {
@@ -1810,6 +1847,72 @@ input_offload_tables (bool do_force_output)
 		 may be no refs to var_decl in offload LTO mode.  */
 	      if (do_force_output)
 		varpool_node::get (var_decl)->force_output = 1;
+	      tmp_decl = var_decl;
+	    }
+	  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;
+		  requires_decl = tmp_decl;
+		  requires_fn = file_data->file_name;
+		}
+	      else if (omp_requires_mask != val && !error_emitted)
+		{
+		  const char *fn1 = requires_fn;
+		  if (requires_decl != NULL_TREE)
+		    {
+		      while (DECL_CONTEXT (requires_decl) != NULL_TREE
+			     && TREE_CODE (requires_decl) != TRANSLATION_UNIT_DECL)
+			requires_decl = DECL_CONTEXT (requires_decl);
+		      if (requires_decl != NULL_TREE)
+			fn1 = IDENTIFIER_POINTER (DECL_NAME (requires_decl));
+		    }
+
+		  const char *fn2 = file_data->file_name;
+		  if (tmp_decl != NULL_TREE)
+		    {
+		      while (DECL_CONTEXT (tmp_decl) != NULL_TREE
+			     && TREE_CODE (tmp_decl) != TRANSLATION_UNIT_DECL)
+			tmp_decl = DECL_CONTEXT (tmp_decl);
+		      if (tmp_decl != NULL_TREE)
+			fn2 = IDENTIFIER_POINTER (DECL_NAME (requires_decl));
+		    }
+
+		  char buf1[sizeof ("unified_address, unified_shared_memory, "
+				    "reverse_offload")];
+		  char buf2[sizeof ("unified_address, unified_shared_memory, "
+				    "reverse_offload")];
+		  omp_requires_to_name (buf2, sizeof (buf2),
+					val != OMP_REQUIRES_TARGET_USED
+					? val
+					: (HOST_WIDE_INT) omp_requires_mask);
+		  if (val != OMP_REQUIRES_TARGET_USED
+		      && omp_requires_mask != OMP_REQUIRES_TARGET_USED)
+		    {
+		      omp_requires_to_name (buf1, sizeof (buf1),
+					    omp_requires_mask);
+		      error ("OpenMP %<requires%> directive with non-identical "
+			     "clauses in multiple compilation units: %qs vs. "
+			     "%qs", buf1, buf2);
+		      inform (UNKNOWN_LOCATION, "%qs has %qs", fn1, buf1);
+		      inform (UNKNOWN_LOCATION, "%qs has %qs", fn2, buf2);
+		    }
+		  else
+		    {
+		      error ("OpenMP %<requires%> directive with %qs specified "
+			     "only in some compilation units", buf2);
+		      inform (UNKNOWN_LOCATION, "%qs has %qs",
+			      val != OMP_REQUIRES_TARGET_USED ? fn2 : fn1,
+			      buf2);
+		      inform (UNKNOWN_LOCATION, "but %qs has not",
+			      val != OMP_REQUIRES_TARGET_USED ? fn1 : fn2);
+		    }
+		  error_emitted = true;
+		}
 	    }
 	  else
 	    fatal_error (input_location,
@@ -1821,6 +1924,18 @@ input_offload_tables (bool do_force_output)
       lto_destroy_simple_input_block (file_data, LTO_section_offload_table,
 				      ib, data, len);
     }
+#ifdef ACCEL_COMPILER
+  char *omp_requires_file = getenv ("GCC_OFFLOAD_OMP_REQUIRES_FILE");
+  if (omp_requires_file == NULL || omp_requires_file[0] == '\0')
+    fatal_error (input_location, "GCC_OFFLOAD_OMP_REQUIRES_FILE unset");
+  FILE *f = fopen (omp_requires_file, "wb");
+  if (!f)
+    fatal_error (input_location, "Cannot open omp_requires file %qs",
+		 omp_requires_file);
+  uint32_t req_mask = omp_requires_mask;
+  fwrite (&req_mask, sizeof (req_mask), 1, f);
+  fclose (f);
+#endif
 }
 
 /* True when we need optimization summary for NODE.  */
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index b9d5529..d73c165 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12701,6 +12701,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       gcc_unreachable ();
     }
 
+  /* Ensure that requires map is written via output_offload_tables, even if only
+     'target (enter/exit) data' is used in the translation unit.  */
+  if (ENABLE_OFFLOADING && (omp_requires_mask & OMP_REQUIRES_TARGET_USED))
+    g->have_offload = true;
+
   clauses = gimple_omp_target_clauses (stmt);
 
   gimple_seq dep_ilist = NULL;
diff --git a/gcc/testsuite/c-c++-common/gomp/requires-4.c b/gcc/testsuite/c-c++-common/gomp/requires-4.c
index 88ba774..8f45d83 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 9a3fa52..3452156 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 cf05c50..b16e701 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 ea6e5a0..d165901 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 b4b5620..241234f 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 b6ff84b..b1520ff 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/requires-8.f90 b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
index e84d609..583c5a5 100644
--- a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90
@@ -1,3 +1,7 @@
+module m0
+  integer :: x
+end module m0
+
 module m  !  { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES UNIFIED_SHARED_MEMORY but other program units do" }
   !$omp requires reverse_offload
 contains
@@ -13,10 +17,14 @@ contains
  end subroutine foo
 end module m
 
-subroutine bar  ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" }
+subroutine bar
   !use m
-  !$omp requires unified_shared_memory
+  !$omp requires unified_shared_memory  ! Possibly OK - needs OpenMP Lang Spec clarification (-> #3240)
   !$omp declare target
 end subroutine bar
 
-! { dg-prune-output "not yet supported" }
+subroutine foobar  ! { dg-error "has OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" }
+  use m0
+  !$omp requires unified_shared_memory
+  !$omp target enter data map(to:x)
+end subroutine foobar
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 117a1d0..230c690 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 0000000..feb76fe
--- /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 f1145bd..e8975e6 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 63872fa..ab56e2d 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 06a11eb..ca8d4b2 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)