OpenMP: Fortran support for metadirectives and dynamic selectors

gcc/fortran/ChangeLog
	PR middle-end/112779
	PR middle-end/113904
	* decl.cc (gfc_match_end): Handle COMP_OMP_BEGIN_METADIRECTIVE and
	COMP_OMP_METADIRECTIVE.
	* dump-parse-tree.cc (show_omp_node): Handle EXEC_OMP_METADIRECTIVE.
	(show_code_node): Likewise.
	* gfortran.h (enum gfc_statement): Add ST_OMP_METADIRECTIVE,
	ST_OMP_BEGIN_METADIRECTIVE, and ST_OMP_END_METADIRECTIVE.
	(struct gfc_omp_clauses): Rename target_first_st_is_teams to
	target_first_st_is_teams_or_meta.
	(struct gfc_omp_variant): New.
	(gfc_get_omp_variant): New.
	(struct gfc_st_label): Add omp_region field.
	(enum gfc_exec_op): Add EXEC_OMP_METADIRECTIVE.
	(struct gfc_code): Add omp_variants fields.
	(gfc_free_omp_variants): Declare.
	(match_omp_directive): Declare.
	(is_omp_declarative_stmt): Declare.
	* io.cc (format_asterisk): Adjust initializer.
	* match.h (gfc_match_omp_begin_metadirective): Declare.
	(gfc_match_omp_metadirective): Declare.
	* openmp.cc (gfc_omp_directives): Uncomment metadirective.
	(gfc_match_omp_eos): Adjust to match context selectors.
	(gfc_free_omp_variants): New.
	(gfc_match_omp_clauses): Remove context_selector parameter and adjust
	to use gfc_match_omp_eos instead.
	(match_omp): Adjust call to gfc_match_omp_clauses.
	(gfc_match_omp_context_selector): Add metadirective_p parameter and
	adjust error-checking.  Adjust matching of simd clauses.
	(gfc_match_omp_context_selector_specification): Adjust parameters
	so it can be used for metadirective as well as declare variant.
	(match_omp_metadirective): New.
	(gfc_match_omp_begin_metadirective): New.
	(gfc_match_omp_metadirective): New.
	(resolve_omp_metadirective): New.
	(resolve_omp_target): Handle metadirectives.
	(gfc_resolve_omp_directive): Handle EXEC_OMP_METADIRECTIVE.
	* parse.cc (gfc_matching_omp_context_selector): New.
	(gfc_in_omp_metadirective_body): New.
	(gfc_omp_region_count): New.
	(decode_omp_directive): Handle ST_OMP_BEGIN_METADIRECTIVE and
	ST_OMP_METADIRECTIVE.
	(match_omp_directive): New.
	(case_omp_structured_block): Define.
	(case_omp_do): Define.
	(gfc_ascii_statement): Handle ST_OMP_BEGIN_METADIRECTIVE,
	ST_OMP_END_METADIRECTIVE, and ST_OMP_METADIRECTIVE.
	(accept_statement):  Handle ST_OMP_METADIRECTIVE and
	ST_OMP_BEGIN_METADIRECTIVE.
	(gfc_omp_end_stmt): New, split from...
	(parse_omp_do): ...here, and...
	(parse_omp_structured_block): ...here.  Handle metadirectives,
	plus "allocate", "atomic", and "dispatch" which were missing.
	(parse_omp_oacc_atomic): Handle "end metadirective".
	(parse_openmp_allocate_block): Likewise.
	(parse_omp_dispatch): Likewise.
	(parse_omp_metadirective_body): New.
	(parse_executable): Handle metadirective.  Use new case macros
	defined above.
	(gfc_parse_file): Initialize metadirective state.
	(is_omp_declarative_stmt): New.
	* parse.h (enum gfc_compile_state): Add COMP_OMP_METADIRECTIVE
	and COMP_OMP_BEGIN_METADIRECTIVE.
	(gfc_omp_end_stmt): Declare.
	(gfc_matching_omp_context_selector): Declare.
	(gfc_in_omp_metadirective_body): Declare.
	(gfc_omp_metadirective_region_count): Declare.
	* resolve.cc (gfc_resolve_code): Handle EXEC_OMP_METADIRECTIVE.
	* st.cc (gfc_free_statement): Likewise.
	* symbol.cc (compare_st_labels): Handle labels within a metadirective
	body.
	(gfc_get_st_label): Likewise.
	* trans-decl.cc (gfc_get_label_decl): Encode the metadirective region
	in the label_name.
	* trans-openmp.cc (gfc_trans_omp_directive): Handle
	EXEC_OMP_METADIRECTIVE.
	(gfc_trans_omp_set_selector): New, split/adapted from code....
	(gfc_trans_omp_declare_variant): ...here.
	(gfc_trans_omp_metadirective): New.
	* trans-stmt.h 	(gfc_trans_omp_metadirective): Declare.
	* trans.cc (trans_code): Handle EXEC_OMP_METADIRECTIVE.

gcc/testsuite/ChangeLog
	PR middle-end/112779
	PR middle-end/113904
	* gfortran.dg/gomp/metadirective-1.f90: New.
	* gfortran.dg/gomp/metadirective-10.f90: New.
	* gfortran.dg/gomp/metadirective-11.f90: New.
	* gfortran.dg/gomp/metadirective-12.f90: New.
	* gfortran.dg/gomp/metadirective-13.f90: New.
	* gfortran.dg/gomp/metadirective-2.f90: New.
	* gfortran.dg/gomp/metadirective-3.f90: New.
	* gfortran.dg/gomp/metadirective-4.f90: New.
	* gfortran.dg/gomp/metadirective-5.f90: New.
	* gfortran.dg/gomp/metadirective-6.f90: New.
	* gfortran.dg/gomp/metadirective-7.f90: New.
	* gfortran.dg/gomp/metadirective-8.f90: New.
	* gfortran.dg/gomp/metadirective-9.f90: New.
	* gfortran.dg/gomp/metadirective-construct.f90: New.
	* gfortran.dg/gomp/metadirective-no-score.f90: New.
	* gfortran.dg/gomp/pure-1.f90 (func_metadirective): New.
	(func_metadirective_2): New.
	(func_metadirective_3): New.
	* gfortran.dg/gomp/pure-2.f90 (func_metadirective): Delete.

libgomp/ChangeLog
	PR middle-end/112779
	PR middle-end/113904
	* testsuite/libgomp.fortran/metadirective-1.f90: New.
	* testsuite/libgomp.fortran/metadirective-2.f90: New.
	* testsuite/libgomp.fortran/metadirective-3.f90: New.
	* testsuite/libgomp.fortran/metadirective-4.f90: New.
	* testsuite/libgomp.fortran/metadirective-5.f90: New.
	* testsuite/libgomp.fortran/metadirective-6.f90: New.

Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
Co-Authored-By: Paul-Antoine Arras <pa@codesourcery.com>

38 files changed, 2101 insertions, 370 deletions

diff --git a/gcc/fortran/decl.cc b/gcc/fortran/decl.cc
index 0c59760..7954a84 100644
--- a/gcc/fortran/decl.cc
+++ b/gcc/fortran/decl.cc
@@ -8457,6 +8457,7 @@ gfc_match_end (gfc_statement *st)
 
     case COMP_CONTAINS:
     case COMP_DERIVED_CONTAINS:
+    case COMP_OMP_BEGIN_METADIRECTIVE:
       state = gfc_state_stack->previous->state;
       block_name = gfc_state_stack->previous->sym == NULL
 		 ? NULL : gfc_state_stack->previous->sym->name;
@@ -8464,6 +8465,28 @@ gfc_match_end (gfc_statement *st)
 		&& gfc_state_stack->previous->sym->abr_modproc_decl;
       break;
 
+    case COMP_OMP_METADIRECTIVE:
+      {
+	/* Metadirectives can be nested, so we need to drill down to the
+	   first state that is not COMP_OMP_METADIRECTIVE.  */
+	gfc_state_data *state_data = gfc_state_stack;
+
+	do
+	  {
+	    state_data = state_data->previous;
+	    state = state_data->state;
+	    block_name = (state_data->sym == NULL
+			  ? NULL : state_data->sym->name);
+	    abbreviated_modproc_decl = (state_data->sym
+					&& state_data->sym->abr_modproc_decl);
+	  }
+	while (state == COMP_OMP_METADIRECTIVE);
+
+	if (block_name && startswith (block_name, "block@"))
+	  block_name = NULL;
+      }
+      break;
+
     default:
       break;
     }
@@ -8609,6 +8632,12 @@ gfc_match_end (gfc_statement *st)
       gfc_free_enum_history ();
       break;
 
+    case COMP_OMP_BEGIN_METADIRECTIVE:
+      *st = ST_OMP_END_METADIRECTIVE;
+      target = " metadirective";
+      eos_ok = 0;
+      break;
+
     default:
       gfc_error ("Unexpected END statement at %C");
       goto cleanup;
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 0ae1350..adc07da 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -2377,6 +2377,7 @@ show_omp_node (int level, gfc_code *c)
     case EXEC_OMP_MASTER: name = "MASTER"; break;
     case EXEC_OMP_MASTER_TASKLOOP: name = "MASTER TASKLOOP"; break;
     case EXEC_OMP_MASTER_TASKLOOP_SIMD: name = "MASTER TASKLOOP SIMD"; break;
+    case EXEC_OMP_METADIRECTIVE: name = "METADIRECTIVE"; break;
     case EXEC_OMP_ORDERED: name = "ORDERED"; break;
     case EXEC_OMP_DEPOBJ: name = "DEPOBJ"; break;
     case EXEC_OMP_PARALLEL: name = "PARALLEL"; break;
@@ -2581,6 +2582,24 @@ show_omp_node (int level, gfc_code *c)
 	  d = d->block;
 	}
     }
+  else if (c->op == EXEC_OMP_METADIRECTIVE)
+    {
+      gfc_omp_variant *variant = c->ext.omp_variants;
+
+      while (variant)
+	{
+	  code_indent (level + 1, 0);
+	  if (variant->selectors)
+	    fputs ("WHEN ()\n", dumpfile);
+	  else
+	    fputs ("DEFAULT ()\n", dumpfile);
+	  /* TODO: Print selector.  */
+	  show_code (level + 2, variant->code);
+	  if (variant->next)
+	    fputs ("\n", dumpfile);
+	  variant = variant->next;
+	}
+    }
   else
     show_code (level + 1, c->block->next);
   if (c->op == EXEC_OMP_ATOMIC)
@@ -3821,6 +3840,7 @@ show_code_node (int level, gfc_code *c)
     case EXEC_OMP_MASTER:
     case EXEC_OMP_MASTER_TASKLOOP:
     case EXEC_OMP_MASTER_TASKLOOP_SIMD:
+    case EXEC_OMP_METADIRECTIVE:
     case EXEC_OMP_ORDERED:
     case EXEC_OMP_PARALLEL:
     case EXEC_OMP_PARALLEL_DO:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 83e4f3f..5fe1276 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -318,6 +318,7 @@ enum gfc_statement
   ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD, ST_OMP_MASKED_TASKLOOP,
   ST_OMP_END_MASKED_TASKLOOP, ST_OMP_MASKED_TASKLOOP_SIMD,
   ST_OMP_END_MASKED_TASKLOOP_SIMD, ST_OMP_SCOPE, ST_OMP_END_SCOPE,
+  ST_OMP_METADIRECTIVE, ST_OMP_BEGIN_METADIRECTIVE, ST_OMP_END_METADIRECTIVE,
   ST_OMP_ERROR, ST_OMP_ASSUME, ST_OMP_END_ASSUME, ST_OMP_ASSUMES,
   ST_OMP_ALLOCATE, ST_OMP_ALLOCATE_EXEC,
   ST_OMP_ALLOCATORS, ST_OMP_END_ALLOCATORS,
@@ -1634,7 +1635,7 @@ typedef struct gfc_omp_clauses
   unsigned order_unconstrained:1, order_reproducible:1, capture:1;
   unsigned grainsize_strict:1, num_tasks_strict:1, compare:1, weak:1;
   unsigned non_rectangular:1, order_concurrent:1;
-  unsigned contains_teams_construct:1, target_first_st_is_teams:1;
+  unsigned contains_teams_construct:1, target_first_st_is_teams_or_meta:1;
   unsigned contained_in_target_construct:1, indirect:1;
   unsigned full:1, erroneous:1;
   ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3;
@@ -1757,6 +1758,17 @@ typedef struct gfc_omp_declare_variant
 gfc_omp_declare_variant;
 #define gfc_get_omp_declare_variant() XCNEW (gfc_omp_declare_variant)
 
+typedef struct gfc_omp_variant
+{
+  struct gfc_omp_variant *next;
+  locus where; /* Where the metadirective clause occurred.  */
+
+  gfc_omp_set_selector *selectors;
+  enum gfc_statement stmt;
+  struct gfc_code *code;
+
+} gfc_omp_variant;
+#define gfc_get_omp_variant() XCNEW (gfc_omp_variant)
 
 typedef struct gfc_omp_udr
 {
@@ -1805,6 +1817,7 @@ typedef struct gfc_st_label
   locus where;
 
   gfc_namespace *ns;
+  int omp_region;
 }
 gfc_st_label;
 
@@ -3108,7 +3121,7 @@ enum gfc_exec_op
   EXEC_OMP_TARGET_TEAMS_LOOP, EXEC_OMP_MASKED, EXEC_OMP_PARALLEL_MASKED,
   EXEC_OMP_PARALLEL_MASKED_TASKLOOP, EXEC_OMP_PARALLEL_MASKED_TASKLOOP_SIMD,
   EXEC_OMP_MASKED_TASKLOOP, EXEC_OMP_MASKED_TASKLOOP_SIMD, EXEC_OMP_SCOPE,
-  EXEC_OMP_UNROLL, EXEC_OMP_TILE, EXEC_OMP_INTEROP,
+  EXEC_OMP_UNROLL, EXEC_OMP_TILE, EXEC_OMP_INTEROP, EXEC_OMP_METADIRECTIVE,
   EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS, EXEC_OMP_DISPATCH
 };
 
@@ -3154,6 +3167,7 @@ typedef struct gfc_code
     gfc_omp_clauses *omp_clauses;
     const char *omp_name;
     gfc_omp_namelist *omp_namelist;
+    gfc_omp_variant *omp_variants;
     bool omp_bool;
     int stop_code;
 
@@ -3802,6 +3816,7 @@ void gfc_free_omp_declare_variant_list (gfc_omp_declare_variant *list);
 void gfc_free_omp_declare_simd (gfc_omp_declare_simd *);
 void gfc_free_omp_declare_simd_list (gfc_omp_declare_simd *);
 void gfc_free_omp_udr (gfc_omp_udr *);
+void gfc_free_omp_variants (gfc_omp_variant *);
 gfc_omp_udr *gfc_omp_udr_find (gfc_symtree *, gfc_typespec *);
 void gfc_resolve_omp_allocate (gfc_namespace *, gfc_omp_namelist *);
 void gfc_resolve_omp_assumptions (gfc_omp_assumptions *);
@@ -4089,6 +4104,8 @@ void debug (gfc_expr *);
 bool gfc_parse_file (void);
 void gfc_global_used (gfc_gsymbol *, locus *);
 gfc_namespace* gfc_build_block_ns (gfc_namespace *);
+gfc_statement match_omp_directive (void);
+bool is_omp_declarative_stmt (gfc_statement);
 
 /* dependency.cc */
 int gfc_dep_compare_functions (gfc_expr *, gfc_expr *, bool);
diff --git a/gcc/fortran/io.cc b/gcc/fortran/io.cc
index 48f4359..b5c9d33 100644
--- a/gcc/fortran/io.cc
+++ b/gcc/fortran/io.cc
@@ -29,7 +29,7 @@ along with GCC; see the file COPYING3.  If not see
 
 gfc_st_label
 format_asterisk = {0, NULL, NULL, -1, ST_LABEL_FORMAT, ST_LABEL_FORMAT, NULL,
-		   0, {NULL, NULL}, NULL};
+		   0, {NULL, NULL}, NULL, 0};
 
 typedef struct
 {
diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h
index 4041613..410361c 100644
--- a/gcc/fortran/match.h
+++ b/gcc/fortran/match.h
@@ -155,6 +155,7 @@ match gfc_match_omp_assume (void);
 match gfc_match_omp_assumes (void);
 match gfc_match_omp_atomic (void);
 match gfc_match_omp_barrier (void);
+match gfc_match_omp_begin_metadirective (void);
 match gfc_match_omp_cancel (void);
 match gfc_match_omp_cancellation_point (void);
 match gfc_match_omp_critical (void);
@@ -180,6 +181,7 @@ match gfc_match_omp_masked_taskloop_simd (void);
 match gfc_match_omp_master (void);
 match gfc_match_omp_master_taskloop (void);
 match gfc_match_omp_master_taskloop_simd (void);
+match gfc_match_omp_metadirective (void);
 match gfc_match_omp_nothing (void);
 match gfc_match_omp_ordered (void);
 match gfc_match_omp_ordered_depend (void);
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 5eef5eb..b1684f8 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -83,7 +83,7 @@ static const struct gfc_omp_directive gfc_omp_directives[] = {
   {"interop", GFC_OMP_DIR_EXECUTABLE, ST_OMP_INTEROP},
   {"loop", GFC_OMP_DIR_EXECUTABLE, ST_OMP_LOOP},
   {"masked", GFC_OMP_DIR_EXECUTABLE, ST_OMP_MASKED},
-  /* {"metadirective", GFC_OMP_DIR_META, ST_OMP_METADIRECTIVE}, */
+  {"metadirective", GFC_OMP_DIR_META, ST_OMP_METADIRECTIVE},
   /* Note: gfc_match_omp_nothing returns ST_NONE.  */
   {"nothing", GFC_OMP_DIR_UTILITY, ST_OMP_NOTHING},
   /* Special case; for now map to the first one.
@@ -116,7 +116,8 @@ static const struct gfc_omp_directive gfc_omp_directives[] = {
 
 
 /* Match an end of OpenMP directive.  End of OpenMP directive is optional
-   whitespace, followed by '\n' or comment '!'.  */
+   whitespace, followed by '\n' or comment '!'.  In the special case where a
+   context selector is being matched, match against ')' instead.  */
 
 static match
 gfc_match_omp_eos (void)
@@ -127,17 +128,25 @@ gfc_match_omp_eos (void)
   old_loc = gfc_current_locus;
   gfc_gobble_whitespace ();
 
-  c = gfc_next_ascii_char ();
-  switch (c)
+  if (gfc_matching_omp_context_selector)
     {
-    case '!':
-      do
-	c = gfc_next_ascii_char ();
-      while (c != '\n');
-      /* Fall through */
+      if (gfc_peek_ascii_char () == ')')
+	return MATCH_YES;
+    }
+  else
+    {
+      c = gfc_next_ascii_char ();
+      switch (c)
+	{
+	case '!':
+	  do
+	    c = gfc_next_ascii_char ();
+	  while (c != '\n');
+	  /* Fall through */
 
-    case '\n':
-      return MATCH_YES;
+	case '\n':
+	  return MATCH_YES;
+	}
     }
 
   gfc_current_locus = old_loc;
@@ -349,6 +358,19 @@ gfc_free_omp_udr (gfc_omp_udr *omp_udr)
     }
 }
 
+/* Free variants of an !$omp metadirective construct.  */
+
+void
+gfc_free_omp_variants (gfc_omp_variant *variant)
+{
+  while (variant)
+    {
+      gfc_omp_variant *next_variant = variant->next;
+      gfc_free_omp_set_selector_list (variant->selectors);
+      free (variant);
+      variant = next_variant;
+    }
+}
 
 static gfc_omp_udr *
 gfc_find_omp_udr (gfc_namespace *ns, const char *name, gfc_typespec *ts)
@@ -2321,8 +2343,7 @@ gfc_match_dupl_atomic (bool not_dupl, const char *name)
 static match
 gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 		       bool first = true, bool needs_space = true,
-		       bool openacc = false, bool context_selector = false,
-		       bool openmp_target = false)
+		       bool openacc = false, bool openmp_target = false)
 {
   bool error = false;
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
@@ -4384,9 +4405,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
     }
 
 end:
-  if (error
-      || (context_selector && gfc_peek_ascii_char () != ')')
-      || (!context_selector && gfc_match_omp_eos () != MATCH_YES))
+  if (error || gfc_match_omp_eos () != MATCH_YES)
     {
       if (!gfc_error_flag_test ())
 	gfc_error ("Failed to match clause at %C");
@@ -5100,7 +5119,7 @@ static match
 match_omp (gfc_exec_op op, const omp_mask mask)
 {
   gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, mask, true, true, false, false,
+  if (gfc_match_omp_clauses (&c, mask, true, true, false,
 			     op == EXEC_OMP_TARGET) != MATCH_YES)
     return MATCH_ERROR;
   new_st.op = op;
@@ -6295,7 +6314,8 @@ gfc_match_omp_interop (void)
      score(score-expression)  */
 
 match
-gfc_match_omp_context_selector (gfc_omp_set_selector *oss)
+gfc_match_omp_context_selector (gfc_omp_set_selector *oss,
+				bool metadirective_p)
 {
   do
     {
@@ -6455,14 +6475,31 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss)
 		  || (property_kind == OMP_TRAIT_PROPERTY_DEV_NUM_EXPR
 		      && otp->expr->ts.type != BT_INTEGER)
 		  || otp->expr->rank != 0
-		  || otp->expr->expr_type != EXPR_CONSTANT)
+		  || (!metadirective_p
+		      && otp->expr->expr_type != EXPR_CONSTANT))
 		{
-		  if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR)
-		    gfc_error ("property must be a constant logical expression "
-			       "at %C");
+		  if (metadirective_p)
+		    {
+		      if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR)
+			gfc_error ("property must be a "
+				   "logical expression at %L",
+				   &otp->expr->where);
+		      else
+			gfc_error ("property must be an "
+				   "integer expression at %L",
+				   &otp->expr->where);
+		    }
 		  else
-		    gfc_error ("property must be a constant integer expression "
-			       "at %C");
+		    {
+		      if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR)
+			gfc_error ("property must be a constant "
+				   "logical expression at %L",
+				   &otp->expr->where);
+		      else
+			gfc_error ("property must be a constant "
+				   "integer expression at %L",
+				   &otp->expr->where);
+		    }
 		  return MATCH_ERROR;
 		}
 	      /* Device number must be conforming, which includes
@@ -6482,14 +6519,17 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss)
 	      {
 		if (os->code == OMP_TRAIT_CONSTRUCT_SIMD)
 		  {
+		    gfc_matching_omp_context_selector = true;
 		    if (gfc_match_omp_clauses (&otp->clauses,
 					       OMP_DECLARE_SIMD_CLAUSES,
-					       true, false, false, true)
+					       true, false, false)
 			!= MATCH_YES)
 		      {
+			gfc_matching_omp_context_selector = false;
 			gfc_error ("expected simd clause at %C");
 			return MATCH_ERROR;
 		      }
+		    gfc_matching_omp_context_selector = false;
 		  }
 		else if (os->code == OMP_TRAIT_IMPLEMENTATION_REQUIRES)
 		  {
@@ -6546,7 +6586,8 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss)
      user  */
 
 match
-gfc_match_omp_context_selector_specification (gfc_omp_declare_variant *odv)
+gfc_match_omp_context_selector_specification (gfc_omp_set_selector **oss_head,
+					      bool metadirective_p)
 {
   do
     {
@@ -6579,11 +6620,11 @@ gfc_match_omp_context_selector_specification (gfc_omp_declare_variant *odv)
 	}
 
       gfc_omp_set_selector *oss = gfc_get_omp_set_selector ();
-      oss->next = odv->set_selectors;
+      oss->next = *oss_head;
       oss->code = set;
-      odv->set_selectors = oss;
+      *oss_head = oss;
 
-      if (gfc_match_omp_context_selector (oss) != MATCH_YES)
+      if (gfc_match_omp_context_selector (oss, metadirective_p) != MATCH_YES)
 	return MATCH_ERROR;
 
       m = gfc_match (" }");
@@ -6714,7 +6755,8 @@ gfc_match_omp_declare_variant (void)
 	      return MATCH_ERROR;
 	    }
 	  has_match = true;
-	  if (gfc_match_omp_context_selector_specification (odv)
+	  if (gfc_match_omp_context_selector_specification (&odv->set_selectors,
+							    false)
 	      != MATCH_YES)
 	    return MATCH_ERROR;
 	  if (gfc_match (" )") != MATCH_YES)
@@ -6831,6 +6873,167 @@ gfc_match_omp_declare_variant (void)
 }
 
 
+static match
+match_omp_metadirective (bool begin_p)
+{
+  locus old_loc = gfc_current_locus;
+  gfc_omp_variant *variants_head;
+  gfc_omp_variant **next_variant = &variants_head;
+  bool default_seen = false;
+
+  /* Parse the context selectors.  */
+  for (;;)
+    {
+      bool default_p = false;
+      gfc_omp_set_selector *selectors = NULL;
+
+      gfc_gobble_whitespace ();
+      if (gfc_match_eos () == MATCH_YES)
+	break;
+      gfc_match_char (',');
+      gfc_gobble_whitespace ();
+
+      locus variant_locus = gfc_current_locus;
+
+      if (gfc_match (" default ( ") == MATCH_YES)
+	default_p = true;
+      else if (gfc_match (" otherwise ( ") == MATCH_YES)
+	default_p = true;
+      else if (gfc_match (" when ( ") != MATCH_YES)
+	{
+	  gfc_error ("expected %<when%>, %<otherwise%>, or %<default%> at %C");
+	  gfc_current_locus = old_loc;
+	  return MATCH_ERROR;
+	}
+
+      if (default_p && default_seen)
+	{
+	  gfc_error ("too many %<otherwise%> or %<default%> clauses "
+		     "in %<metadirective%> at %C");
+	  gfc_current_locus = old_loc;
+	  return MATCH_ERROR;
+	}
+      else if (default_seen)
+	{
+	  gfc_error ("%<otherwise%> or %<default%> clause "
+		     "must appear last in %<metadirective%> at %C");
+	  gfc_current_locus = old_loc;
+	  return MATCH_ERROR;
+	}
+
+      if (!default_p)
+	{
+	  if (gfc_match_omp_context_selector_specification (&selectors, true)
+	      != MATCH_YES)
+	    return MATCH_ERROR;
+
+	  if (gfc_match (" : ") != MATCH_YES)
+	    {
+	      gfc_error ("expected %<:%> at %C");
+	      gfc_current_locus = old_loc;
+	      return MATCH_ERROR;
+	    }
+
+	  gfc_commit_symbols ();
+	}
+
+      gfc_matching_omp_context_selector = true;
+      gfc_statement directive = match_omp_directive ();
+      gfc_matching_omp_context_selector = false;
+
+      if (is_omp_declarative_stmt (directive))
+	sorry ("declarative directive variants are not supported");
+
+      if (gfc_error_flag_test ())
+	{
+	  gfc_current_locus = old_loc;
+	  return MATCH_ERROR;
+	}
+
+      if (gfc_match (" )") != MATCH_YES)
+	{
+	  gfc_error ("Expected %<)%> at %C");
+	  gfc_current_locus = old_loc;
+	  return MATCH_ERROR;
+	}
+
+      gfc_commit_symbols ();
+
+      if (begin_p
+	  && directive != ST_NONE
+	  && gfc_omp_end_stmt (directive) == ST_NONE)
+	{
+	  gfc_error ("variant directive used in OMP BEGIN METADIRECTIVE "
+		     "at %C must have a corresponding end directive");
+	  gfc_current_locus = old_loc;
+	  return MATCH_ERROR;
+	}
+
+      if (default_p)
+	default_seen = true;
+
+      gfc_omp_variant *omv = gfc_get_omp_variant ();
+      omv->selectors = selectors;
+      omv->stmt = directive;
+      omv->where = variant_locus;
+
+      if (directive == ST_NONE)
+	{
+	  /* The directive was a 'nothing' directive.  */
+	  omv->code = gfc_get_code (EXEC_CONTINUE);
+	  omv->code->ext.omp_clauses = NULL;
+	}
+      else
+	{
+	  omv->code = gfc_get_code (new_st.op);
+	  omv->code->ext.omp_clauses = new_st.ext.omp_clauses;
+	  /* Prevent the OpenMP clauses from being freed via NEW_ST.  */
+	  new_st.ext.omp_clauses = NULL;
+	}
+
+      *next_variant = omv;
+      next_variant = &omv->next;
+    }
+
+  if (gfc_match_omp_eos () != MATCH_YES)
+    {
+      gfc_error ("Unexpected junk after OMP METADIRECTIVE at %C");
+      gfc_current_locus = old_loc;
+      return MATCH_ERROR;
+    }
+
+  /* Add a 'default (nothing)' clause if no default is explicitly given.  */
+  if (!default_seen)
+    {
+      gfc_omp_variant *omv = gfc_get_omp_variant ();
+      omv->stmt = ST_NONE;
+      omv->code = gfc_get_code (EXEC_CONTINUE);
+      omv->code->ext.omp_clauses = NULL;
+      omv->where = old_loc;
+      omv->selectors = NULL;
+
+      *next_variant = omv;
+      next_variant = &omv->next;
+    }
+
+  new_st.op = EXEC_OMP_METADIRECTIVE;
+  new_st.ext.omp_variants = variants_head;
+
+  return MATCH_YES;
+}
+
+match
+gfc_match_omp_begin_metadirective (void)
+{
+  return match_omp_metadirective (true);
+}
+
+match
+gfc_match_omp_metadirective (void)
+{
+  return match_omp_metadirective (false);
+}
+
 match
 gfc_match_omp_threadprivate (void)
 {
@@ -11987,6 +12190,19 @@ resolve_omp_do (gfc_code *code)
 				  non_generated_count);
 }
 
+static void
+resolve_omp_metadirective (gfc_code *code, gfc_namespace *ns)
+{
+  gfc_omp_variant *variant = code->ext.omp_variants;
+
+  while (variant)
+    {
+      gfc_code *variant_code = variant->code;
+      gfc_resolve_code (variant_code, ns);
+      variant = variant->next;
+    }
+}
+
 
 static gfc_statement
 omp_code_to_statement (gfc_code *code)
@@ -12538,13 +12754,32 @@ resolve_omp_target (gfc_code *code)
   gfc_code *c = code->block->next;
   if (c->op == EXEC_BLOCK)
     c = c->ext.block.ns->code;
-  if (code->ext.omp_clauses->target_first_st_is_teams
-      && ((GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL)
-	  || (c->op == EXEC_BLOCK
-	      && c->next
-	      && GFC_IS_TEAMS_CONSTRUCT (c->next->op)
-	      && c->next->next == NULL)))
-    return;
+  if (code->ext.omp_clauses->target_first_st_is_teams_or_meta)
+    {
+      if (c->op == EXEC_OMP_METADIRECTIVE)
+	{
+	  struct gfc_omp_variant *mc
+	    = c->ext.omp_variants;
+	  /* All mc->(next...->)code should be identical with regards
+	     to the diagnostic below.  */
+	  do
+	    {
+	      if (mc->stmt != ST_NONE
+		  && GFC_IS_TEAMS_CONSTRUCT (mc->code->op))
+		{
+		  if (c->next == NULL && mc->code->next == NULL)
+		    return;
+		  c = mc->code;
+		  break;
+		}
+	      mc = mc->next;
+	    }
+	  while (mc);
+	}
+      else if (GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL)
+	return;
+    }
+
   while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op))
     c = c->next;
   if (c)
@@ -12714,6 +12949,9 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns)
 	resolve_omp_clauses (code, code->ext.omp_clauses, ns);
       resolve_omp_dispatch (code);
       break;
+    case EXEC_OMP_METADIRECTIVE:
+      resolve_omp_metadirective (code, ns);
+      break;
     default:
       break;
     }
diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc
index a75284e..00cd23d 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -48,6 +48,16 @@ gfc_state_data *gfc_state_stack;
 static bool last_was_use_stmt = false;
 bool in_exec_part;
 
+/* True when matching an OpenMP context selector.  */
+bool gfc_matching_omp_context_selector;
+
+/* True when parsing the body of an OpenMP metadirective.  */
+bool gfc_in_omp_metadirective_body;
+
+/* Each metadirective body in the translation unit is given a unique
+   number, used to ensure that labels in the body have unique names.  */
+int gfc_omp_metadirective_region_count;
+
 /* TODO: Re-order functions to kill these forward decls.  */
 static void check_statement_label (gfc_statement);
 static void undo_new_statement (void);
@@ -993,6 +1003,12 @@ decode_omp_directive (void)
       matcho ("assumes", gfc_match_omp_assumes, ST_OMP_ASSUMES);
       matchs ("assume", gfc_match_omp_assume, ST_OMP_ASSUME);
       break;
+
+    case 'b':
+      matcho ("begin metadirective", gfc_match_omp_begin_metadirective,
+	      ST_OMP_BEGIN_METADIRECTIVE);
+      break;
+
     case 'd':
       matchds ("declare reduction", gfc_match_omp_declare_reduction,
 	       ST_OMP_DECLARE_REDUCTION);
@@ -1005,11 +1021,19 @@ decode_omp_directive (void)
       break;
     case 'e':
       matchs ("end assume", gfc_match_omp_eos_error, ST_OMP_END_ASSUME);
+      matcho ("end metadirective", gfc_match_omp_eos_error,
+	      ST_OMP_END_METADIRECTIVE);
       matchs ("end simd", gfc_match_omp_eos_error, ST_OMP_END_SIMD);
       matchs ("end tile", gfc_match_omp_eos_error, ST_OMP_END_TILE);
       matchs ("end unroll", gfc_match_omp_eos_error, ST_OMP_END_UNROLL);
       matcho ("error", gfc_match_omp_error, ST_OMP_ERROR);
       break;
+
+    case 'm':
+      matcho ("metadirective", gfc_match_omp_metadirective,
+	      ST_OMP_METADIRECTIVE);
+      break;
+
     case 'n':
       matcho ("nothing", gfc_match_omp_nothing, ST_NONE);
       break;
@@ -1309,6 +1333,10 @@ decode_omp_directive (void)
 	gfc_error_now ("Unclassifiable OpenMP directive at %C");
     }
 
+  /* If parsing a metadirective, let the caller deal with the cleanup.  */
+  if (gfc_matching_omp_context_selector)
+    return ST_NONE;
+
   reject_statement ();
 
   gfc_error_recovery ();
@@ -1430,6 +1458,12 @@ decode_omp_directive (void)
   return ST_GET_FCN_CHARACTERISTICS;
 }
 
+gfc_statement
+match_omp_directive (void)
+{
+  return decode_omp_directive ();
+}
+
 static gfc_statement
 decode_gcc_attribute (void)
 {
@@ -1955,6 +1989,44 @@ next_statement (void)
   case ST_OMP_DECLARE_VARIANT: case ST_OMP_ALLOCATE: case ST_OMP_ASSUMES: \
   case ST_OMP_REQUIRES: case ST_OACC_ROUTINE: case ST_OACC_DECLARE
 
+/* OpenMP statements that are followed by a structured block.  */
+
+#define case_omp_structured_block case ST_OMP_ASSUME: case ST_OMP_PARALLEL: \
+  case ST_OMP_PARALLEL_MASKED: case ST_OMP_PARALLEL_MASTER: \
+  case ST_OMP_PARALLEL_SECTIONS: case ST_OMP_ORDERED: \
+  case ST_OMP_CRITICAL: case ST_OMP_MASKED: case ST_OMP_MASTER: \
+  case ST_OMP_SCOPE: case ST_OMP_SECTIONS: case ST_OMP_SINGLE: \
+  case ST_OMP_TARGET: case ST_OMP_TARGET_DATA: case ST_OMP_TARGET_PARALLEL: \
+  case ST_OMP_TARGET_TEAMS: case ST_OMP_TEAMS: case ST_OMP_TASK: \
+  case ST_OMP_TASKGROUP: \
+  case ST_OMP_WORKSHARE: case ST_OMP_PARALLEL_WORKSHARE
+
+/* OpenMP statements that are followed by a do loop.  */
+
+#define case_omp_do case ST_OMP_DISTRIBUTE: \
+  case ST_OMP_DISTRIBUTE_PARALLEL_DO: \
+  case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: case ST_OMP_DISTRIBUTE_SIMD: \
+  case ST_OMP_DO: case ST_OMP_DO_SIMD: case ST_OMP_LOOP: \
+  case ST_OMP_PARALLEL_DO: case ST_OMP_PARALLEL_DO_SIMD: \
+  case ST_OMP_PARALLEL_LOOP: case ST_OMP_PARALLEL_MASKED_TASKLOOP: \
+  case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD: \
+  case ST_OMP_PARALLEL_MASTER_TASKLOOP: \
+  case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD: \
+  case ST_OMP_MASKED_TASKLOOP: case ST_OMP_MASKED_TASKLOOP_SIMD: \
+  case ST_OMP_MASTER_TASKLOOP: case ST_OMP_MASTER_TASKLOOP_SIMD: \
+  case ST_OMP_SIMD: \
+  case ST_OMP_TARGET_PARALLEL_DO: case ST_OMP_TARGET_PARALLEL_DO_SIMD: \
+  case ST_OMP_TARGET_PARALLEL_LOOP: case ST_OMP_TARGET_SIMD: \
+  case ST_OMP_TARGET_TEAMS_DISTRIBUTE: \
+  case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: \
+  case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: \
+  case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: case ST_OMP_TARGET_TEAMS_LOOP: \
+  case ST_OMP_TASKLOOP: case ST_OMP_TASKLOOP_SIMD: \
+  case ST_OMP_TEAMS_DISTRIBUTE: case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: \
+  case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: \
+  case ST_OMP_TEAMS_DISTRIBUTE_SIMD: case ST_OMP_TEAMS_LOOP: \
+  case ST_OMP_TILE: case ST_OMP_UNROLL
+
 /* Block end statements.  Errors associated with interchanging these
    are detected in gfc_match_end().  */
 
@@ -2592,6 +2664,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel)
     case ST_OMP_BARRIER:
       p = "!$OMP BARRIER";
       break;
+    case ST_OMP_BEGIN_METADIRECTIVE:
+      p = "!$OMP BEGIN METADIRECTIVE";
+      break;
     case ST_OMP_CANCEL:
       p = "!$OMP CANCEL";
       break;
@@ -2697,6 +2772,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel)
     case ST_OMP_END_MASTER_TASKLOOP_SIMD:
       p = "!$OMP END MASTER TASKLOOP SIMD";
       break;
+    case ST_OMP_END_METADIRECTIVE:
+      p = "!$OMP END METADIRECTIVE";
+      break;
     case ST_OMP_END_ORDERED:
       p = "!$OMP END ORDERED";
       break;
@@ -2850,6 +2928,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel)
     case ST_OMP_MASTER_TASKLOOP_SIMD:
       p = "!$OMP MASTER TASKLOOP SIMD";
       break;
+    case ST_OMP_METADIRECTIVE:
+      p = "!$OMP METADIRECTIVE";
+      break;
     case ST_OMP_ORDERED:
     case ST_OMP_ORDERED_DEPEND:
       p = "!$OMP ORDERED";
@@ -3116,6 +3197,8 @@ accept_statement (gfc_statement st)
       break;
 
     case ST_ENTRY:
+    case ST_OMP_METADIRECTIVE:
+    case ST_OMP_BEGIN_METADIRECTIVE:
     case_executable:
     case_exec_markers:
       add_statement ();
@@ -5511,6 +5594,150 @@ loop:
   accept_statement (st);
 }
 
+/* Get the corresponding ending statement type for the OpenMP directive
+   OMP_ST.  If it does not have one, return ST_NONE.  */
+
+gfc_statement
+gfc_omp_end_stmt (gfc_statement omp_st,
+		  bool omp_do_p, bool omp_structured_p)
+{
+  if (omp_do_p)
+    {
+      switch (omp_st)
+	{
+	case ST_OMP_DISTRIBUTE: return ST_OMP_END_DISTRIBUTE;
+	case ST_OMP_DISTRIBUTE_PARALLEL_DO:
+	  return ST_OMP_END_DISTRIBUTE_PARALLEL_DO;
+	case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+	  return ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD;
+	case ST_OMP_DISTRIBUTE_SIMD:
+	  return ST_OMP_END_DISTRIBUTE_SIMD;
+	case ST_OMP_DO: return ST_OMP_END_DO;
+	case ST_OMP_DO_SIMD: return ST_OMP_END_DO_SIMD;
+	case ST_OMP_LOOP: return ST_OMP_END_LOOP;
+	case ST_OMP_PARALLEL_DO: return ST_OMP_END_PARALLEL_DO;
+	case ST_OMP_PARALLEL_DO_SIMD:
+	  return ST_OMP_END_PARALLEL_DO_SIMD;
+	case ST_OMP_PARALLEL_LOOP:
+	  return ST_OMP_END_PARALLEL_LOOP;
+	case ST_OMP_SIMD: return ST_OMP_END_SIMD;
+	case ST_OMP_TARGET_PARALLEL_DO:
+	  return ST_OMP_END_TARGET_PARALLEL_DO;
+	case ST_OMP_TARGET_PARALLEL_DO_SIMD:
+	  return ST_OMP_END_TARGET_PARALLEL_DO_SIMD;
+	case ST_OMP_TARGET_PARALLEL_LOOP:
+	  return ST_OMP_END_TARGET_PARALLEL_LOOP;
+	case ST_OMP_TARGET_SIMD: return ST_OMP_END_TARGET_SIMD;
+	case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
+	  return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE;
+	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	  return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO;
+	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	  return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
+	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+	  return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD;
+	case ST_OMP_TARGET_TEAMS_LOOP:
+	  return ST_OMP_END_TARGET_TEAMS_LOOP;
+	case ST_OMP_TASKLOOP: return ST_OMP_END_TASKLOOP;
+	case ST_OMP_TASKLOOP_SIMD: return ST_OMP_END_TASKLOOP_SIMD;
+	case ST_OMP_MASKED_TASKLOOP: return ST_OMP_END_MASKED_TASKLOOP;
+	case ST_OMP_MASKED_TASKLOOP_SIMD:
+	  return ST_OMP_END_MASKED_TASKLOOP_SIMD;
+	case ST_OMP_MASTER_TASKLOOP: return ST_OMP_END_MASTER_TASKLOOP;
+	case ST_OMP_MASTER_TASKLOOP_SIMD:
+	  return ST_OMP_END_MASTER_TASKLOOP_SIMD;
+	case ST_OMP_PARALLEL_MASKED_TASKLOOP:
+	  return ST_OMP_END_PARALLEL_MASKED_TASKLOOP;
+	case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD:
+	  return ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD;
+	case ST_OMP_PARALLEL_MASTER_TASKLOOP:
+	  return ST_OMP_END_PARALLEL_MASTER_TASKLOOP;
+	case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD:
+	  return ST_OMP_END_PARALLEL_MASTER_TASKLOOP_SIMD;
+	case ST_OMP_TEAMS_DISTRIBUTE:
+	  return ST_OMP_END_TEAMS_DISTRIBUTE;
+	case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	  return ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO;
+	case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	  return ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
+	case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
+	  return ST_OMP_END_TEAMS_DISTRIBUTE_SIMD;
+	case ST_OMP_TEAMS_LOOP:
+	  return ST_OMP_END_TEAMS_LOOP;
+	case ST_OMP_TILE:
+	  return ST_OMP_END_TILE;
+	case ST_OMP_UNROLL:
+	  return ST_OMP_END_UNROLL;
+	default:
+	  break;
+	}
+    }
+
+  if (omp_structured_p)
+    {
+      switch (omp_st)
+	{
+	case ST_OMP_ALLOCATORS:
+	  return ST_OMP_END_ALLOCATORS;
+	case ST_OMP_ASSUME:
+	  return ST_OMP_END_ASSUME;
+	case ST_OMP_ATOMIC:
+	  return ST_OMP_END_ATOMIC;
+	case ST_OMP_DISPATCH:
+	  return ST_OMP_END_DISPATCH;
+	case ST_OMP_PARALLEL:
+	  return ST_OMP_END_PARALLEL;
+	case ST_OMP_PARALLEL_MASKED:
+	  return ST_OMP_END_PARALLEL_MASKED;
+	case ST_OMP_PARALLEL_MASTER:
+	  return ST_OMP_END_PARALLEL_MASTER;
+	case ST_OMP_PARALLEL_SECTIONS:
+	  return ST_OMP_END_PARALLEL_SECTIONS;
+	case ST_OMP_SCOPE:
+	  return ST_OMP_END_SCOPE;
+	case ST_OMP_SECTIONS:
+	  return ST_OMP_END_SECTIONS;
+	case ST_OMP_ORDERED:
+	  return ST_OMP_END_ORDERED;
+	case ST_OMP_CRITICAL:
+	  return ST_OMP_END_CRITICAL;
+	case ST_OMP_MASKED:
+	  return ST_OMP_END_MASKED;
+	case ST_OMP_MASTER:
+	  return ST_OMP_END_MASTER;
+	case ST_OMP_SINGLE:
+	  return ST_OMP_END_SINGLE;
+	case ST_OMP_TARGET:
+	  return ST_OMP_END_TARGET;
+	case ST_OMP_TARGET_DATA:
+	  return ST_OMP_END_TARGET_DATA;
+	case ST_OMP_TARGET_PARALLEL:
+	  return ST_OMP_END_TARGET_PARALLEL;
+	case ST_OMP_TARGET_TEAMS:
+	  return ST_OMP_END_TARGET_TEAMS;
+	case ST_OMP_TASK:
+	  return ST_OMP_END_TASK;
+	case ST_OMP_TASKGROUP:
+	  return ST_OMP_END_TASKGROUP;
+	case ST_OMP_TEAMS:
+	  return ST_OMP_END_TEAMS;
+	case ST_OMP_TEAMS_DISTRIBUTE:
+	  return ST_OMP_END_TEAMS_DISTRIBUTE;
+	case ST_OMP_DISTRIBUTE:
+	  return ST_OMP_END_DISTRIBUTE;
+	case ST_OMP_WORKSHARE:
+	  return ST_OMP_END_WORKSHARE;
+	case ST_OMP_PARALLEL_WORKSHARE:
+	  return ST_OMP_END_PARALLEL_WORKSHARE;
+	case ST_OMP_BEGIN_METADIRECTIVE:
+	  return ST_OMP_END_METADIRECTIVE;
+	default:
+	  break;
+	}
+    }
+
+  return ST_NONE;
+}
 
 /* Parse the statements of OpenMP do/parallel do.  */
 
@@ -5571,94 +5798,16 @@ parse_omp_do (gfc_statement omp_st, int nested)
 
   st = next_statement ();
 do_end:
-  gfc_statement omp_end_st = ST_OMP_END_DO;
-  switch (omp_st)
-    {
-    case ST_OMP_DISTRIBUTE: omp_end_st = ST_OMP_END_DISTRIBUTE; break;
-    case ST_OMP_DISTRIBUTE_PARALLEL_DO:
-      omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO;
-      break;
-    case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
-      omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD;
-      break;
-    case ST_OMP_DISTRIBUTE_SIMD:
-      omp_end_st = ST_OMP_END_DISTRIBUTE_SIMD;
-      break;
-    case ST_OMP_DO: omp_end_st = ST_OMP_END_DO; break;
-    case ST_OMP_DO_SIMD: omp_end_st = ST_OMP_END_DO_SIMD; break;
-    case ST_OMP_LOOP: omp_end_st = ST_OMP_END_LOOP; break;
-    case ST_OMP_PARALLEL_DO: omp_end_st = ST_OMP_END_PARALLEL_DO; break;
-    case ST_OMP_PARALLEL_DO_SIMD:
-      omp_end_st = ST_OMP_END_PARALLEL_DO_SIMD;
-      break;
-    case ST_OMP_PARALLEL_LOOP:
-      omp_end_st = ST_OMP_END_PARALLEL_LOOP;
-      break;
-    case ST_OMP_SIMD: omp_end_st = ST_OMP_END_SIMD; break;
-    case ST_OMP_TARGET_PARALLEL_DO:
-      omp_end_st = ST_OMP_END_TARGET_PARALLEL_DO;
-      break;
-    case ST_OMP_TARGET_PARALLEL_DO_SIMD:
-      omp_end_st = ST_OMP_END_TARGET_PARALLEL_DO_SIMD;
-      break;
-    case ST_OMP_TARGET_PARALLEL_LOOP:
-      omp_end_st = ST_OMP_END_TARGET_PARALLEL_LOOP;
-      break;
-    case ST_OMP_TARGET_SIMD: omp_end_st = ST_OMP_END_TARGET_SIMD; break;
-    case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
-      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE;
-      break;
-    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
-      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO;
-      break;
-    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
-      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
-      break;
-    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
-      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD;
-      break;
-    case ST_OMP_TARGET_TEAMS_LOOP:
-      omp_end_st = ST_OMP_END_TARGET_TEAMS_LOOP;
-      break;
-    case ST_OMP_TASKLOOP: omp_end_st = ST_OMP_END_TASKLOOP; break;
-    case ST_OMP_TASKLOOP_SIMD: omp_end_st = ST_OMP_END_TASKLOOP_SIMD; break;
-    case ST_OMP_MASKED_TASKLOOP: omp_end_st = ST_OMP_END_MASKED_TASKLOOP; break;
-    case ST_OMP_MASKED_TASKLOOP_SIMD:
-      omp_end_st = ST_OMP_END_MASKED_TASKLOOP_SIMD;
-      break;
-    case ST_OMP_MASTER_TASKLOOP: omp_end_st = ST_OMP_END_MASTER_TASKLOOP; break;
-    case ST_OMP_MASTER_TASKLOOP_SIMD:
-      omp_end_st = ST_OMP_END_MASTER_TASKLOOP_SIMD;
-      break;
-    case ST_OMP_PARALLEL_MASKED_TASKLOOP:
-      omp_end_st = ST_OMP_END_PARALLEL_MASKED_TASKLOOP;
-      break;
-    case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD:
-      omp_end_st = ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD;
-      break;
-    case ST_OMP_PARALLEL_MASTER_TASKLOOP:
-      omp_end_st = ST_OMP_END_PARALLEL_MASTER_TASKLOOP;
-      break;
-    case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD:
-      omp_end_st = ST_OMP_END_PARALLEL_MASTER_TASKLOOP_SIMD;
-      break;
-    case ST_OMP_TEAMS_DISTRIBUTE:
-      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE;
-      break;
-    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
-      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO;
-      break;
-    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
-      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
-      break;
-    case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
-      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_SIMD;
-      break;
-    case ST_OMP_TEAMS_LOOP: omp_end_st = ST_OMP_END_TEAMS_LOOP; break;
-    case ST_OMP_TILE: omp_end_st = ST_OMP_END_TILE; break;
-    case ST_OMP_UNROLL: omp_end_st = ST_OMP_END_UNROLL; break;
-    default: gcc_unreachable ();
-    }
+  gfc_statement omp_end_st = gfc_omp_end_stmt (omp_st, true, false);
+  if (omp_st == ST_NONE)
+    gcc_unreachable ();
+
+  /* If handling a metadirective variant, treat 'omp end metadirective'
+     as the expected end statement for the current construct.  */
+  if (st == ST_OMP_END_METADIRECTIVE
+      && gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE)
+    st = omp_end_st;
+
   if (st == omp_end_st)
     {
       if (new_st.op == EXEC_OMP_END_NOWAIT)
@@ -5693,7 +5842,10 @@ parse_omp_oacc_atomic (bool omp_p)
   if (omp_p)
     {
       st_atomic = ST_OMP_ATOMIC;
-      st_end_atomic = ST_OMP_END_ATOMIC;
+      if (gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE)
+	st_end_atomic = ST_OMP_END_METADIRECTIVE;
+      else
+	st_end_atomic = ST_OMP_END_ATOMIC;
     }
   else
     {
@@ -5944,7 +6096,10 @@ parse_openmp_allocate_block (gfc_statement omp_st)
   accept_statement (st);
   pop_state ();
   st = next_statement ();
-  if (omp_st == ST_OMP_ALLOCATORS && st == ST_OMP_END_ALLOCATORS)
+  if (omp_st == ST_OMP_ALLOCATORS
+      && (st == ST_OMP_END_ALLOCATORS
+	  || (st == ST_OMP_END_METADIRECTIVE
+	      && gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE)))
     {
       accept_statement (st);
       st = next_statement ();
@@ -5970,80 +6125,15 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
   np->op = cp->op;
   np->block = NULL;
 
-  switch (omp_st)
-    {
-    case ST_OMP_ASSUME:
-      omp_end_st = ST_OMP_END_ASSUME;
-      break;
-    case ST_OMP_PARALLEL:
-      omp_end_st = ST_OMP_END_PARALLEL;
-      break;
-    case ST_OMP_PARALLEL_MASKED:
-      omp_end_st = ST_OMP_END_PARALLEL_MASKED;
-      break;
-    case ST_OMP_PARALLEL_MASTER:
-      omp_end_st = ST_OMP_END_PARALLEL_MASTER;
-      break;
-    case ST_OMP_PARALLEL_SECTIONS:
-      omp_end_st = ST_OMP_END_PARALLEL_SECTIONS;
-      break;
-    case ST_OMP_SCOPE:
-      omp_end_st = ST_OMP_END_SCOPE;
-      break;
-    case ST_OMP_SECTIONS:
-      omp_end_st = ST_OMP_END_SECTIONS;
-      break;
-    case ST_OMP_ORDERED:
-      omp_end_st = ST_OMP_END_ORDERED;
-      break;
-    case ST_OMP_CRITICAL:
-      omp_end_st = ST_OMP_END_CRITICAL;
-      break;
-    case ST_OMP_MASKED:
-      omp_end_st = ST_OMP_END_MASKED;
-      break;
-    case ST_OMP_MASTER:
-      omp_end_st = ST_OMP_END_MASTER;
-      break;
-    case ST_OMP_SINGLE:
-      omp_end_st = ST_OMP_END_SINGLE;
-      break;
-    case ST_OMP_TARGET:
-      omp_end_st = ST_OMP_END_TARGET;
-      break;
-    case ST_OMP_TARGET_DATA:
-      omp_end_st = ST_OMP_END_TARGET_DATA;
-      break;
-    case ST_OMP_TARGET_PARALLEL:
-      omp_end_st = ST_OMP_END_TARGET_PARALLEL;
-      break;
-    case ST_OMP_TARGET_TEAMS:
-      omp_end_st = ST_OMP_END_TARGET_TEAMS;
-      break;
-    case ST_OMP_TASK:
-      omp_end_st = ST_OMP_END_TASK;
-      break;
-    case ST_OMP_TASKGROUP:
-      omp_end_st = ST_OMP_END_TASKGROUP;
-      break;
-    case ST_OMP_TEAMS:
-      omp_end_st = ST_OMP_END_TEAMS;
-      break;
-    case ST_OMP_TEAMS_DISTRIBUTE:
-      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE;
-      break;
-    case ST_OMP_DISTRIBUTE:
-      omp_end_st = ST_OMP_END_DISTRIBUTE;
-      break;
-    case ST_OMP_WORKSHARE:
-      omp_end_st = ST_OMP_END_WORKSHARE;
-      break;
-    case ST_OMP_PARALLEL_WORKSHARE:
-      omp_end_st = ST_OMP_END_PARALLEL_WORKSHARE;
-      break;
-    default:
-      gcc_unreachable ();
-    }
+  omp_end_st = gfc_omp_end_stmt (omp_st, false, true);
+  if (omp_end_st == ST_NONE)
+    gcc_unreachable ();
+
+  /* If handling a metadirective variant, treat 'omp end metadirective'
+     as the expected end statement for the current construct.  */
+  if (gfc_state_stack->previous != NULL
+      && gfc_state_stack->previous->state == COMP_OMP_BEGIN_METADIRECTIVE)
+    omp_end_st = ST_OMP_END_METADIRECTIVE;
 
   bool block_construct = false;
   gfc_namespace *my_ns = NULL;
@@ -6089,11 +6179,13 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
       case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
       case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
       case ST_OMP_TEAMS_LOOP:
+      case ST_OMP_METADIRECTIVE:
+      case ST_OMP_BEGIN_METADIRECTIVE:
 	{
 	  gfc_state_data *stk = gfc_state_stack->previous;
 	  if (stk->state == COMP_OMP_STRICTLY_STRUCTURED_BLOCK)
 	    stk = stk->previous;
-	  stk->tail->ext.omp_clauses->target_first_st_is_teams = true;
+	  stk->tail->ext.omp_clauses->target_first_st_is_teams_or_meta = true;
 	  break;
 	}
       default:
@@ -6266,7 +6358,6 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
   return st;
 }
 
-
 static gfc_statement
 parse_omp_dispatch (void)
 {
@@ -6295,7 +6386,9 @@ parse_omp_dispatch (void)
     }
   pop_state ();
   st = next_statement ();
-  if (st == ST_OMP_END_DISPATCH)
+  if (st == ST_OMP_END_DISPATCH
+      || (st == ST_OMP_END_METADIRECTIVE
+	  && gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE))
     {
       if (cp->ext.omp_clauses->nowait && new_st.ext.omp_bool)
 	gfc_error_now ("Duplicated NOWAIT clause on !$OMP DISPATCH and !$OMP "
@@ -6307,6 +6400,98 @@ parse_omp_dispatch (void)
   return st;
 }
 
+static gfc_statement
+parse_omp_metadirective_body (gfc_statement omp_st)
+{
+  gfc_omp_variant *variant
+    = new_st.ext.omp_variants;
+  locus body_locus = gfc_current_locus;
+
+  accept_statement (omp_st);
+
+  gfc_statement next_st = ST_NONE;
+
+  while (variant)
+    {
+      gfc_current_locus = body_locus;
+      gfc_state_data s;
+      bool workshare_p
+	= (variant->stmt == ST_OMP_WORKSHARE
+	   || variant->stmt == ST_OMP_PARALLEL_WORKSHARE);
+      enum gfc_compile_state new_state
+	= (omp_st == ST_OMP_METADIRECTIVE
+	   ? COMP_OMP_METADIRECTIVE : COMP_OMP_BEGIN_METADIRECTIVE);
+
+      new_st = *variant->code;
+      push_state (&s, new_state, NULL);
+
+      gfc_statement st;
+      bool old_in_metadirective_body = gfc_in_omp_metadirective_body;
+      gfc_in_omp_metadirective_body = true;
+
+      gfc_omp_metadirective_region_count++;
+      switch (variant->stmt)
+	{
+	case_omp_structured_block:
+	  st = parse_omp_structured_block (variant->stmt, workshare_p);
+	  break;
+	case_omp_do:
+	  st = parse_omp_do (variant->stmt, 0);
+	  /* TODO: Does st == ST_IMPLIED_ENDDO need special handling?  */
+	  break;
+	case ST_OMP_ALLOCATORS:
+	  st = parse_openmp_allocate_block (variant->stmt);
+	  break;
+	case ST_OMP_ATOMIC:
+	  st = parse_omp_oacc_atomic (true);
+	  break;
+	case ST_OMP_DISPATCH:
+	  st = parse_omp_dispatch ();
+	  break;
+	default:
+	  accept_statement (variant->stmt);
+	  st = parse_executable (next_statement ());
+	  break;
+	}
+
+      if (gfc_state_stack->state == COMP_OMP_METADIRECTIVE
+	  && startswith (gfc_ascii_statement (st), "!$OMP END "))
+	{
+	  for (gfc_state_data *p = gfc_state_stack; p; p = p->previous)
+	    if (p->state == COMP_OMP_STRUCTURED_BLOCK
+		|| p->state == COMP_OMP_BEGIN_METADIRECTIVE)
+	      goto finish;
+	  gfc_error ("Unexpected %s statement in OMP METADIRECTIVE "
+		     "block at %C",
+		     gfc_ascii_statement (st));
+	  reject_statement ();
+	  st = next_statement ();
+	}
+    finish:
+
+      gfc_in_omp_metadirective_body = old_in_metadirective_body;
+
+      if (gfc_state_stack->head)
+	*variant->code = *gfc_state_stack->head;
+      pop_state ();
+
+      gfc_commit_symbols ();
+      gfc_warning_check ();
+      if (variant->next)
+	gfc_clear_new_st ();
+
+      /* Sanity-check that each variant finishes parsing at the same place.  */
+      if (next_st == ST_NONE)
+	next_st = st;
+      else
+	gcc_assert (st == next_st);
+
+      variant = variant->next;
+    }
+
+  return next_st;
+}
+
 /* Accept a series of executable statements.  We return the first
    statement that doesn't fit to the caller.  Any block statements are
    passed on to the correct handler, which usually passes the buck
@@ -6316,6 +6501,7 @@ static gfc_statement
 parse_executable (gfc_statement st)
 {
   int close_flag;
+  bool one_stmt_p = false;
   in_exec_part = true;
 
   if (st == ST_NONE)
@@ -6323,6 +6509,12 @@ parse_executable (gfc_statement st)
 
   for (;;)
     {
+      /* Only parse one statement for the form of metadirective without
+	 an explicit begin..end.  */
+      if (gfc_state_stack->state == COMP_OMP_METADIRECTIVE && one_stmt_p)
+	return st;
+      one_stmt_p = true;
+
       close_flag = check_do_closure ();
       if (close_flag)
 	switch (st)
@@ -6432,70 +6624,13 @@ parse_executable (gfc_statement st)
 	  st = parse_openmp_allocate_block (st);
 	  continue;
 
-	case ST_OMP_ASSUME:
-	case ST_OMP_PARALLEL:
-	case ST_OMP_PARALLEL_MASKED:
-	case ST_OMP_PARALLEL_MASTER:
-	case ST_OMP_PARALLEL_SECTIONS:
-	case ST_OMP_ORDERED:
-	case ST_OMP_CRITICAL:
-	case ST_OMP_MASKED:
-	case ST_OMP_MASTER:
-	case ST_OMP_SCOPE:
-	case ST_OMP_SECTIONS:
-	case ST_OMP_SINGLE:
-	case ST_OMP_TARGET:
-	case ST_OMP_TARGET_DATA:
-	case ST_OMP_TARGET_PARALLEL:
-	case ST_OMP_TARGET_TEAMS:
-	case ST_OMP_TEAMS:
-	case ST_OMP_TASK:
-	case ST_OMP_TASKGROUP:
-	  st = parse_omp_structured_block (st, false);
+	case_omp_structured_block:
+	  st = parse_omp_structured_block (st,
+					   st == ST_OMP_WORKSHARE
+					   || st == ST_OMP_PARALLEL_WORKSHARE);
 	  continue;
 
-	case ST_OMP_WORKSHARE:
-	case ST_OMP_PARALLEL_WORKSHARE:
-	  st = parse_omp_structured_block (st, true);
-	  continue;
-
-	case ST_OMP_DISTRIBUTE:
-	case ST_OMP_DISTRIBUTE_PARALLEL_DO:
-	case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
-	case ST_OMP_DISTRIBUTE_SIMD:
-	case ST_OMP_DO:
-	case ST_OMP_DO_SIMD:
-	case ST_OMP_LOOP:
-	case ST_OMP_PARALLEL_DO:
-	case ST_OMP_PARALLEL_DO_SIMD:
-	case ST_OMP_PARALLEL_LOOP:
-	case ST_OMP_PARALLEL_MASKED_TASKLOOP:
-	case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD:
-	case ST_OMP_PARALLEL_MASTER_TASKLOOP:
-	case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD:
-	case ST_OMP_MASKED_TASKLOOP:
-	case ST_OMP_MASKED_TASKLOOP_SIMD:
-	case ST_OMP_MASTER_TASKLOOP:
-	case ST_OMP_MASTER_TASKLOOP_SIMD:
-	case ST_OMP_SIMD:
-	case ST_OMP_TARGET_PARALLEL_DO:
-	case ST_OMP_TARGET_PARALLEL_DO_SIMD:
-	case ST_OMP_TARGET_PARALLEL_LOOP:
-	case ST_OMP_TARGET_SIMD:
-	case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
-	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
-	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
-	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
-	case ST_OMP_TARGET_TEAMS_LOOP:
-	case ST_OMP_TASKLOOP:
-	case ST_OMP_TASKLOOP_SIMD:
-	case ST_OMP_TEAMS_DISTRIBUTE:
-	case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
-	case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
-	case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
-	case ST_OMP_TEAMS_LOOP:
-	case ST_OMP_TILE:
-	case ST_OMP_UNROLL:
+	case_omp_do:
 	  st = parse_omp_do (st, 0);
 	  if (st == ST_IMPLIED_ENDDO)
 	    return st;
@@ -6513,6 +6648,17 @@ parse_executable (gfc_statement st)
 	  st = parse_omp_dispatch ();
 	  continue;
 
+	case ST_OMP_METADIRECTIVE:
+	case ST_OMP_BEGIN_METADIRECTIVE:
+	  st = parse_omp_metadirective_body (st);
+	  continue;
+
+	case ST_OMP_END_METADIRECTIVE:
+	  if (gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE)
+	    return next_statement ();
+	  else
+	    return st;
+
 	default:
 	  return st;
 	}
@@ -7278,6 +7424,10 @@ gfc_parse_file (void)
 
   gfc_statement_label = NULL;
 
+  gfc_omp_metadirective_region_count = 0;
+  gfc_in_omp_metadirective_body = false;
+  gfc_matching_omp_context_selector = false;
+
   if (setjmp (eof_buf))
     return false;	/* Come here on unexpected EOF */
 
@@ -7589,3 +7739,16 @@ is_oacc (gfc_state_data *sd)
       return false;
     }
 }
+
+/* Return true if ST is a declarative OpenMP statement.  */
+bool
+is_omp_declarative_stmt (gfc_statement st)
+{
+  switch (st)
+    {
+      case_omp_decl:
+	return true;
+      default:
+	return false;
+    }
+}
diff --git a/gcc/fortran/parse.h b/gcc/fortran/parse.h
index 448fa0f..722e94c 100644
--- a/gcc/fortran/parse.h
+++ b/gcc/fortran/parse.h
@@ -31,7 +31,8 @@ enum gfc_compile_state
   COMP_STRUCTURE, COMP_UNION, COMP_MAP,
   COMP_DO, COMP_SELECT, COMP_FORALL, COMP_WHERE, COMP_CONTAINS, COMP_ENUM,
   COMP_SELECT_TYPE, COMP_SELECT_RANK, COMP_OMP_STRUCTURED_BLOCK, COMP_CRITICAL,
-  COMP_DO_CONCURRENT, COMP_OMP_STRICTLY_STRUCTURED_BLOCK
+  COMP_DO_CONCURRENT, COMP_OMP_STRICTLY_STRUCTURED_BLOCK,
+  COMP_OMP_METADIRECTIVE, COMP_OMP_BEGIN_METADIRECTIVE
 };
 
 /* Stack element for the current compilation state.  These structures
@@ -67,10 +68,15 @@ bool gfc_check_do_variable (gfc_symtree *);
 bool gfc_find_state (gfc_compile_state);
 gfc_state_data *gfc_enclosing_unit (gfc_compile_state *);
 const char *gfc_ascii_statement (gfc_statement, bool strip_sentinel = false) ;
+gfc_statement gfc_omp_end_stmt (gfc_statement, bool = true, bool = true);
 match gfc_match_enum (void);
 match gfc_match_enumerator_def (void);
 void gfc_free_enum_history (void);
 extern bool gfc_matching_function;
+extern bool gfc_matching_omp_context_selector;
+extern bool gfc_in_omp_metadirective_body;
+extern int gfc_omp_metadirective_region_count;
+
 match gfc_match_prefix (gfc_typespec *);
 bool is_oacc (gfc_state_data *);
 #endif  /* GFC_PARSE_H  */
diff --git a/gcc/fortran/resolve.cc b/gcc/fortran/resolve.cc
index 7f73d53..12a623da 100644
--- a/gcc/fortran/resolve.cc
+++ b/gcc/fortran/resolve.cc
@@ -13806,6 +13806,11 @@ gfc_resolve_code (gfc_code *code, gfc_namespace *ns)
 	  gfc_resolve_forall (code, ns, forall_save);
 	  forall_flag = 2;
 	}
+      else if (code->op == EXEC_OMP_METADIRECTIVE)
+	for (gfc_omp_variant *variant
+	       = code->ext.omp_variants;
+	     variant; variant = variant->next)
+	  gfc_resolve_code (variant->code, ns);
       else if (code->block)
 	{
 	  omp_workshare_save = -1;
@@ -14379,6 +14384,7 @@ start:
 	case EXEC_OMP_MASKED:
 	case EXEC_OMP_MASKED_TASKLOOP:
 	case EXEC_OMP_MASKED_TASKLOOP_SIMD:
+	case EXEC_OMP_METADIRECTIVE:
 	case EXEC_OMP_ORDERED:
 	case EXEC_OMP_SCAN:
 	case EXEC_OMP_SCOPE:
diff --git a/gcc/fortran/st.cc b/gcc/fortran/st.cc
index 509d28c..f7f67b1 100644
--- a/gcc/fortran/st.cc
+++ b/gcc/fortran/st.cc
@@ -306,6 +306,10 @@ gfc_free_statement (gfc_code *p)
     case EXEC_OMP_TASKYIELD:
       break;
 
+    case EXEC_OMP_METADIRECTIVE:
+      gfc_free_omp_variants (p->ext.omp_variants);
+      break;
+
     default:
       gfc_internal_error ("gfc_free_statement(): Bad statement");
     }
diff --git a/gcc/fortran/symbol.cc b/gcc/fortran/symbol.cc
index e6535fa..c689481 100644
--- a/gcc/fortran/symbol.cc
+++ b/gcc/fortran/symbol.cc
@@ -2697,10 +2697,13 @@ free_components (gfc_component *p)
 static int
 compare_st_labels (void *a1, void *b1)
 {
-  int a = ((gfc_st_label *) a1)->value;
-  int b = ((gfc_st_label *) b1)->value;
+  gfc_st_label *a = (gfc_st_label *) a1;
+  gfc_st_label *b = (gfc_st_label *) b1;
 
-  return (b - a);
+  if (a->omp_region == b->omp_region)
+    return b->value - a->value;
+  else
+    return b->omp_region - a->omp_region;
 }
 
 
@@ -2750,6 +2753,8 @@ gfc_get_st_label (int labelno)
 {
   gfc_st_label *lp;
   gfc_namespace *ns;
+  int omp_region = (gfc_in_omp_metadirective_body
+		    ? gfc_omp_metadirective_region_count : 0);
 
   if (gfc_current_state () == COMP_DERIVED)
     ns = gfc_current_block ()->f2k_derived;
@@ -2766,10 +2771,16 @@ gfc_get_st_label (int labelno)
   lp = ns->st_labels;
   while (lp)
     {
-      if (lp->value == labelno)
-	return lp;
-
-      if (lp->value < labelno)
+      if (lp->omp_region == omp_region)
+	{
+	  if (lp->value == labelno)
+	    return lp;
+	  if (lp->value < labelno)
+	    lp = lp->left;
+	  else
+	    lp = lp->right;
+	}
+      else if (lp->omp_region < omp_region)
 	lp = lp->left;
       else
 	lp = lp->right;
@@ -2781,6 +2792,7 @@ gfc_get_st_label (int labelno)
   lp->defined = ST_LABEL_UNKNOWN;
   lp->referenced = ST_LABEL_UNKNOWN;
   lp->ns = ns;
+  lp->omp_region = omp_region;
 
   gfc_insert_bbt (&ns->st_labels, lp, compare_st_labels);
 
diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc
index b8fc9a1..f6a65cf 100644
--- a/gcc/fortran/trans-decl.cc
+++ b/gcc/fortran/trans-decl.cc
@@ -342,7 +342,10 @@ gfc_get_label_decl (gfc_st_label * lp)
       gcc_assert (lp != NULL && lp->value <= MAX_LABEL_VALUE);
 
       /* Build a mangled name for the label.  */
-      sprintf (label_name, "__label_%.6d", lp->value);
+      if (lp->omp_region)
+	sprintf (label_name, "__label_%d_%.6d", lp->omp_region, lp->value);
+      else
+	sprintf (label_name, "__label_%.6d", lp->value);
 
       /* Build the LABEL_DECL node.  */
       label_decl = gfc_build_label_decl (get_identifier (label_name));
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 160dc84..a593f5a 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -8494,6 +8494,8 @@ gfc_trans_omp_directive (gfc_code *code)
     case EXEC_OMP_MASTER_TASKLOOP:
     case EXEC_OMP_MASTER_TASKLOOP_SIMD:
       return gfc_trans_omp_master_masked_taskloop (code, code->op);
+    case EXEC_OMP_METADIRECTIVE:
+      return gfc_trans_omp_metadirective (code);
     case EXEC_OMP_ORDERED:
       return gfc_trans_omp_ordered (code);
     case EXEC_OMP_PARALLEL:
@@ -8587,6 +8589,100 @@ gfc_trans_omp_declare_simd (gfc_namespace *ns)
     }
 }
 
+/* Translate the context selector list GFC_SELECTORS, using WHERE as the
+   locus for error messages.  */
+
+static tree
+gfc_trans_omp_set_selector (gfc_omp_set_selector *gfc_selectors, locus where)
+{
+  tree set_selectors = NULL_TREE;
+  gfc_omp_set_selector *oss;
+
+  for (oss = gfc_selectors; oss; oss = oss->next)
+    {
+      tree selectors = NULL_TREE;
+      gfc_omp_selector *os;
+      enum omp_tss_code set = oss->code;
+      gcc_assert (set != OMP_TRAIT_SET_INVALID);
+
+      for (os = oss->trait_selectors; os; os = os->next)
+	{
+	  tree scoreval = NULL_TREE;
+	  tree properties = NULL_TREE;
+	  gfc_omp_trait_property *otp;
+	  enum omp_ts_code sel = os->code;
+
+	  /* Per the spec, "Implementations can ignore specified
+	     selectors that are not those described in this section";
+	     however, we  must record such selectors because they
+	     cause match failures.  */
+	  if (sel == OMP_TRAIT_INVALID)
+	    {
+	      selectors = make_trait_selector (sel, NULL_TREE, NULL_TREE,
+					       selectors);
+	      continue;
+	    }
+
+	  for (otp = os->properties; otp; otp = otp->next)
+	    {
+	      switch (otp->property_kind)
+		{
+		case OMP_TRAIT_PROPERTY_DEV_NUM_EXPR:
+		case OMP_TRAIT_PROPERTY_BOOL_EXPR:
+		  {
+		    tree expr = NULL_TREE;
+		    gfc_se se;
+		    gfc_init_se (&se, NULL);
+		    gfc_conv_expr (&se, otp->expr);
+		    expr = se.expr;
+		    properties = make_trait_property (NULL_TREE, expr,
+						      properties);
+		  }
+		  break;
+		case OMP_TRAIT_PROPERTY_ID:
+		  properties
+		    = make_trait_property (get_identifier (otp->name),
+					   NULL_TREE, properties);
+		  break;
+		case OMP_TRAIT_PROPERTY_NAME_LIST:
+		  {
+		    tree prop = OMP_TP_NAMELIST_NODE;
+		    tree value = NULL_TREE;
+		    if (otp->is_name)
+		      value = get_identifier (otp->name);
+		    else
+		      value = gfc_conv_constant_to_tree (otp->expr);
+
+		    properties = make_trait_property (prop, value,
+						      properties);
+		  }
+		  break;
+		case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
+		  properties = gfc_trans_omp_clauses (NULL, otp->clauses,
+						      where, true);
+		  break;
+		default:
+		  gcc_unreachable ();
+		}
+	    }
+
+	  if (os->score)
+	    {
+	      gfc_se se;
+	      gfc_init_se (&se, NULL);
+	      gfc_conv_expr (&se, os->score);
+	      scoreval = se.expr;
+	    }
+
+	  selectors = make_trait_selector (sel, scoreval,
+					   properties, selectors);
+	}
+      set_selectors = make_trait_set_selector (set, selectors, set_selectors);
+    }
+  return set_selectors;
+}
+
+
 void
 gfc_trans_omp_declare_variant (gfc_namespace *ns)
 {
@@ -8662,90 +8758,8 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
 	      && strcmp (odv->base_proc_symtree->name, ns->proc_name->name)))
 	continue;
 
-      tree set_selectors = NULL_TREE;
-      gfc_omp_set_selector *oss;
-
-      for (oss = odv->set_selectors; oss; oss = oss->next)
-	{
-	  tree selectors = NULL_TREE;
-	  gfc_omp_selector *os;
-	  enum omp_tss_code set = oss->code;
-	  gcc_assert (set != OMP_TRAIT_SET_INVALID);
-
-	  for (os = oss->trait_selectors; os; os = os->next)
-	    {
-	      tree scoreval = NULL_TREE;
-	      tree properties = NULL_TREE;
-	      gfc_omp_trait_property *otp;
-	      enum omp_ts_code sel = os->code;
-
-	      /* Per the spec, "Implementations can ignore specified
-		 selectors that are not those described in this section";
-		 however, we  must record such selectors because they
-		 cause match failures.  */
-	      if (sel == OMP_TRAIT_INVALID)
-		{
-		  selectors = make_trait_selector (sel, NULL_TREE, NULL_TREE,
-						   selectors);
-		  continue;
-		}
-
-	      for (otp = os->properties; otp; otp = otp->next)
-		{
-		  switch (otp->property_kind)
-		    {
-		    case OMP_TRAIT_PROPERTY_DEV_NUM_EXPR:
-		    case OMP_TRAIT_PROPERTY_BOOL_EXPR:
-		      {
-			gfc_se se;
-			gfc_init_se (&se, NULL);
-			gfc_conv_expr (&se, otp->expr);
-			properties = make_trait_property (NULL_TREE, se.expr,
-							  properties);
-		      }
-		      break;
-		    case OMP_TRAIT_PROPERTY_ID:
-		      properties
-			= make_trait_property (get_identifier (otp->name),
-					       NULL_TREE, properties);
-		      break;
-		    case OMP_TRAIT_PROPERTY_NAME_LIST:
-		      {
-			tree prop = OMP_TP_NAMELIST_NODE;
-			tree value = NULL_TREE;
-			if (otp->is_name)
-			  value = get_identifier (otp->name);
-			else
-			  value = gfc_conv_constant_to_tree (otp->expr);
-
-			properties = make_trait_property (prop, value,
-							  properties);
-		      }
-		      break;
-		    case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
-		      properties = gfc_trans_omp_clauses (NULL, otp->clauses,
-							  odv->where, true);
-		      break;
-		    default:
-		      gcc_unreachable ();
-		    }
-		}
-
-	      if (os->score)
-		{
-		  gfc_se se;
-		  gfc_init_se (&se, NULL);
-		  gfc_conv_expr (&se, os->score);
-		  scoreval = se.expr;
-		}
-
-	      selectors	= make_trait_selector (sel, scoreval,
-					       properties, selectors);
-	    }
-	  set_selectors = make_trait_set_selector (set, selectors,
-						   set_selectors);
-	}
-
+      tree set_selectors = gfc_trans_omp_set_selector (odv->set_selectors,
+						       odv->where);
       const char *variant_proc_name = odv->variant_proc_symtree->name;
       gfc_symbol *variant_proc_sym = odv->variant_proc_symtree->n.sym;
       if (variant_proc_sym == NULL || variant_proc_sym->attr.implicit_type)
@@ -9048,3 +9062,54 @@ gfc_omp_call_is_alloc (tree ptr)
     }
   return build_call_expr_loc (input_location, fn, 1, ptr);
 }
+
+tree
+gfc_trans_omp_metadirective (gfc_code *code)
+{
+  gfc_omp_variant *variant = code->ext.omp_variants;
+
+  tree metadirective_tree = make_node (OMP_METADIRECTIVE);
+  SET_EXPR_LOCATION (metadirective_tree, gfc_get_location (&code->loc));
+  TREE_TYPE (metadirective_tree) = void_type_node;
+  OMP_METADIRECTIVE_VARIANTS (metadirective_tree) = NULL_TREE;
+
+  tree tree_body = NULL_TREE;
+
+  while (variant)
+    {
+      tree ctx = gfc_trans_omp_set_selector (variant->selectors,
+					     variant->where);
+      ctx = omp_check_context_selector (gfc_get_location (&variant->where),
+					ctx, true);
+      if (ctx == error_mark_node)
+	return error_mark_node;
+
+      /* If the selector doesn't match, drop the whole variant.  */
+      if (!omp_context_selector_matches (ctx, NULL_TREE, false))
+	{
+	  variant = variant->next;
+	  continue;
+	}
+
+      gfc_code *next_code = variant->code->next;
+      if (next_code && tree_body == NULL_TREE)
+	tree_body = gfc_trans_code (next_code);
+
+      if (next_code)
+	variant->code->next = NULL;
+      tree directive = gfc_trans_code (variant->code);
+      if (next_code)
+	variant->code->next = next_code;
+
+      tree body = next_code ? tree_body : NULL_TREE;
+      tree omp_variant = make_omp_metadirective_variant (ctx, directive, body);
+      OMP_METADIRECTIVE_VARIANTS (metadirective_tree)
+	= chainon (OMP_METADIRECTIVE_VARIANTS (metadirective_tree),
+		   omp_variant);
+      variant = variant->next;
+    }
+
+  /* TODO: Resolve the metadirective here if possible.   */
+
+  return metadirective_tree;
+}
diff --git a/gcc/fortran/trans-stmt.h b/gcc/fortran/trans-stmt.h
index 544c2f9..36cabaf 100644
--- a/gcc/fortran/trans-stmt.h
+++ b/gcc/fortran/trans-stmt.h
@@ -71,6 +71,7 @@ tree gfc_trans_deallocate (gfc_code *);
 tree gfc_trans_omp_directive (gfc_code *);
 void gfc_trans_omp_declare_simd (gfc_namespace *);
 void gfc_trans_omp_declare_variant (gfc_namespace *);
+tree gfc_trans_omp_metadirective (gfc_code *code);
 tree gfc_trans_oacc_directive (gfc_code *);
 tree gfc_trans_oacc_declare (gfc_namespace *);
 
diff --git a/gcc/fortran/trans.cc b/gcc/fortran/trans.cc
index 3834986..b03dcc1 100644
--- a/gcc/fortran/trans.cc
+++ b/gcc/fortran/trans.cc
@@ -2588,6 +2588,7 @@ trans_code (gfc_code * code, tree cond)
 	case EXEC_OMP_MASTER:
 	case EXEC_OMP_MASTER_TASKLOOP:
 	case EXEC_OMP_MASTER_TASKLOOP_SIMD:
+	case EXEC_OMP_METADIRECTIVE:
 	case EXEC_OMP_ORDERED:
 	case EXEC_OMP_PARALLEL:
 	case EXEC_OMP_PARALLEL_DO:
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90
new file mode 100644
index 0000000..15671f2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90
@@ -0,0 +1,80 @@
+! { dg-do compile }
+
+program main
+  integer, parameter :: N = 10
+  integer, dimension(N) :: a
+  integer, dimension(N) :: b
+  integer, dimension(N) :: c
+  integer :: i
+
+  do i = 1, N
+    a(i) = i * 2
+    b(i) = i * 3
+  end do
+
+  !$omp metadirective &
+  !$omp&	default (teams loop) &
+  !$omp&	default (parallel loop)	! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." }
+    do i = 1, N
+      c(i) = a(i) * b(i)
+    end do
+
+  !$omp metadirective &
+  !$omp&	otherwise (teams loop) &
+  !$omp&	default (parallel loop)	! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." }
+    do i = 1, N
+      c(i) = a(i) * b(i)
+    end do
+
+  !$omp metadirective &
+  !$omp&	otherwise (teams loop) &
+  !$omp&	otherwise (parallel loop)	! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." }
+    do i = 1, N
+      c(i) = a(i) * b(i)
+    end do
+
+  !$omp metadirective default (xyz) ! { dg-error "Unclassifiable OpenMP directive at .1." }
+    do i = 1, N
+      c(i) = a(i) * b(i)
+    end do
+
+  !$omp metadirective &
+  !$omp&	default (teams loop) &
+  !$omp&	where (device={arch("nvptx")}: parallel loop) ! { dg-error "expected 'when', 'otherwise', or 'default' at .1." }
+    do i = 1, N
+      c(i) = a(i) * b(i)
+    end do
+
+  !$omp metadirective &
+  !$omp&	otherwise (teams loop) &
+  !$omp&	when (device={arch("nvptx")}: parallel loop) ! { dg-error "'otherwise' or 'default' clause must appear last" }
+    do i = 1, N
+      c(i) = a(i) * b(i)
+    end do
+
+  !$omp metadirective &
+  !$omp&	when (device={arch("nvptx")} parallel loop) & ! { dg-error "expected .:." } 
+  !$omp&	default (teams loop)
+    do i = 1, N
+      c(i) = a(i) * b(i)
+    end do
+
+  ! Test improperly nested metadirectives - even though the second
+  ! metadirective resolves to 'omp nothing', that is not the same as there
+  ! being literally nothing there.
+  !$omp metadirective &
+  !$omp&    when (implementation={vendor("gnu")}: parallel do)
+    !$omp metadirective &
+    !$omp& when (implementation={vendor("cray")}: parallel do) ! { dg-error "Unexpected !.OMP METADIRECTIVE statement" }
+      do i = 1, N
+        c(i) = a(i) * b(i)
+      end do
+
+!$omp begin metadirective &
+  !$omp&	when (device={arch("nvptx")}: parallel do) &
+  !$omp&	default (barrier) ! { dg-error "variant directive used in OMP BEGIN METADIRECTIVE at .1. must have a corresponding end directive" }
+    do i = 1, N
+      c(i) = a(i) * b(i)
+    end do
+  !$omp end metadirective ! { dg-error "Unexpected !.OMP END METADIRECTIVE statement at .1." }
+end program
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-10.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-10.f90
new file mode 100644
index 0000000..5dad5d2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-10.f90
@@ -0,0 +1,40 @@
+! { dg-do compile }
+
+program metadirectives
+   implicit none
+   logical :: UseDevice
+
+   !$OMP metadirective &
+   !$OMP   when ( user = { condition ( UseDevice ) } &
+   !$OMP     : parallel ) &
+   !$OMP   default ( parallel )
+   block
+      call bar()
+   end block
+
+   !$OMP metadirective &
+   !$OMP   when ( user = { condition ( UseDevice ) } &
+   !$OMP     : parallel ) &
+   !$OMP   default ( parallel )
+   call bar()
+   !$omp end parallel  ! Accepted, because all cases have 'parallel'
+   
+   !$OMP begin metadirective &
+   !$OMP   when ( user = { condition ( UseDevice ) } &
+   !$OMP     : nothing ) &
+   !$OMP   default ( parallel )
+   call bar()
+   block
+      call foo()
+   end block
+   !$OMP end metadirective
+
+   !$OMP begin metadirective &
+   !$OMP   when ( user = { condition ( UseDevice ) } &
+   !$OMP     : parallel ) &
+   !$OMP   default ( parallel )
+   call bar()
+   !$omp end parallel  ! { dg-error "Unexpected !.OMP END PARALLEL statement at .1." }
+end program ! { dg-error "Unexpected END statement at .1." }
+
+! { dg-error "Unexpected end of file" "" { target *-*-* } 0 }
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-11.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-11.f90
new file mode 100644
index 0000000..e7de70e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-11.f90
@@ -0,0 +1,33 @@
+! { dg-do compile }
+! { dg-ice "Statements following a block in a metadirective" }
+! PR fortran/107067
+
+program metadirectives
+   implicit none
+   logical :: UseDevice
+
+   !$OMP begin metadirective &
+   !$OMP   when ( user = { condition ( UseDevice ) } &
+   !$OMP     : nothing ) &
+   !$OMP   default ( parallel )
+   block
+      call foo()
+   end block
+   call bar()   ! FIXME/XFAIL ICE in parse_omp_metadirective_body()
+   !$omp end metadirective
+
+
+   !$OMP begin metadirective &
+   !$OMP   when ( user = { condition ( UseDevice ) } &
+   !$OMP     : nothing ) &
+   !$OMP   default ( parallel )
+   block
+      call bar()
+   end block
+   block        ! FIXME/XFAIL ICE in parse_omp_metadirective_body()
+      call foo()
+   end block
+   !$omp end metadirective
+end program
+
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-12.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-12.f90
new file mode 100644
index 0000000..fc122cc
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-12.f90
@@ -0,0 +1,18 @@
+! { dg-do compile }
+
+! PR112779 item H; this testcase used to ICE.
+
+program test
+  implicit none
+  integer, parameter :: N = 100
+  integer :: x(N), y(N), z(N)
+  block
+    integer :: i
+    !$omp metadirective &
+                !$omp& when(device={arch("nvptx")}: teams loop) &
+                !$omp& default(parallel loop)
+    do i = 1, N
+          z(i) = x(i) * y(i)
+    enddo
+   end block
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-13.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-13.f90
new file mode 100644
index 0000000..bc69f65
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-13.f90
@@ -0,0 +1,30 @@
+! { dg-do compile }
+
+subroutine foo
+  implicit none
+  external f
+
+  !$omp dispatch
+  call f()
+  !$omp dispatch
+    call f()
+  !$omp end dispatch
+
+  !$omp begin metadirective when(construct={parallel} : nothing) otherwise(dispatch)
+  call f()
+  !$omp end metadirective
+end
+
+subroutine bar
+  implicit none
+  integer :: x
+  !$omp atomic update
+    x = x + 1
+  !$omp atomic update
+    x = x + 1
+  !$omp end atomic
+
+  !$omp begin metadirective when(construct={parallel} : nothing) otherwise(atomic update)
+    x = x + 1
+  !$omp end metadirective
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90
new file mode 100644
index 0000000..bd0c382
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90
@@ -0,0 +1,72 @@
+! { dg-do compile }
+
+program main
+  integer, parameter :: N = 100
+  integer :: x = 0
+  integer :: y = 0
+  integer :: i
+
+  ! Test implicit default directive
+  !$omp metadirective &
+  !$omp&	when (device={arch("nvptx")}: barrier)
+    x = 1
+
+  ! Test implicit default directive combined with a directive that takes a
+  ! do loop.
+  !$omp metadirective &
+  !$omp&	when (device={arch("nvptx")}: parallel do)
+    do i = 1, N
+      x = x + i
+    end do
+
+  ! Test with multiple standalone directives.
+  !$omp metadirective &
+  !$omp&	when (device={arch("nvptx")}: barrier) &
+  !$omp&	default (flush)
+    x = 1
+
+  ! Test combining a standalone directive with one that takes a do loop.
+  !$omp metadirective &
+  !$omp&	when (device={arch("nvptx")}: parallel do) &
+  !$omp&	default (barrier)
+    do i = 1, N
+      x = x + i
+    end do
+
+  ! Test combining a directive that takes a do loop with one that takes
+  ! a statement body.
+  !$omp begin metadirective &
+  !$omp&	when (device={arch("nvptx")}: parallel do) &
+  !$omp&	default (parallel)
+    do i = 1, N
+      x = x + i
+    end do
+  !$omp end metadirective
+  
+  ! Test labels in the body.
+  !$omp begin metadirective &
+  !$omp&	when (device={arch("nvptx")}: parallel do) &
+  !$omp&	when (device={arch("gcn")}: parallel)
+    do i = 1, N
+      x = x + i
+      if (x .gt. N/2) goto 10
+10    x = x + 1
+      goto 20
+      x = x + 2
+20    continue
+    end do
+  !$omp end metadirective
+
+  ! Test that commas are permitted before each clause.
+  !$omp begin metadirective, &
+  !$omp&	when (device={arch("nvptx")}: parallel do) &
+  !$omp&	, when (device={arch("gcn")}: parallel) &
+  !$omp&	, default (parallel)
+    do i = 1, N
+      x = x + i
+    end do
+  !$omp end metadirective
+
+  ! Test empty metadirective.
+  !$omp metadirective
+end program
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90
new file mode 100644
index 0000000..c5e25e5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90
@@ -0,0 +1,25 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module test
+  integer, parameter :: N = 100
+contains
+  subroutine f (x, y, z)
+    integer :: x(N), y(N), z(N)
+
+    !$omp target map (to: v1, v2) map(from: v3)
+      !$omp metadirective &
+		!$omp& when(device={arch("nvptx")}: teams loop) &
+		!$omp& default(parallel loop)
+	do i = 1, N
+	  z(i) = x(i) * y(i)
+	enddo
+    !$omp end target
+  end subroutine
+end module
+
+! If offload device "nvptx" isn't supported, the front end can eliminate
+!  that alternative and not produce a metadirective at all.  Otherwise this
+!  won't be resolved until late.
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" { target { ! offload_nvptx } } } }
+! { dg-final { scan-tree-dump "#pragma omp metadirective" "gimple" { target { offload_nvptx } } } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90
new file mode 100644
index 0000000..1da4a0c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90
@@ -0,0 +1,37 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program test
+  implicit none
+  integer, parameter :: N = 100
+  real :: a(N)
+  
+  !$omp target map(from: a)
+    call f (a, 3.14159)
+  !$omp end target
+
+  call f (a, 2.71828)
+contains
+  subroutine f (a, x)
+    integer :: i
+    real :: a(N), x
+    !$omp declare target
+
+    !$omp metadirective &
+    !$omp&  when (construct={target}: distribute parallel do ) &
+    !$omp&  default(parallel do simd)
+      do i = 1, N
+ 	a(i) = x * i
+      end do
+  end subroutine
+end program
+
+! The metadirective should be resolved during Gimplification.
+
+! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } }
+! { dg-final { scan-tree-dump-times "when \\(construct = .*target.*\\):" 1 "original" } }
+! { dg-final { scan-tree-dump-times "otherwise:" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } }
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90
new file mode 100644
index 0000000..0397039
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90
@@ -0,0 +1,30 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module test
+  integer, parameter :: N = 100
+contains
+  subroutine f (a, flag)
+    integer :: a(N)
+    logical :: flag
+    integer :: i
+    
+   !$omp metadirective &
+   !$omp&  when (user={condition(flag)}: &
+   !$omp&	 target teams distribute parallel do map(from: a(1:N))) &
+   !$omp&  default(parallel do)
+     do i = 1, N
+       a(i) = i
+     end do
+  end subroutine
+end module
+
+! The metadirective should be resolved at parse time, but is currently
+! resolved during Gimplification
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times  "#pragma omp distribute" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90
new file mode 100644
index 0000000..a38d8cf
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90
@@ -0,0 +1,26 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+module test
+  integer, parameter :: N = 100
+contains
+  subroutine f (a, run_parallel, run_guided)
+    integer :: a(N)
+    logical :: run_parallel, run_guided
+    integer :: i
+
+    !$omp begin metadirective when(user={condition(run_parallel)}: parallel)
+      !$omp metadirective &
+      !$omp&  when(construct={parallel}, user={condition(run_guided)}: &
+      !$omp&       do schedule(guided)) &
+      !$omp&  when(construct={parallel}: do schedule(static))
+	do i = 1, N
+	  a(i) = i
+	end do
+    !$omp end metadirective
+  end subroutine
+end module
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90
new file mode 100644
index 0000000..37825e6
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90
@@ -0,0 +1,42 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple -fdump-tree-ompdevlow" }
+
+subroutine f (a, num)
+  integer, parameter :: N = 256
+  integer :: a(N)
+  integer :: num
+  integer :: i
+
+  !$omp metadirective &
+  !$omp& when (target_device={device_num(num), kind("gpu"), arch("nvptx")}: &
+  !$omp&       target parallel do map(tofrom: a(1:N))) &
+  !$omp& when (target_device={device_num(num), kind("gpu"), &
+  !$omp&                      arch("amdgcn"), isa("gfx906")}: &
+  !$omp&       target parallel do) &
+  !$omp& when (target_device={device_num(num), kind("cpu"), arch("x86_64")}: &
+  !$omp&       parallel do)
+    do i = 1, N
+      a(i) = a(i) + i
+    end do
+
+  !$omp metadirective &
+  !$omp& when (target_device={kind("gpu"), arch("nvptx")}: &
+  !$omp&       target parallel do map(tofrom: a(1:N)))
+    do i = 1, N
+      a(i) = a(i) + i
+    end do
+end subroutine
+
+! For configurations with offloading, we expect one "pragma omp target"
+! with "device(num)" for each target_device selector that specifies
+! "device_num(num)".  Without offloading, there should be zero as the
+!  resolution happens during gimplification.
+! { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(" 3 "gimple" { target offloading_enabled } } }
+! { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(" 0 "gimple" { target { ! offloading_enabled } } } }
+
+! For configurations with offloading, expect one OMP_TARGET_DEVICE_MATCHES
+! for each kind/arch/isa selector.  These are supposed to go away after
+!  ompdevlow.
+! { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 9 "gimple" { target offloading_enabled } } }
+! { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "gimple" { target { ! offloading_enabled } } } }
+! { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "ompdevlow" { target offloading_enabled } } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90
new file mode 100644
index 0000000..1ebcd33
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90
@@ -0,0 +1,22 @@
+! { dg-do compile }
+
+program test
+  integer :: i
+  integer, parameter :: N = 100
+  integer :: sum = 0
+  
+  ! The compiler should never consider a situation where both metadirectives
+  ! match, but that does not matter because the spec says "Replacement of
+  ! the metadirective with the directive variant associated with any of the
+  ! dynamic replacement candidates must result in a conforming OpenMP
+  ! program.  So the second metadirective is rejected as not being
+  ! a valid loop-nest even if the first one does not match.
+  
+!$omp metadirective when (implementation={vendor("ibm")}: &
+  !$omp&  target teams distribute)
+    !$omp metadirective when (implementation={vendor("gnu")}: parallel do) ! { dg-error "Unexpected !.OMP METADIRECTIVE statement" }
+      do i = 1, N
+	sum = sum + i
+      end do
+end program
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-9.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-9.f90
new file mode 100644
index 0000000..9a63de8
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-9.f90
@@ -0,0 +1,30 @@
+! { dg-do compile }
+
+program OpenMP_Metadirective_WrongEnd_Test
+  implicit none
+
+  integer :: &
+  iaVS, iV, jV, kV
+  integer, dimension ( 3 ) :: &
+    lV, uV
+  logical :: &
+    UseDevice
+
+    !$OMP metadirective &
+    !$OMP   when ( user = { condition ( UseDevice ) } &
+    !$OMP     : target teams distribute parallel do simd collapse ( 3 ) &
+    !$OMP         private ( iaVS ) ) &
+    !$OMP   default ( parallel do simd collapse ( 3 ) private ( iaVS ) )
+    do kV = lV ( 3 ), uV ( 3 )
+      do jV = lV ( 2 ), uV ( 2 )
+        do iV = lV ( 1 ), uV ( 1 )
+
+
+        end do
+      end do
+    end do
+    !$OMP end target teams distribute parallel do simd ! { dg-error "Unexpected !.OMP END TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD statement in OMP METADIRECTIVE block at .1." }
+
+
+end program
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-construct.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-construct.f90
new file mode 100644
index 0000000..ec1f0ee
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-construct.f90
@@ -0,0 +1,260 @@
+! { dg-do compile }
+! { dg-additional-options "-foffload=disable -fdump-tree-original -fdump-tree-gimple" }
+
+program main
+implicit none
+
+integer, parameter :: N = 10
+double precision, parameter :: S = 2.0
+double precision :: a(N)
+
+call init (N, a)
+call f1 (N, a, S)
+call check (N, a, S)
+
+call init (N, a)
+call f2 (N, a, S)
+call check (N, a, S)
+
+call init (N, a)
+call f3 (N, a, S)
+call check (N, a, S)
+
+call init (N, a)
+call f4 (N, a, S)
+call check (N, a, S)
+
+call init (N, a)
+call f5 (N, a, S)
+call check (N, a, S)
+
+call init (N, a)
+call f6 (N, a, S)
+call check (N, a, S)
+
+call init (N, a)
+call f7 (N, a, S)
+call check (N, a, S)
+
+call init (N, a)
+call f8 (N, a, S)
+call check (N, a, S)
+
+call init (N, a)
+call f9 (N, a, S)
+call check (N, a, S)
+
+contains
+
+subroutine init (n, a)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  integer :: i
+  do i = 1, n
+    a(i) = i
+  end do
+end subroutine
+
+subroutine check (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+  do i = 1, n
+    if (a(i) /= i * s) error stop
+  end do
+end subroutine
+
+! Check various combinations for enforcing correct ordering of 
+! construct matches.
+subroutine f1 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target teams
+!$omp parallel  
+!$omp metadirective &
+!$omp &  when (construct={target} &
+!$omp &	: do) &
+!$omp &  default (error at(execution) message("f1 match failed"))
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end parallel  
+!$omp end target teams
+end subroutine
+
+subroutine f2 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target teams
+!$omp parallel
+!$omp metadirective &
+!$omp &  when (construct={teams, parallel} &
+!$omp &	: do) &
+!$omp &  default (error at(execution) message("f2 match failed"))
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end parallel  
+!$omp end target teams
+end subroutine
+
+subroutine f3 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target teams
+!$omp parallel
+!$omp metadirective &
+!$omp &  when (construct={target, teams, parallel} &
+!$omp &	: do) &
+!$omp &  default (error at(execution) message("f3 match failed"))
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end parallel  
+!$omp end target teams
+end subroutine
+
+subroutine f4 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target teams
+!$omp parallel
+!$omp metadirective &
+!$omp &  when (construct={target, parallel} &
+!$omp &	: do) &
+!$omp &  default (error at(execution) message("f4 match failed"))
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end parallel  
+!$omp end target teams
+end subroutine
+
+subroutine f5 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target teams
+!$omp parallel
+!$omp metadirective &
+!$omp &  when (construct={target, teams} &
+!$omp &	: do) &
+!$omp &  default (error at(execution) message("f5 match failed"))
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end parallel  
+!$omp end target teams
+end subroutine
+
+! Next batch is for things where the construct doesn't match the context.
+subroutine f6 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target
+!$omp teams
+!$omp metadirective &
+!$omp &  when (construct={parallel} &
+!$omp &	: error at(execution) message("f6 match failed")) &
+!$omp &  default (parallel do)
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end teams
+!$omp end target
+end subroutine
+
+subroutine f7 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target
+!$omp teams
+!$omp metadirective &
+!$omp &  when (construct={target, parallel} &
+!$omp &	: error at(execution) message("f7 match failed")) &
+!$omp &  default (parallel do)
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end teams
+!$omp end target
+end subroutine
+
+subroutine f8 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target
+!$omp teams
+!$omp metadirective &
+!$omp &  when (construct={parallel, target} &
+!$omp &	: error at(execution) message("f8 match failed")) &
+!$omp &  default (parallel do)
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end teams
+!$omp end target
+end subroutine
+
+! Next test choosing the best alternative when there are multiple
+! matches.
+subroutine f9 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(n)
+  double precision :: s
+  integer :: i
+!$omp target teams
+!$omp parallel
+!$omp metadirective &
+!$omp &  when (construct={teams, parallel} &
+!$omp &	: error at(execution) message("f9 match incorrect 1")) &
+!$omp &  when (construct={target, teams, parallel} &
+!$omp &	: do) &
+!$omp &  when (construct={target, teams} &
+!$omp &	: error at(execution) message("f9 match incorrect 2")) &
+!$omp &  default (error at(execution) message("f9 match failed"))
+  do i = 1, n
+    a(i) = a(i) * s
+  end do
+!$omp end parallel  
+!$omp end target teams
+end subroutine
+
+end program
+
+! Note there are no tests for the matching the extended simd clause
+! syntax, which is only useful for "declare variant".
+
+
+! After parsing, there should be a runtime error call for each of the
+! failure cases, but they should all be optimized away during OMP 
+! lowering.
+! { dg-final { scan-tree-dump-times "__builtin_GOMP_error" 11 "original" } }
+! { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-no-score.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-no-score.f90
new file mode 100644
index 0000000..968ce60
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-no-score.f90
@@ -0,0 +1,122 @@
+! { dg-do compile { target x86_64-*-* } }
+! { dg-additional-options "-foffload=disable" }
+
+! This test is expected to fail with compile-time errors:
+! "A trait-score cannot be specified in traits from the construct,
+!  device or target_device trait-selector-sets."
+
+
+subroutine f1 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(*)
+  double precision :: s
+  integer :: i
+!$omp metadirective &
+!$omp&  when (device={kind (score(5) : host)} &
+!$omp&	: parallel do)
+  ! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 }
+  do i = 1, n
+    a(i) = a(i) * s;
+  end do
+end subroutine
+
+subroutine f2 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(*)
+  double precision :: s
+  integer :: i
+!$omp metadirective &
+!$omp&  when (device={kind (host), arch (score(6) : x86_64), isa (avx512f)} &
+!$omp&	: parallel do)
+  ! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 }
+  do i = 1, n
+    a(i) = a(i) * s;
+  end do
+end subroutine
+
+subroutine f3 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(*)
+  double precision :: s
+  integer :: i
+!$omp metadirective &
+!$omp&  when (device={kind (host), arch (score(6) : x86_64), &
+!$omp&		  isa (score(7): avx512f)} &
+!$omp&	: parallel do)
+  ! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-3 }
+  do i = 1, n
+    a(i) = a(i) * s;
+  end do
+end subroutine
+
+subroutine f4 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(*)
+  double precision :: s
+  integer :: i
+  integer, parameter :: omp_initial_device = -1
+!$omp metadirective &
+!$omp&  when (target_device={device_num (score(42) : omp_initial_device), &
+!$omp&			 kind (host)} &
+!$omp&	: parallel do)
+  ! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 }
+  do i = 1, n
+    a(i) = a(i) * s;
+  end do
+end subroutine
+
+subroutine f5 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(*)
+  double precision :: s
+  integer :: i
+  integer, parameter :: omp_initial_device = -1
+!$omp metadirective &
+!$omp&  when (target_device={device_num(omp_initial_device), &
+!$omp&			 kind (score(5) : host)} &
+!$omp&	: parallel do)
+  ! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 }
+  do i = 1, n
+    a(i) = a(i) * s;
+  end do
+end subroutine
+
+subroutine f6 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(*)
+  double precision :: s
+  integer :: i
+  integer, parameter :: omp_initial_device = -1
+!$omp metadirective &
+!$omp&  when (target_device={device_num(omp_initial_device), kind (host), &
+!$omp&			 arch (score(6) : x86_64), isa (avx512f)} &
+!$omp&	: parallel do)
+  ! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 }
+  do i = 1, n
+    a(i) = a(i) * s;
+  end do
+end subroutine
+
+subroutine f7 (n, a, s)
+  implicit none
+  integer :: n
+  double precision :: a(*)
+  double precision :: s
+  integer :: i
+  integer, parameter :: omp_initial_device = -1
+!$omp metadirective &
+!$omp&  when (target_device={device_num(omp_initial_device), kind (host), &
+!$omp&			 arch (score(6) : x86_64), &
+!$omp&			 isa (score(7): avx512f)} &
+!$omp&	: parallel do)
+  ! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 }
+  do i = 1, n
+    a(i) = a(i) * s;
+  end do
+end subroutine
diff --git a/gcc/testsuite/gfortran.dg/gomp/pure-1.f90 b/gcc/testsuite/gfortran.dg/gomp/pure-1.f90
index cdbebe2..d7bc6cc 100644
--- a/gcc/testsuite/gfortran.dg/gomp/pure-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/pure-1.f90
@@ -110,3 +110,32 @@ pure integer function func_tile(n)
   end do
   func_tile = r
 end
+
+pure logical function func_metadirective()
+  implicit none
+  !$omp metadirective
+  func_metadirective = .false.
+end
+
+! not 'parallel' not pure -> invalid in 5.2; + in general invalid in 5.1
+pure logical function func_metadirective_2 ()
+  implicit none
+  integer :: i, n
+  n = 0
+  !$omp metadirective when (device={arch("nvptx")} : parallel do)     ! { dg-error "OpenMP directive at .1. is not pure and thus may not appear in a PURE procedure" }
+  do i = 1, 5
+    n = n + i
+  end do
+end
+
+! unroll is supposed to be pure, so this case is OK
+pure logical function func_metadirective_3()
+  implicit none
+  integer :: i, n
+
+  n = 0
+  !$omp metadirective when(device={arch("nvptx")} : unroll full)
+  do i = 1, 5
+    n = n + i
+  end do
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pure-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pure-2.f90
index 35503c6..f602218 100644
--- a/gcc/testsuite/gfortran.dg/gomp/pure-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/pure-2.f90
@@ -26,14 +26,6 @@ logical function func_interchange(n)
   end do
 end
 
-
-!pure logical function func_metadirective()
-logical function func_metadirective()
-  implicit none
-  !$omp metadirective  ! { dg-error "Unclassifiable OpenMP directive" }
-  func_metadirective = .false.
-end
-
 !pure logical function func_reverse(n)
 logical function func_reverse(n)
   implicit none
diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90
new file mode 100644
index 0000000..7b3e09f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90
@@ -0,0 +1,61 @@
+! { dg-do run }
+
+program test
+  implicit none
+
+  integer, parameter :: N = 100
+  integer :: x(N), y(N), z(N)
+  integer :: i
+
+  do i = 1, N
+    x(i) = i;
+    y(i) = -i;
+  end do
+
+  call f (x, y, z)
+
+  do i = 1, N
+    if (z(i) .ne. x(i) * y(i)) stop 1
+  end do
+
+  ! -----
+  do i = 1, N
+    x(i) = i;
+    y(i) = -i;
+  end do
+
+  call g (x, y, z)
+
+  do i = 1, N
+    if (z(i) .ne. x(i) * y(i)) stop 1
+  end do
+
+contains
+  subroutine f (x, y, z)
+    integer :: x(N), y(N), z(N)
+
+    !$omp target map (to: x, y) map(from: z)
+      block
+      !$omp metadirective &
+		!$omp& when(device={arch("nvptx")}: teams loop) &
+		!$omp& default(parallel loop)
+	do i = 1, N
+	  z(i) = x(i) * y(i)
+	enddo
+      end block
+  end subroutine
+  subroutine g (x, y, z)
+    integer :: x(N), y(N), z(N)
+
+    !$omp target map (to: x, y) map(from: z)
+    block
+      !$omp metadirective &
+		!$omp& when(device={arch("nvptx")}: teams loop) &
+		!$omp& default(parallel loop)
+	do i = 1, N
+	  z(i) = x(i) * y(i)
+	enddo
+    end block
+    !$omp end target
+  end subroutine
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90
new file mode 100644
index 0000000..8018b1a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90
@@ -0,0 +1,38 @@
+! { dg-do run }
+
+program test
+  implicit none
+  integer, parameter :: N = 100
+  real, parameter :: PI_CONST = 2.0*acos(0.0)
+  real, parameter :: E_CONST = exp(1.0)
+  real, parameter :: EPSILON = 0.001
+  integer :: i
+  real :: a(N)
+
+  !$omp target map(from: a)
+    call f (a, PI_CONST)
+  !$omp end target
+
+  do i = 1, N
+    if (abs (a(i) - (PI_CONST * i)) .gt. EPSILON) stop 1
+  end do
+
+  call f (a, E_CONST)
+
+  do i = 1, N
+    if (abs (a(i) - (E_CONST * i)) .gt. EPSILON) stop 2
+  end do
+contains
+  subroutine f (a, x)
+    integer :: i
+    real :: a(N), x
+    !$omp declare target
+
+    !$omp metadirective &
+    !$omp&  when (construct={target}: distribute parallel do ) &
+    !$omp&  default(parallel do simd)
+      do i = 1, N
+	a(i) = x * i
+      end do
+  end subroutine
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-3.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-3.f90
new file mode 100644
index 0000000..693c40b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/metadirective-3.f90
@@ -0,0 +1,29 @@
+! { dg-do run }
+
+program test
+  implicit none
+
+  integer, parameter :: N = 100
+  integer :: a(N)
+  integer :: res
+
+  if (f (a, .false.)) stop 1
+  if (.not. f (a, .true.)) stop 2
+contains
+  logical function f (a, flag)
+    integer :: a(N)
+    logical :: flag
+    logical :: res = .false.
+    integer :: i
+    f = .false.
+    !$omp metadirective &
+    !$omp&  when (user={condition(.not. flag)}: &
+    !$omp&	 target teams distribute parallel do &
+    !$omp&		map(from: a(1:N)) private(res)) &
+    !$omp&  default(parallel do)
+      do i = 1, N
+	a(i) = i
+	f = .true.
+     end do
+  end function
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-4.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-4.f90
new file mode 100644
index 0000000..04fdf61
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/metadirective-4.f90
@@ -0,0 +1,46 @@
+! { dg-do run }
+
+program test
+  use omp_lib
+
+  implicit none
+  integer, parameter :: N = 100
+  integer :: a(N)
+  logical :: is_parallel, is_static
+
+  ! is_static is always set if run_parallel is false.
+  call f (a, .false., .false., is_parallel, is_static)
+  if (is_parallel .or. .not. is_static) stop 1
+
+  call f (a, .false., .true., is_parallel, is_static)
+  if (is_parallel .or. .not. is_static) stop 2
+
+  call f (a, .true., .false., is_parallel, is_static)
+  if (.not. is_parallel .or. is_static) stop 3
+
+  call f (a, .true., .true., is_parallel, is_static)
+  if (.not. is_parallel .or. .not. is_static) stop 4
+contains
+  subroutine f (a, run_parallel, run_static, is_parallel, is_static)
+    integer :: a(N)
+    logical, intent(in) :: run_parallel, run_static
+    logical, intent(out) :: is_parallel, is_static
+    integer :: i
+
+    is_parallel = .false.
+    is_static = .false.
+
+    !$omp begin metadirective when(user={condition(run_parallel)}: parallel)
+      if (omp_in_parallel ()) is_parallel = .true.
+
+      !$omp metadirective &
+      !$omp&  when(construct={parallel}, user={condition(.not. run_static)}: &
+      !$omp&       do schedule(guided) private(is_static)) &
+      !$omp&  when(construct={parallel}: do schedule(static))
+	do i = 1, N
+	  a(i) = i
+	  is_static = .true.
+	end do
+    !$omp end metadirective
+  end subroutine
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-5.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-5.f90
new file mode 100644
index 0000000..3992286
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/metadirective-5.f90
@@ -0,0 +1,44 @@
+! { dg-do run }
+
+program main
+  use omp_lib
+
+  implicit none
+
+  integer, parameter :: N = 100
+  integer :: a(N)
+  integer :: on_device_count = 0
+  integer :: i
+
+  do i = 1, N
+    a(i) = i
+  end do
+
+  do i = 0, omp_get_num_devices ()
+    on_device_count = on_device_count + f (a, i)
+  end do
+
+  if (on_device_count .ne. omp_get_num_devices ()) stop 1
+
+  do i = 1, N
+    if (a(i) .ne. 2 * i) stop 2;
+  end do
+contains
+  integer function f (a, num)
+    integer, intent(inout) :: a(N)
+    integer, intent(in) :: num
+    integer :: on_device
+    integer :: i
+
+    on_device = 0
+    !$omp metadirective &
+    !$omp&  when (target_device={device_num(num), kind("gpu")}: &
+    !$omp&    target parallel do map(to: a(1:N)), map(from: on_device)) &
+    !$omp&  default (parallel do private(on_device))
+      do i = 1, N
+        a(i) = a(i) + i
+        on_device = 1
+      end do
+    f = on_device;
+  end function
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-6.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-6.f90
new file mode 100644
index 0000000..436fdba
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/metadirective-6.f90
@@ -0,0 +1,58 @@
+! { dg-do compile }
+
+program test
+  implicit none
+
+  integer, parameter :: N = 100
+  integer :: x(N), y(N), z(N)
+  integer :: i
+
+contains
+  subroutine f (x, y, z)
+    integer :: x(N), y(N), z(N)
+
+    !$omp target map (to: x, y) map(from: z)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+      block
+      !$omp metadirective &
+		!$omp& when(device={arch("nvptx")}: teams loop) &
+		!$omp& default(parallel loop)  ! { dg-error "\\(1\\)" }
+ ! FIXME: The line above should be the same error as above but some fails here with -fno-diagnostics-show-caret
+ ! Seems as if some gcc/testsuite/ fix is missing for libgomp/testsuite
+	do i = 1, N
+	  z(i) = x(i) * y(i)
+	enddo
+       z(N) = z(N) + 1  ! <<< invalid
+      end block
+  end subroutine
+
+  subroutine f2 (x, y, z)
+    integer :: x(N), y(N), z(N)
+
+    !$omp target map (to: x, y) map(from: z)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+      block
+      integer :: i ! << invalid
+      !$omp metadirective &
+		!$omp& when(device={arch("nvptx")}: teams loop) &
+		!$omp& default(parallel loop)
+	do i = 1, N
+	  z(i) = x(i) * y(i)
+	enddo
+      end block
+  end subroutine
+  subroutine g (x, y, z)
+    integer :: x(N), y(N), z(N)
+
+    !$omp target map (to: x, y) map(from: z)  ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
+    block
+      !$omp metadirective &   ! <<<< invalid
+		!$omp& when(device={arch("nvptx")}: flush) &
+		!$omp& default(nothing)
+       !$omp teams loop
+	do i = 1, N
+	  z(i) = x(i) * y(i)
+	enddo
+    end block
+    !$omp end target
+  end subroutine
+
+end program