diff --git a/gcc/fortran/decl.cc b/gcc/fortran/decl.cc
index 503ecb8d9b5..0ab2e710e78 100644
--- a/gcc/fortran/decl.cc
+++ b/gcc/fortran/decl.cc
@@ -8404,6 +8404,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;
@@ -8411,6 +8412,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;
     }
@@ -8556,6 +8579,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 1563b810b98..204c6b60a1f 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -2162,6 +2162,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;
@@ -2359,6 +2360,25 @@ show_omp_node (int level, gfc_code *c)
 	  d = d->block;
 	}
     }
+  else if (c->op == EXEC_OMP_METADIRECTIVE)
+    {
+      gfc_omp_metadirective_variant *variant
+	= c->ext.omp_metadirective_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)
@@ -3493,6 +3513,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 82f388c05f8..7adc98fccf1 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -317,6 +317,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,
@@ -1583,7 +1584,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;
   ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3;
   ENUM_BITFIELD (gfc_omp_device_type) device_type:2;
@@ -1703,6 +1704,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_metadirective_variant
+{
+  struct gfc_omp_metadirective_variant *next;
+  locus where; /* Where the metadirective clause occurred.  */
+
+  gfc_omp_set_selector *selectors;
+  enum gfc_statement stmt;
+  struct gfc_code *code;
+
+} gfc_omp_metadirective_variant;
+#define gfc_get_omp_metadirective_variant() XCNEW (gfc_omp_metadirective_variant)
 
 typedef struct gfc_omp_udr
 {
@@ -1751,6 +1763,7 @@ typedef struct gfc_st_label
   locus where;
 
   gfc_namespace *ns;
+  int omp_region;
 }
 gfc_st_label;
 
@@ -3009,6 +3022,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_METADIRECTIVE,
   EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS
 };
 
@@ -3067,6 +3081,7 @@ typedef struct gfc_code
     gfc_omp_clauses *omp_clauses;
     const char *omp_name;
     gfc_omp_namelist *omp_namelist;
+    gfc_omp_metadirective_variant *omp_metadirective_variants;
     bool omp_bool;
   }
   ext;		/* Points to additional structures required by statement */
@@ -3661,6 +3676,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_metadirective_variants (gfc_omp_metadirective_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 *);
@@ -3942,6 +3958,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);
@@ -4038,4 +4056,8 @@ bool gfc_is_reallocatable_lhs (gfc_expr *);
 void finish_oacc_declare (gfc_namespace *, gfc_symbol *, bool);
 void gfc_adjust_builtins (void);
 
+/* trans-openmp.c */
+
+bool gfc_skip_omp_metadirective_variant (gfc_omp_metadirective_variant *);
+
 #endif /* GCC_GFORTRAN_H  */
diff --git a/gcc/fortran/io.cc b/gcc/fortran/io.cc
index 6fd69f7c9a8..9dfb4e1ef25 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 b09921357fd..b0f5528ee72 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);
@@ -178,6 +179,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 0af80d54fad..b602d9aa34d 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -113,7 +113,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)
@@ -124,17 +125,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;
@@ -339,6 +348,19 @@ gfc_free_omp_udr (gfc_omp_udr *omp_udr)
     }
 }
 
+/* Free variants of an !$omp metadirective construct.  */
+
+void
+gfc_free_omp_metadirective_variants (gfc_omp_metadirective_variant *variant)
+{
+  while (variant)
+    {
+      gfc_omp_metadirective_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)
@@ -1866,8 +1888,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 ();
@@ -3789,9 +3810,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");
@@ -4479,7 +4498,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;
@@ -5810,14 +5829,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)
 		  {
@@ -5874,7 +5896,7 @@ 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)
 {
   do
     {
@@ -5907,9 +5929,9 @@ 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)
 	return MATCH_ERROR;
@@ -6010,7 +6032,8 @@ gfc_match_omp_declare_variant (void)
 	  return MATCH_ERROR;
 	}
 
-      if (gfc_match_omp_context_selector_specification (odv) != MATCH_YES)
+      if (gfc_match_omp_context_selector_specification (&odv->set_selectors)
+	  != MATCH_YES)
 	return MATCH_ERROR;
 
       if (gfc_match (" )") != MATCH_YES)
@@ -6026,6 +6049,157 @@ gfc_match_omp_declare_variant (void)
 }
 
 
+static match
+match_omp_metadirective (bool begin_p)
+{
+  locus old_loc = gfc_current_locus;
+  gfc_omp_metadirective_variant *variants_head;
+  gfc_omp_metadirective_variant **next_variant = &variants_head;
+  bool default_seen = false;
+
+  /* Parse the context selectors.  */
+  for (;;)
+    {
+      bool default_p = false;
+      gfc_omp_set_selector *selectors = NULL;
+      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_eos () == MATCH_YES)
+	break;
+      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;
+	}
+
+      if (!default_p)
+	{
+	  if (gfc_match_omp_context_selector_specification (&selectors)
+	      != 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_metadirective_variant *omv = gfc_get_omp_metadirective_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;
+	}
+
+      if (!gfc_skip_omp_metadirective_variant (omv))
+	{
+	  *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_metadirective_variant *omv = gfc_get_omp_metadirective_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_metadirective_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)
 {
@@ -10919,6 +11093,19 @@ resolve_omp_do (gfc_code *code)
   restructure_intervening_code (&(code->block->next), code, count);
 }
 
+static void
+resolve_omp_metadirective (gfc_code *code, gfc_namespace *ns)
+{
+  gfc_omp_metadirective_variant *variant = code->ext.omp_metadirective_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)
@@ -11462,13 +11649,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_metadirective_variant *mc
+	    = c->ext.omp_metadirective_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)
@@ -11595,6 +11801,9 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns)
       code->ext.omp_clauses->if_present = false;
       resolve_omp_clauses (code, code->ext.omp_clauses, ns);
       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 d8b38cfb5ac..f23045636c9 100644
--- a/gcc/fortran/parse.cc
+++ b/gcc/fortran/parse.cc
@@ -48,6 +48,10 @@ gfc_state_data *gfc_state_stack;
 static bool last_was_use_stmt = false;
 bool in_exec_part;
 
+bool gfc_matching_omp_context_selector;
+bool gfc_in_metadirective_body;
+int gfc_omp_region_count;
+
 /* TODO: Re-order functions to kill these forward decls.  */
 static void check_statement_label (gfc_statement);
 static void undo_new_statement (void);
@@ -1041,6 +1045,8 @@ decode_omp_directive (void)
       break;
     case 'b':
       matcho ("barrier", gfc_match_omp_barrier, ST_OMP_BARRIER);
+      matcho ("begin metadirective", gfc_match_omp_begin_metadirective,
+	      ST_OMP_BEGIN_METADIRECTIVE);
       break;
     case 'c':
       matcho ("cancellation% point", gfc_match_omp_cancellation_point,
@@ -1085,6 +1091,8 @@ decode_omp_directive (void)
       matcho ("end master taskloop", gfc_match_omp_eos_error,
 	      ST_OMP_END_MASTER_TASKLOOP);
       matcho ("end master", gfc_match_omp_eos_error, ST_OMP_END_MASTER);
+      matcho ("end metadirective", gfc_match_omp_eos_error,
+	      ST_OMP_END_METADIRECTIVE);
       matchs ("end ordered", gfc_match_omp_eos_error, ST_OMP_END_ORDERED);
       matchs ("end parallel do simd", gfc_match_omp_eos_error,
 	      ST_OMP_END_PARALLEL_DO_SIMD);
@@ -1168,6 +1176,8 @@ decode_omp_directive (void)
       matcho ("master taskloop", gfc_match_omp_master_taskloop,
 	      ST_OMP_MASTER_TASKLOOP);
       matcho ("master", gfc_match_omp_master, ST_OMP_MASTER);
+      matcho ("metadirective", gfc_match_omp_metadirective,
+	      ST_OMP_METADIRECTIVE);
       break;
     case 'n':
       matcho ("nothing", gfc_match_omp_nothing, ST_NONE);
@@ -1296,6 +1306,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 ();
@@ -1413,6 +1427,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)
 {
@@ -1935,6 +1955,43 @@ 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
+
 /* Block end statements.  Errors associated with interchanging these
    are detected in gfc_match_end().  */
 
@@ -2572,6 +2629,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;
@@ -2671,6 +2731,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;
@@ -2815,6 +2878,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";
@@ -3075,6 +3141,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 ();
@@ -5386,6 +5454,140 @@ 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;
+	default:
+	  break;
+	}
+    }
+
+  if (omp_structured_p)
+    {
+      switch (omp_st)
+	{
+	case ST_OMP_ASSUME:
+	  return ST_OMP_END_ASSUME;
+	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.  */
 
@@ -5436,94 +5638,16 @@ parse_omp_do (gfc_statement omp_st)
   pop_state ();
 
   st = next_statement ();
-  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;
-    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)
@@ -5835,80 +5959,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;
@@ -5954,11 +6013,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:
@@ -6131,6 +6192,88 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
   return st;
 }
 
+static gfc_statement
+parse_omp_metadirective_body (gfc_statement omp_st)
+{
+  gfc_omp_metadirective_variant *variant
+    = new_st.ext.omp_metadirective_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_metadirective_body;
+      gfc_in_metadirective_body = true;
+
+      gfc_omp_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);
+	  /* TODO: Does st == ST_IMPLIED_ENDDO need special handling?  */
+	  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 an OMP METADIRECTIVE block at %C",
+	    gfc_ascii_statement (st));
+	  reject_statement ();
+	  st = next_statement ();
+	}
+    finish:
+
+      gfc_in_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
@@ -6141,6 +6284,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)
@@ -6148,6 +6292,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)
@@ -6257,68 +6407,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_omp_do:
 	  st = parse_omp_do (st);
 	  if (st == ST_IMPLIED_ENDDO)
 	    return st;
@@ -6332,6 +6427,19 @@ parse_executable (gfc_statement st)
 	  st = parse_omp_oacc_atomic (true);
 	  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)
+	    {
+	      st = next_statement ();
+	      return st;
+	    }
+	  /* FALLTHRU */
+
 	default:
 	  return st;
 	}
@@ -7097,6 +7205,10 @@ gfc_parse_file (void)
 
   gfc_statement_label = NULL;
 
+  gfc_omp_region_count = 0;
+  gfc_in_metadirective_body = false;
+  gfc_matching_omp_context_selector = false;
+
   if (setjmp (eof_buf))
     return false;	/* Come here on unexpected EOF */
 
@@ -7409,3 +7521,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 ce19d4deb07..2c5b3adec3b 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_metadirective_body;
+extern int gfc_omp_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 6148ed94fe1..5858b3a88dc 100644
--- a/gcc/fortran/resolve.cc
+++ b/gcc/fortran/resolve.cc
@@ -12261,6 +12261,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_metadirective_variant *variant
+	       = code->ext.omp_metadirective_variants;
+	     variant; variant = variant->next)
+	  gfc_resolve_code (variant->code, ns);
       else if (code->block)
 	{
 	  omp_workshare_save = -1;
@@ -12793,6 +12798,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 6a605ad91d4..9d6cc974083 100644
--- a/gcc/fortran/st.cc
+++ b/gcc/fortran/st.cc
@@ -299,6 +299,10 @@ gfc_free_statement (gfc_code *p)
     case EXEC_OMP_TASKYIELD:
       break;
 
+    case EXEC_OMP_METADIRECTIVE:
+      gfc_free_omp_metadirective_variants (p->ext.omp_metadirective_variants);
+      break;
+
     default:
       gfc_internal_error ("gfc_free_statement(): Bad statement");
     }
diff --git a/gcc/fortran/symbol.cc b/gcc/fortran/symbol.cc
index fddf68f8398..cd1de989eab 100644
--- a/gcc/fortran/symbol.cc
+++ b/gcc/fortran/symbol.cc
@@ -2630,10 +2630,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;
 }
 
 
@@ -2683,6 +2686,7 @@ gfc_get_st_label (int labelno)
 {
   gfc_st_label *lp;
   gfc_namespace *ns;
+  int omp_region = gfc_in_metadirective_body ? gfc_omp_region_count : 0;
 
   if (gfc_current_state () == COMP_DERIVED)
     ns = gfc_current_block ()->f2k_derived;
@@ -2699,10 +2703,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;
@@ -2714,6 +2724,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 de162f6cc75..a17f6582d81 100644
--- a/gcc/fortran/trans-decl.cc
+++ b/gcc/fortran/trans-decl.cc
@@ -327,7 +327,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 439ae5f9f55..2b72862d62b 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -8228,6 +8228,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:
@@ -8319,6 +8321,104 @@ gfc_trans_omp_declare_simd (gfc_namespace *ns)
     }
 }
 
+/* Translate the context selector list GFC_SELECTORS, using WHERE as the
+   locus for error messages.  CONV_EXPR_P is true for normal translation
+   and false for early conversion during parsing.  */
+
+static tree
+gfc_trans_omp_set_selector (gfc_omp_set_selector *gfc_selectors, locus where,
+			    bool conv_expr_p)
+{
+  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_EXPR:
+		  {
+		    tree expr = NULL_TREE;
+		    if (conv_expr_p)
+		      {
+			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)
 {
@@ -8394,89 +8494,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_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, true);
       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)
@@ -8577,3 +8596,57 @@ 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_metadirective_variant *variant = code->ext.omp_metadirective_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, true);
+      ctx = omp_check_context_selector (gfc_get_location (&variant->where),
+					ctx, true);
+      if (ctx == error_mark_node)
+	return error_mark_node;
+
+      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;
+}
+
+/* Hook to check during parsing whether a metadirective is both resolvable
+   and non-matching; if so, it can be eliminated immediately.  */
+bool
+gfc_skip_omp_metadirective_variant (gfc_omp_metadirective_variant *variant)
+{
+  tree selector = gfc_trans_omp_set_selector (variant->selectors,
+					      variant->where, false);
+  return omp_context_selector_matches (selector, true, true) == 0;
+}
diff --git a/gcc/fortran/trans-stmt.h b/gcc/fortran/trans-stmt.h
index 0f0f99931ca..d19e161cf11 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 1a03a4a91d4..7e3089dfc8b 100644
--- a/gcc/fortran/trans.cc
+++ b/gcc/fortran/trans.cc
@@ -2611,6 +2611,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 00000000000..c5b3946341d
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90
@@ -0,0 +1,55 @@
+! { 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) & ! { dg-error "expected 'when', 'otherwise', or 'default' at .1." }
+  !$omp&	where (device={arch("nvptx")}: parallel loop)
+    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 00000000000..5dad5d29eb6
--- /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 00000000000..e7de70e6259
--- /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-2.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90
new file mode 100644
index 00000000000..cdd5e85068e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90
@@ -0,0 +1,62 @@
+! { 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 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 00000000000..eca389a7842
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90
@@ -0,0 +1,34 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-gimple" }
+! { dg-additional-options "-fdump-tree-optimized" }
+
+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
+
+! The metadirective should be resolved after Gimplification.
+
+! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } }
+! { dg-final { scan-tree-dump-times "when \\(device =.*arch.*nvptx.*\\):" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } }
+! { dg-final { scan-tree-dump-times "default:" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } }
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 00000000000..2fdc7f3c515
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90
@@ -0,0 +1,39 @@
+! { 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
+
+  ! TODO: This does not execute a version of f with the default clause
+  ! active as might be expected.
+  call f (a, 2.71828) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" }
+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 "default:" 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 00000000000..03970393eb4
--- /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 00000000000..9b6c371296f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90
@@ -0,0 +1,31 @@
+! { 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
+
+! The outer metadirective should be resolved at parse time, but is
+! currently resolved during Gimplification.
+
+! The inner metadirective should be resolved during Gimplificiation.
+
+! { 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 00000000000..870ea192fbc
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90
@@ -0,0 +1,36 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+  integer, parameter :: N = 256
+contains
+  subroutine f (a, num)
+    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
+end program
+
+! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"gpu\"\\\[0\\\], &\"amdgcn\"\\\[0\\\], &\"gfx906\"\\\[0\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } }
+! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"cpu\"\\\[0\\\], &\"x86_64\"\\\[0\\\], 0B\\)" "gimple" } }
+! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(-1, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } }
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 00000000000..3f46801e044
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90
@@ -0,0 +1,22 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+program test
+  integer :: i
+  integer, parameter :: N = 100
+  integer :: sum = 0
+  
+  ! The compiler should never consider a situation where both metadirectives
+  ! match.  If it does, then the nested metadirective would be an error
+  ! as it is not a loop-nest as per the OpenMP specification.
+
+  !$omp metadirective when (implementation={vendor("ibm")}: &
+  !$omp&  target teams distribute)
+    !$omp metadirective when (implementation={vendor("gnu")}: parallel do)
+      do i = 1, N
+	sum = sum + i
+      end do
+end program
+
+! { dg-final { scan-tree-dump-not "when \\(implementation = .*vendor.*ibm.*\\):" "original" } }
+! { dg-final { scan-tree-dump-times "when \\(implementation = .*vendor.*gnu.*\\):" 1 "original" } }
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 00000000000..e6ab3fc0a65
--- /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 an 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 00000000000..ec1f0ee3d9d
--- /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 00000000000..968ce609b10
--- /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 598e455d2e9..15681ac604a 100644
--- a/gcc/testsuite/gfortran.dg/gomp/pure-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/pure-1.f90
@@ -86,3 +86,10 @@ pure integer function func_simd(n)
   end do
   func_simd = r
 end
+
+!pure logical function func_metadirective()
+logical function func_metadirective()
+  implicit none
+  !$omp metadirective
+  func_metadirective = .false.
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/pure-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pure-2.f90
index 1e3cf8c9416..25d6069b236 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 00000000000..7b3e09f7c2a
--- /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 00000000000..d83474cf2db
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90
@@ -0,0 +1,40 @@
+! { dg-do run }
+
+program test
+  implicit none
+  integer, parameter :: N = 100
+  real, parameter :: PI_CONST = 3.14159
+  real, parameter :: E_CONST = 2.71828
+  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
+
+  ! TODO: This does not execute a version of f with the default clause
+  ! active as might be expected.
+  call f (a, E_CONST) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" }
+
+  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 00000000000..693c40bca5a
--- /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 00000000000..04fdf61489c
--- /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 00000000000..3992286dc08
--- /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 00000000000..436fdbade2f
--- /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
