public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [WIP, OpenMP] OpenMP metadirectives support
@ 2021-07-09 11:16 Kwok Cheung Yeung
  2021-07-26 11:38 ` Kwok Cheung Yeung
                   ` (2 more replies)
  0 siblings, 3 replies; 29+ messages in thread
From: Kwok Cheung Yeung @ 2021-07-09 11:16 UTC (permalink / raw)
  To: GCC Patches, Jakub Jelinek

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

Hello

This is a WIP implementation of metadirectives as defined in the OpenMP 5.0 
spec. I intend to add support for metadirectives as specified in OpenMP 5.1 
later (where the directive can be selected dynamically at runtime), but am 
concentrating on the static part for now. Parsing has only been implemented in 
the C frontend so far. I am especially interested in feedback regarding certain 
aspects of the implementation before I become too committed to the current design.

1) When parsing each directive variant, a vector of tokens is constructed and 
populated with the tokens for a regular equivalent pragma, along with the tokens 
for its clauses and the body. The parser routine for that pragma type is then 
called with these tokens, and the entire resulting parse tree is stored as a 
sub-tree of the metadirective tree structure.

This results in the body being parsed and stored once for each directive 
variant. I believe this is necessary because the body is parsed differently if 
there is a 'for' in the directive (using c_parser_omp_for_loop) compared to if 
there is not, plus clauses in the directive (e.g. tile, collapse) can change how 
the for loop is parsed.

As an optimisation, identical body trees could be merged together, but that can 
come later.

2) Selectors in the device set (i.e. kind, isa, arch) resolve differently 
depending on whether the program is running on a target or on the host. Since we 
don't keep multiple versions of a function for each target on the host compiler, 
resolving metadirectives with these selectors needs to be delayed until after 
LTO streaming, at which point the host or offload compiler can make the 
appropriate decision.

One negative of this is that the metadirective Gimple representation lasts 
beyond the OMP expand stage, when generally we would expect all OMP directives 
to have been expanded to something else.

3) In the OpenMP examples (version 5.0.1), section 9.7, the example 
metadirective.3.c does not work as expected.

#pragma omp declare target
void exp_pi_diff(double *d, double my_pi){
    #pragma omp metadirective \
                when( construct={target}: distribute parallel for ) \
                default( parallel for simd)
...
int main()
{
    ...
    #pragma omp target teams map(tofrom: d[0:N])
    exp_pi_diff(d,my_pi);
    ...
    exp_pi_diff(d,my_pi);

In the first call to exp_pi_diff in an '#pragma omp target' construct, the 
metadirective is expected to expand to 'distribute parallel for', but in the 
second (without the '#pragma omp target'), it should expand to 'parallel for simd'.

During OMP expansion of the 'omp target', it creates a child function that calls 
exp_pi_diff:

__attribute__((omp target entrypoint))
void main._omp_fn.0 (const struct .omp_data_t.12 & restrict .omp_data_i)
{
   ...
   <bb 4> :
   __builtin_GOMP_teams (0, 0);
   exp_pi_diff (d.13, my_pi);

This is not a problem on the offload compiler (since by definition its copy of 
exp_pi_diff must be in a 'target'), but if the host device is used, the same 
version of exp_pi_diff is called in both target and non-target contexts.

What would be the best way to solve this? Offhand, I can think of two solutions:

(a) Recursively go through all functions that can be reached via a target region 
and create clones for each, redirecting all function calls in the clones to the 
new cloned versions. Resolve the metadirectives in the clones and originals 
separately.

(b) Make the construct selector a dynamic selector when OpenMP 5.1 metadirective 
support is implemented. Keep track of the current construct list every time an 
OpenMP construct is entered or exited, and make the decision at runtime.


Thanks

Kwok

[-- Attachment #2: omp_metadirectives_wip.patch --]
[-- Type: text/plain, Size: 57697 bytes --]

diff --git a/gcc/Makefile.in b/gcc/Makefile.in
index 1164554e6d6..28e29fab93d 100644
--- a/gcc/Makefile.in
+++ b/gcc/Makefile.in
@@ -1505,6 +1505,7 @@ OBJS = \
 	omp-general.o \
 	omp-low.o \
 	omp-oacc-kernels-decompose.o \
+        omp-expand-metadirective.o \
 	omp-simd-clone.o \
 	opt-problem.o \
 	optabs.o \
diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index 4f8e8e0128c..01dc1e6d9c0 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1312,12 +1312,14 @@ static const struct omp_pragma_def omp_pragmas[] = {
   { "allocate", PRAGMA_OMP_ALLOCATE },
   { "atomic", PRAGMA_OMP_ATOMIC },
   { "barrier", PRAGMA_OMP_BARRIER },
+  { "begin", PRAGMA_OMP_BEGIN },
   { "cancel", PRAGMA_OMP_CANCEL },
   { "cancellation", PRAGMA_OMP_CANCELLATION_POINT },
   { "critical", PRAGMA_OMP_CRITICAL },
   { "depobj", PRAGMA_OMP_DEPOBJ },
-  { "end", PRAGMA_OMP_END_DECLARE_TARGET },
+  { "end", PRAGMA_OMP_END },
   { "flush", PRAGMA_OMP_FLUSH },
+  { "metadirective", PRAGMA_OMP_METADIRECTIVE },
   { "requires", PRAGMA_OMP_REQUIRES },
   { "section", PRAGMA_OMP_SECTION },
   { "sections", PRAGMA_OMP_SECTIONS },
@@ -1387,6 +1389,41 @@ c_pp_lookup_pragma (unsigned int id, const char **space, const char **name)
   gcc_unreachable ();
 }
 
+static int
+c_pp_lookup_pragma_by_name_1 (const void *name, const void *elem)
+{
+  const struct omp_pragma_def *pragma_def
+    = (const struct omp_pragma_def *) elem;
+
+  return strcmp ((const char *) name, pragma_def->name);
+}
+
+enum pragma_kind
+c_pp_lookup_pragma_by_name (const char *name)
+{
+  const int n_omp_pragmas = sizeof (omp_pragmas) / sizeof (*omp_pragmas);
+  const int n_omp_pragmas_simd = sizeof (omp_pragmas_simd)
+				 / sizeof (*omp_pragmas_simd);
+
+  void *result = bsearch (name, omp_pragmas, n_omp_pragmas,
+			  sizeof (*omp_pragmas),
+			  c_pp_lookup_pragma_by_name_1);
+  if (!result)
+    result = bsearch (name, omp_pragmas_simd, n_omp_pragmas_simd,
+		      sizeof (*omp_pragmas_simd),
+		      c_pp_lookup_pragma_by_name_1);
+
+  if (result)
+    {
+      const struct omp_pragma_def *def
+	= (const struct omp_pragma_def *) result;
+
+      return (enum pragma_kind) def->id;
+    }
+
+  return PRAGMA_NONE;
+}
+
 /* Front-end wrappers for pragma registration to avoid dragging
    cpplib.h in almost everywhere.  */
 
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 6c34ffa5be4..6d4698d41ba 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -45,17 +45,19 @@ enum pragma_kind {
   PRAGMA_OMP_ALLOCATE,
   PRAGMA_OMP_ATOMIC,
   PRAGMA_OMP_BARRIER,
+  PRAGMA_OMP_BEGIN,
   PRAGMA_OMP_CANCEL,
   PRAGMA_OMP_CANCELLATION_POINT,
   PRAGMA_OMP_CRITICAL,
   PRAGMA_OMP_DECLARE,
   PRAGMA_OMP_DEPOBJ,
   PRAGMA_OMP_DISTRIBUTE,
-  PRAGMA_OMP_END_DECLARE_TARGET,
+  PRAGMA_OMP_END,
   PRAGMA_OMP_FLUSH,
   PRAGMA_OMP_FOR,
   PRAGMA_OMP_LOOP,
   PRAGMA_OMP_MASTER,
+  PRAGMA_OMP_METADIRECTIVE,
   PRAGMA_OMP_ORDERED,
   PRAGMA_OMP_PARALLEL,
   PRAGMA_OMP_REQUIRES,
@@ -252,6 +254,7 @@ extern enum cpp_ttype c_lex_with_flags (tree *, location_t *, unsigned char *,
 					int);
 
 extern void c_pp_lookup_pragma (unsigned int, const char **, const char **);
+extern enum pragma_kind c_pp_lookup_pragma_by_name (const char *);
 
 extern GTY(()) tree pragma_extern_prefix;
 
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b9930d487fd..fa807530ca7 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -1583,8 +1583,12 @@ enum pragma_context { pragma_external, pragma_struct, pragma_param,
 static bool c_parser_pragma (c_parser *, enum pragma_context, bool *);
 static void c_parser_omp_cancellation_point (c_parser *, enum pragma_context);
 static bool c_parser_omp_target (c_parser *, enum pragma_context, bool *);
+static void c_parser_omp_begin (c_parser *, bool *);
+static void c_parser_omp_end (c_parser *);
 static void c_parser_omp_end_declare_target (c_parser *);
 static void c_parser_omp_declare (c_parser *, enum pragma_context);
+static tree c_parser_omp_metadirective (location_t, c_parser *, char *,
+					omp_clause_mask, tree *, bool *, bool);
 static void c_parser_omp_requires (c_parser *);
 static bool c_parser_omp_ordered (c_parser *, enum pragma_context, bool *);
 static void c_parser_oacc_routine (c_parser *, enum pragma_context);
@@ -12402,8 +12406,12 @@ c_parser_pragma (c_parser *parser, enum pragma_context context, bool *if_p)
     case PRAGMA_OMP_TARGET:
       return c_parser_omp_target (parser, context, if_p);
 
-    case PRAGMA_OMP_END_DECLARE_TARGET:
-      c_parser_omp_end_declare_target (parser);
+    case PRAGMA_OMP_BEGIN:
+      c_parser_omp_begin (parser, if_p);
+      return false;
+
+    case PRAGMA_OMP_END:
+      c_parser_omp_end (parser);
       return false;
 
     case PRAGMA_OMP_SCAN:
@@ -18195,6 +18203,7 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
   location_t for_loc;
   bool tiling = false;
   bool inscan = false;
+
   vec<tree, va_gc> *for_block = make_tree_vector ();
 
   for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
@@ -20934,6 +20943,60 @@ c_parser_omp_end_declare_target (c_parser *parser)
     current_omp_declare_target_attribute--;
 }
 
+static void
+c_parser_omp_begin (c_parser *parser, bool *if_p)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+  c_parser_consume_pragma(parser);
+  if (c_parser_peek_token (parser)->type == CPP_NAME)
+    {
+      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+
+      if (strcmp (p, "metadirective") == 0)
+	{
+	  char p_name[sizeof "#pragma omp teams distribute parallel for simd"];
+	  omp_clause_mask mask (0);
+
+	  c_parser_consume_token (parser);
+	  c_parser_omp_metadirective (loc, parser, p_name, mask, NULL, if_p,
+				      true);
+	  return;
+	}
+    }
+
+  error_at (loc, "expected %<begin metadirective%>");
+  c_parser_skip_to_pragma_eol (parser);
+}
+
+static void
+c_parser_omp_end (c_parser *parser)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+
+  if (c_parser_peek_2nd_token (parser)->type == CPP_NAME)
+    {
+      const char *p
+	= IDENTIFIER_POINTER (c_parser_peek_2nd_token (parser)->value);
+
+      if (strcmp (p, "declare") == 0)
+	{
+	  c_parser_omp_end_declare_target (parser);
+	  return;
+	}
+      else if (strcmp (p, "metadirective") == 0)
+	{
+	  /* The pragma 'omp end metadirective' should have been consumed
+	     when processing the metadirective.  */
+	  error_at (loc, "%<#pragma omp end metadirective%> without "
+			 "corresponding %<#pragma omp begin metadirective%>");
+	}
+    }
+  else
+    error_at (loc, "expected %<end declare target%> or %<end metadirective%>");
+
+  c_parser_skip_to_pragma_eol (parser);
+}
+
 
 /* OpenMP 4.0
    #pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -21607,6 +21670,295 @@ c_parser_omp_taskloop (location_t loc, c_parser *parser,
   return ret;
 }
 
+/* OpenMP 5.0:
+
+  # pragma omp metadirective [clause[, clause]]
+
+  # pragma omp begin metadirective [clause[, clause]]
+  # pragma omp end metadirective
+*/
+
+static tree
+c_parser_omp_metadirective (location_t loc, c_parser *parser,
+			    char *p_name, omp_clause_mask mask, tree *cclauses,
+			    bool *if_p,
+			    bool begin_end_p)
+{
+  tree ret;
+  bool all_selectors_resolveable = true;
+  auto_vec<auto_vec<c_token> > directive_tokens;
+  auto_vec<tree> ctxs;
+  bool default_seen = false;
+
+  ret = make_node (OMP_METADIRECTIVE);
+  SET_EXPR_LOCATION (ret, loc);
+  TREE_TYPE (ret) = void_type_node;
+  OMP_METADIRECTIVE_CLAUSES (ret) = NULL_TREE;
+  strcat (p_name, " metadirective");
+
+  while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL))
+    {
+      if (c_parser_next_token_is_not (parser, CPP_NAME)
+	  && c_parser_next_token_is_not (parser, CPP_KEYWORD))
+	{
+	  c_parser_error (parser, "expected %<when%> or %<default%>");
+	  return NULL_TREE;
+	}
+
+      location_t match_loc = c_parser_peek_token (parser)->location;
+      const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+      c_parser_consume_token (parser);
+      bool default_p = strcmp (p, "default") == 0;
+      if (default_p)
+	{
+	  if (default_seen)
+	    {
+	      c_parser_error (parser, "there can only be one default clause");
+	      return NULL_TREE;
+	    }
+	  else
+	    default_seen = true;
+	}
+      if (strcmp (p, "when") == 0 || default_p)
+	{
+	  matching_parens parens;
+	  tree ctx = NULL_TREE;
+	  bool skip = false;
+	  if (!parens.require_open (parser))
+	    return error_mark_node;
+
+	  if (!default_p)
+	    {
+	      ctx = c_parser_omp_context_selector_specification (parser,
+								 NULL_TREE);
+	      if (ctx == error_mark_node)
+		return NULL_TREE;
+	      ctx = c_omp_check_context_selector (match_loc, ctx);
+	      if (ctx == error_mark_node)
+		return NULL_TREE;
+
+	      switch (omp_context_selector_matches (ctx, true))
+		{
+		case -1:
+		  all_selectors_resolveable = false;
+		  break;
+		case 1:
+		  break;
+		case 0:
+		  /* Remove the selector from further consideration.  */
+		  skip = true;
+		  break;
+		}
+
+	      if (c_parser_next_token_is_not (parser, CPP_COLON))
+		{
+		  c_parser_error (parser, "expected colon");
+		  return NULL_TREE;
+		}
+	      c_parser_consume_token (parser);
+	    }
+
+	  /* Read in the directive type and create a dummy pragma token for
+	     it.  */
+	  c_token *token = c_parser_peek_token (parser);
+	  if (token->type != CPP_NAME)
+	    {
+	      c_parser_error (parser, "expected directive name");
+	      return NULL_TREE;
+	    }
+
+	  location_t loc = c_parser_peek_token (parser)->location;
+	  const char *p
+	    = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+	  enum pragma_kind p_kind = c_pp_lookup_pragma_by_name (p);
+
+	  c_parser_consume_token (parser);
+	  if (p_kind == PRAGMA_NONE)
+	    {
+	      c_parser_error (parser, "unknown directive name");
+	      return NULL_TREE;
+	    }
+
+	  if (!skip)
+	    {
+	      c_token pragma_token;
+	      pragma_token.type = CPP_PRAGMA;
+	      pragma_token.location = loc;
+	      pragma_token.pragma_kind = p_kind;
+
+	      directive_tokens.safe_push (auto_vec<c_token> ());
+	      directive_tokens.last ().safe_push (pragma_token);
+
+	      ctxs.safe_push (ctx);
+	    }
+
+	  /* Read in tokens for the directive clauses.  */
+	  auto_vec<c_token> *tokens = skip ? NULL : &directive_tokens.last ();
+	  int nesting_depth = 0;
+	  while (1)
+	    {
+	      c_token *token = c_parser_peek_token (parser);
+	      switch (token->type)
+		{
+		case CPP_EOF:
+		case CPP_PRAGMA_EOL:
+		  break;
+		case CPP_OPEN_PAREN:
+		  ++nesting_depth;
+		  goto add;
+		case CPP_CLOSE_PAREN:
+		  if (nesting_depth-- == 0)
+		    break;
+		  goto add;
+		default:
+		add:
+		  if (!skip)
+		    tokens->safe_push (*token);
+		  c_parser_consume_token (parser);
+		  continue;
+		}
+	      break;
+	    }
+
+	  c_parser_consume_token (parser);
+
+	  if (!skip)
+	    {
+	      c_token eol_token;
+	      memset (&eol_token, 0, sizeof (eol_token));
+	      eol_token.type = CPP_PRAGMA_EOL;
+	      tokens->safe_push (eol_token);
+	    }
+	}
+      else {
+	c_parser_error (parser, "expected %<when%> or %<default%>");
+	return NULL_TREE;
+      }
+    }
+  c_parser_skip_to_pragma_eol (parser);
+
+  /* Add the body tokens to the tokens for each candidate directive.  */
+  int nesting_depth = 0;
+  int bracket_depth = 0;
+  while (1)
+    {
+      int i;
+      auto_vec<c_token> *tokens;
+      c_token *token = c_parser_peek_token (parser);
+      bool stop = false;
+
+      if (begin_end_p)
+	{
+	  /* Keep reading until '#pragma end metadirective' is read.  */
+	  switch (token->type)
+	  {
+	  case CPP_PRAGMA:
+	    if (token->pragma_kind == PRAGMA_OMP_END)
+	      {
+		c_token *next_token = c_parser_peek_2nd_token (parser);
+		if (next_token->type == CPP_NAME
+		    && strcmp (IDENTIFIER_POINTER (next_token->value),
+			       "metadirective") == 0)
+		  {
+		    c_parser_consume_pragma (parser);
+		    c_parser_consume_token (parser);
+		    c_parser_skip_to_pragma_eol (parser);
+		    break;
+		  }
+	      }
+	  default:
+	    FOR_EACH_VEC_ELT (directive_tokens, i, tokens)
+	      tokens->safe_push (*token);
+	    if (token->type == CPP_PRAGMA)
+	      c_parser_consume_pragma (parser);
+	    else if (token->type == CPP_PRAGMA_EOL)
+	      c_parser_skip_to_pragma_eol (parser);
+	    else
+	      c_parser_consume_token (parser);
+	    continue;
+	  }
+	  break;
+	}
+      else
+	{
+	  switch (token->type)
+	    {
+	    case CPP_EOF:
+	      break;
+	    case CPP_OPEN_BRACE:
+	      ++nesting_depth;
+	      goto add2;
+	    case CPP_CLOSE_BRACE:
+	      if (--nesting_depth == 0)
+		stop = true;
+	      goto add2;
+	    case CPP_OPEN_PAREN:
+	      ++bracket_depth;
+	      goto add2;
+	    case CPP_CLOSE_PAREN:
+	      --bracket_depth;
+	      goto add2;
+	    case CPP_SEMICOLON:
+	      if (nesting_depth == 0 && bracket_depth == 0)
+		stop = true;
+	      goto add2;
+	    default:
+	    add2:
+	      FOR_EACH_VEC_ELT (directive_tokens, i, tokens)
+		tokens->safe_push (*token);
+	      if (token->type == CPP_PRAGMA)
+		c_parser_consume_pragma (parser);
+	      else if (token->type == CPP_PRAGMA_EOL)
+		c_parser_skip_to_pragma_eol (parser);
+	      else
+		c_parser_consume_token (parser);
+	      if (stop)
+		break;
+	      continue;
+	    }
+	  break;
+	}
+    }
+
+  /* Process each candidate directive.  */
+  auto_vec<c_token> *tokens;
+  int i;
+  FOR_EACH_VEC_ELT (directive_tokens, i, tokens)
+    {
+      /* Make sure nothing tries to read past the end of the tokens.  */
+      c_token eof_token;
+      memset (&eof_token, 0, sizeof (eof_token));
+      eof_token.type = CPP_EOF;
+      tokens->safe_push (eof_token);
+      tokens->safe_push (eof_token);
+
+      unsigned int tokens_avail = parser->tokens_avail;
+      gcc_assert (parser->tokens == &parser->tokens_buf[0]);
+      parser->tokens = tokens->address ();
+      parser->tokens_avail = tokens->length ();
+
+      tree block = c_begin_compound_stmt (false);
+      c_parser_omp_construct (parser, if_p);
+      block = c_end_compound_stmt (loc, block, false);
+
+      tree variant = build_tree_list (ctxs[i], block);
+      OMP_METADIRECTIVE_CLAUSES (ret)
+	= chainon (OMP_METADIRECTIVE_CLAUSES (ret), variant);
+
+      parser->tokens = &parser->tokens_buf[0];
+      parser->tokens_avail = tokens_avail;
+    }
+
+  if (all_selectors_resolveable)
+    {
+      ret = omp_resolve_metadirective (ret);
+      gcc_assert (ret != NULL_TREE);
+    }
+  add_stmt (ret);
+
+  return ret;
+}
+
 /* Main entry point to parsing most OpenMP pragmas.  */
 
 static void
@@ -21676,6 +22028,11 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
       strcpy (p_name, "#pragma omp");
       stmt = c_parser_omp_master (loc, parser, p_name, mask, NULL, if_p);
       break;
+    case PRAGMA_OMP_METADIRECTIVE:
+      strcpy (p_name, "#pragma omp");
+      stmt = c_parser_omp_metadirective (loc, parser, p_name, mask, NULL,
+					 if_p, false);
+      break;
     case PRAGMA_OMP_PARALLEL:
       strcpy (p_name, "#pragma omp");
       stmt = c_parser_omp_parallel (loc, parser, p_name, mask, NULL, if_p);
@@ -21713,7 +22070,6 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
     gcc_assert (EXPR_LOCATION (stmt) != UNKNOWN_LOCATION);
 }
 
-
 /* OpenMP 2.5:
    # pragma omp threadprivate (variable-list) */
 
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index f3503b13a5a..e2584c18571 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -43543,6 +43543,32 @@ cp_parser_omp_end_declare_target (cp_parser *parser, cp_token *pragma_tok)
     scope_chain->omp_declare_target_attribute--;
 }
 
+static void
+cp_parser_omp_end (cp_parser *parser, cp_token *pragma_tok)
+{
+  const char *p = "";
+  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+    {
+      tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+      p = IDENTIFIER_POINTER (id);
+    }
+  if (strcmp (p, "declare") == 0)
+    {
+      cp_parser_omp_end_declare_target (parser, pragma_tok);
+      return;
+    }
+  else if (strcmp (p, "metadirective") == 0)
+    /* The pragma 'omp end metadirective' should have been consumed
+       when processing the metadirective.  */
+    error_at (pragma_tok->location,
+	      "%<#pragma omp end metadirective%> without corresponding "
+	      "%<#pragma omp begin metadirective%>");
+  else
+    error_at (pragma_tok->location,
+	      "expected %<declare target%> or %<metadirective%>");
+  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+}
+
 /* Helper function of cp_parser_omp_declare_reduction.  Parse the combiner
    expression and optional initializer clause of
    #pragma omp declare reduction.  We store the expression(s) as
@@ -45259,8 +45285,8 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p)
       pop_omp_privatization_clauses (stmt);
       return ret;
 
-    case PRAGMA_OMP_END_DECLARE_TARGET:
-      cp_parser_omp_end_declare_target (parser, pragma_tok);
+    case PRAGMA_OMP_END:
+      cp_parser_omp_end (parser, pragma_tok);
       return false;
 
     case PRAGMA_OMP_SCAN:
diff --git a/gcc/gimple-low.c b/gcc/gimple-low.c
index fa7d4de30c0..b68b0cbb06c 100644
--- a/gcc/gimple-low.c
+++ b/gcc/gimple-low.c
@@ -234,6 +234,39 @@ lower_omp_directive (gimple_stmt_iterator *gsi, struct lower_data *data)
   gsi_next (gsi);
 }
 
+/* Lower the OpenMP metadirective statement pointed by GSI.  */
+
+static void
+lower_omp_metadirective (gimple_stmt_iterator *gsi, struct lower_data *data)
+{
+  gimple *stmt = gsi_stmt (*gsi);
+  gimple *body = gimple_omp_metadirective_bodies (stmt);
+  tree succ_label = create_artificial_label (UNKNOWN_LOCATION);
+  unsigned i;
+
+  for (i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+    {
+      tree label = create_artificial_label (UNKNOWN_LOCATION);
+      gimple *g = gimple_build_label (label);
+
+      gsi_insert_after (gsi, g, GSI_CONTINUE_LINKING);
+      lower_sequence (gimple_omp_body_ptr (body), data);
+      gsi_insert_seq_after (gsi, gimple_omp_body (body), GSI_CONTINUE_LINKING);
+      gsi_insert_after (gsi, gimple_build_goto (succ_label),
+			GSI_CONTINUE_LINKING);
+      gimple_omp_metadirective_set_label (stmt, i, label);
+
+      body = body->next;
+    }
+
+  gsi_insert_after (gsi, gimple_build_label (succ_label),
+		    GSI_CONTINUE_LINKING);
+  gimple_omp_metadirective_set_succ_label (stmt, succ_label);
+  gimple_omp_metadirective_set_bodies (stmt, NULL);
+
+  gsi_next (gsi);
+}
+
 
 /* Lower statement GSI.  DATA is passed through the recursion.  We try to
    track the fallthruness of statements and get rid of unreachable return
@@ -398,6 +431,12 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data)
       data->cannot_fallthru = false;
       return;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      data->cannot_fallthru = false;
+      lower_omp_metadirective (gsi, data);
+      data->cannot_fallthru = false;
+      return;
+
     case GIMPLE_TRANSACTION:
       lower_sequence (gimple_transaction_body_ptr (
 			as_a <gtransaction *> (stmt)),
diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index 0ef01e6420b..e7a4ba1171b 100644
--- a/gcc/gimple-pretty-print.c
+++ b/gcc/gimple-pretty-print.c
@@ -1978,6 +1978,64 @@ dump_gimple_omp_return (pretty_printer *buffer, const gimple *gs, int spc,
     }
 }
 
+/* Dump a GIMPLE_OMP_METADIRECTIVE tuple on the pretty_printer BUFFER.  */
+
+static void
+dump_gimple_omp_metadirective (pretty_printer *buffer, const gimple *gs,
+			       int spc, dump_flags_t flags)
+{
+  if (flags & TDF_RAW)
+    {
+      dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S> >", gs,
+		       gimple_omp_body (gs));
+    }
+  else
+    {
+      pp_string (buffer, "#pragma omp metadirective");
+      newline_and_indent (buffer, spc + 2);
+
+      gimple *body = gimple_omp_metadirective_bodies (gs);
+      bool has_bodies_p = body != NULL;
+      unsigned num_clauses = gimple_omp_metadirective_num_clauses (gs);
+
+      for (unsigned i = 0; i < num_clauses; i++)
+	{
+	  tree selector = gimple_omp_metadirective_selector (gs, i);
+
+	  if (selector == NULL_TREE)
+	    pp_string (buffer, "default:");
+	  else
+	    {
+	      pp_string (buffer, "when (");
+	      dump_generic_node (buffer, selector, spc, flags, false);
+	      pp_string (buffer, "):");
+	    }
+
+	  if (has_bodies_p)
+	    {
+	      newline_and_indent (buffer, spc + 4);
+	      pp_left_brace (buffer);
+	      pp_newline (buffer);
+	      dump_gimple_seq (buffer, gimple_omp_body (body), spc + 6, flags);
+	      newline_and_indent (buffer, spc + 4);
+	      pp_right_brace (buffer);
+
+	      body = body->next;
+	      if (body)
+		newline_and_indent (buffer, spc + 2);
+	    }
+	  else
+	    {
+	      tree label = gimple_omp_metadirective_label (gs, i);
+	      pp_string (buffer, " ");
+	      dump_generic_node (buffer, label, spc, flags, false);
+	      if (i != num_clauses - 1)
+		newline_and_indent (buffer, spc + 2);
+	    }
+	}
+    }
+}
+
 /* Dump a GIMPLE_TRANSACTION tuple on the pretty_printer BUFFER.  */
 
 static void
@@ -2729,6 +2787,12 @@ pp_gimple_stmt_1 (pretty_printer *buffer, const gimple *gs, int spc,
 				flags);
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      dump_gimple_omp_metadirective (buffer,
+				     as_a <const gomp_metadirective *> (gs),
+				     spc, flags);
+      break;
+
     case GIMPLE_CATCH:
       dump_gimple_catch (buffer, as_a <const gcatch *> (gs), spc, flags);
       break;
diff --git a/gcc/gimple-streamer-in.c b/gcc/gimple-streamer-in.c
index 1c979f438a5..66bbc2e8e0a 100644
--- a/gcc/gimple-streamer-in.c
+++ b/gcc/gimple-streamer-in.c
@@ -151,6 +151,7 @@ input_gimple_stmt (class lto_input_block *ib, class data_in *data_in,
     case GIMPLE_COND:
     case GIMPLE_GOTO:
     case GIMPLE_DEBUG:
+    case GIMPLE_OMP_METADIRECTIVE:
       for (i = 0; i < num_ops; i++)
 	{
 	  tree *opp, op = stream_read_tree (ib, data_in);
@@ -188,6 +189,10 @@ input_gimple_stmt (class lto_input_block *ib, class data_in *data_in,
 	  else
 	    gimple_call_set_fntype (call_stmt, stream_read_tree (ib, data_in));
 	}
+      if (gomp_metadirective *metadirective_stmt
+	    = dyn_cast <gomp_metadirective*> (stmt))
+	gimple_omp_metadirective_set_succ_label (metadirective_stmt,
+						 stream_read_tree (ib, data_in));
       break;
 
     case GIMPLE_NOP:
diff --git a/gcc/gimple-streamer-out.c b/gcc/gimple-streamer-out.c
index fcbf92300d4..f0ddd6a81a3 100644
--- a/gcc/gimple-streamer-out.c
+++ b/gcc/gimple-streamer-out.c
@@ -127,6 +127,7 @@ output_gimple_stmt (struct output_block *ob, struct function *fn, gimple *stmt)
     case GIMPLE_COND:
     case GIMPLE_GOTO:
     case GIMPLE_DEBUG:
+    case GIMPLE_OMP_METADIRECTIVE:
       for (i = 0; i < gimple_num_ops (stmt); i++)
 	{
 	  tree op = gimple_op (stmt, i);
@@ -169,6 +170,8 @@ output_gimple_stmt (struct output_block *ob, struct function *fn, gimple *stmt)
 	  else
 	    stream_write_tree (ob, gimple_call_fntype (stmt), true);
 	}
+      if (gimple_code (stmt) == GIMPLE_OMP_METADIRECTIVE)
+	stream_write_tree (ob, gimple_omp_metadirective_succ_label (stmt), true);
       break;
 
     case GIMPLE_NOP:
diff --git a/gcc/gimple-walk.c b/gcc/gimple-walk.c
index e4a55f1eeb6..ad7e1c0839e 100644
--- a/gcc/gimple-walk.c
+++ b/gcc/gimple-walk.c
@@ -674,6 +674,21 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt,
 	return wi->callback_result;
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      {
+	gimple *body = gimple_omp_metadirective_bodies (stmt);
+
+	while (body)
+	  {
+	    ret = walk_gimple_seq_mod (gimple_omp_body_ptr (body),
+				       callback_stmt, callback_op, wi);
+	    if (ret)
+	      return wi->callback_result;
+	    body = body->next;
+	  }
+      }
+      break;
+
     case GIMPLE_WITH_CLEANUP_EXPR:
       ret = walk_gimple_seq_mod (gimple_wce_cleanup_ptr (stmt), callback_stmt,
 			     callback_op, wi);
diff --git a/gcc/gimple.c b/gcc/gimple.c
index f1044e9c630..f94009d39f3 100644
--- a/gcc/gimple.c
+++ b/gcc/gimple.c
@@ -1234,6 +1234,28 @@ gimple_build_omp_atomic_store (tree val, enum omp_memory_order mo)
   return p;
 }
 
+/* Build a GIMPLE_OMP_METADIRECTIVE statement.  */
+
+gomp_metadirective *
+gimple_build_omp_metadirective (int clause_count)
+{
+  gomp_metadirective *p
+    = as_a <gomp_metadirective *> (gimple_alloc (GIMPLE_OMP_METADIRECTIVE,
+						 clause_count * 2));
+  gimple_omp_metadirective_set_bodies (p, NULL);
+  return p;
+}
+
+
+gomp_metadirective_body *
+gimple_build_omp_metadirective_body (gimple_seq body)
+{
+  gomp_metadirective_body *m_body = as_a <gomp_metadirective_body *>
+    (gimple_alloc (GIMPLE_OMP_METADIRECTIVE_BODY, 0));
+  gimple_omp_set_body (m_body, body);
+  return m_body;
+}
+
 /* Build a GIMPLE_TRANSACTION statement.  */
 
 gtransaction *
diff --git a/gcc/gimple.def b/gcc/gimple.def
index 0ac0cf72bfa..1da68c16a91 100644
--- a/gcc/gimple.def
+++ b/gcc/gimple.def
@@ -384,6 +384,13 @@ DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_PARALLEL_LAYOUT)
    CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
 DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_omp_ordered", GSS_OMP_SINGLE_LAYOUT)
 
+/* GIMPLE_OMP_METADIRECTIVE represents #pragma omp metadirective.  */
+DEFGSCODE(GIMPLE_OMP_METADIRECTIVE, "gimple_omp_metadirective",
+	  GSS_OMP_METADIRECTIVE)
+
+DEFGSCODE(GIMPLE_OMP_METADIRECTIVE_BODY, "gimple_omp_metadirective_body",
+	  GSS_OMP_METADIRECTIVE_BODY)
+
 /* GIMPLE_PREDICT <PREDICT, OUTCOME> specifies a hint for branch prediction.
 
    PREDICT is one of the predictors from predict.def.
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 91b92b4a4d1..c5288af2bd7 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -822,6 +822,29 @@ struct GTY((tag("GSS_OMP_ATOMIC_STORE_LAYOUT")))
          stmt->code == GIMPLE_OMP_RETURN.  */
 };
 
+struct GTY((tag("GSS_OMP_METADIRECTIVE_BODY")))
+  gomp_metadirective_body : public gimple_statement_omp_single_layout
+{
+    /* No extra fields; adds invariant:
+       stmt->code == GIMPLE_OMP_METADIRECTIVE_BODY.  */
+};
+
+struct GTY((tag("GSS_OMP_METADIRECTIVE")))
+  gomp_metadirective : public gimple_statement_with_ops_base
+{
+  /* [ WORD 1-7 ] : base class */
+
+  /* [ WORD 8 ]  */
+  gomp_metadirective_body *bodies;
+
+  /* [ WORD 9 ] : a label after the metadirective
+     and all the candidate bodies  */
+  tree succ_label;
+
+  /* [ WORD 10 ] : operand vector.  */
+  tree GTY((length ("%h.num_ops"))) op[1];
+};
+
 /* GIMPLE_TRANSACTION.  */
 
 /* Bits to be stored in the GIMPLE_TRANSACTION subcode.  */
@@ -1233,6 +1256,22 @@ is_a_helper <gomp_task *>::test (gimple *gs)
   return gs->code == GIMPLE_OMP_TASK;
 }
 
+template <>
+template <>
+inline bool
+is_a_helper <gomp_metadirective *>::test (gimple *gs)
+{
+  return gs->code == GIMPLE_OMP_METADIRECTIVE;
+}
+
+template <>
+template <>
+inline bool
+is_a_helper <gomp_metadirective_body *>::test (gimple *gs)
+{
+  return gs->code == GIMPLE_OMP_METADIRECTIVE_BODY;
+}
+
 template <>
 template <>
 inline bool
@@ -1475,6 +1514,22 @@ is_a_helper <const gomp_task *>::test (const gimple *gs)
   return gs->code == GIMPLE_OMP_TASK;
 }
 
+template <>
+template <>
+inline bool
+is_a_helper <const gomp_metadirective *>::test (const gimple *gs)
+{
+  return gs->code == GIMPLE_OMP_METADIRECTIVE;
+}
+
+template <>
+template <>
+inline bool
+is_a_helper <const gomp_metadirective_body *>::test (const gimple *gs)
+{
+  return gs->code == GIMPLE_OMP_METADIRECTIVE_BODY;
+}
+
 template <>
 template <>
 inline bool
@@ -1572,6 +1627,8 @@ gomp_teams *gimple_build_omp_teams (gimple_seq, tree);
 gomp_atomic_load *gimple_build_omp_atomic_load (tree, tree,
 						enum omp_memory_order);
 gomp_atomic_store *gimple_build_omp_atomic_store (tree, enum omp_memory_order);
+gomp_metadirective *gimple_build_omp_metadirective (int clause_count);
+gomp_metadirective_body *gimple_build_omp_metadirective_body (gimple_seq body);
 gtransaction *gimple_build_transaction (gimple_seq);
 extern void gimple_seq_add_stmt (gimple_seq *, gimple *);
 extern void gimple_seq_add_stmt_without_update (gimple_seq *, gimple *);
@@ -1827,6 +1884,7 @@ gimple_has_substatements (gimple *g)
     case GIMPLE_OMP_TARGET:
     case GIMPLE_OMP_TEAMS:
     case GIMPLE_OMP_CRITICAL:
+    case GIMPLE_OMP_METADIRECTIVE:
     case GIMPLE_WITH_CLEANUP_EXPR:
     case GIMPLE_TRANSACTION:
       return true;
@@ -2479,12 +2537,21 @@ gimple_ops (gimple *gs)
 }
 
 
+/* Return true if GIMPLE statement G has any operands, including any that
+   should not be processed by the SSA passes.  */
+
+static inline bool
+gimple_has_ops_1 (const gimple *g)
+{
+  return gimple_has_ops (g) || gimple_code (g) == GIMPLE_OMP_METADIRECTIVE;
+}
+
 /* Return operand I for statement GS.  */
 
 static inline tree
 gimple_op (const gimple *gs, unsigned i)
 {
-  if (gimple_has_ops (gs))
+  if (gimple_has_ops_1 (gs))
     {
       gcc_gimple_checking_assert (i < gimple_num_ops (gs));
       return gimple_ops (CONST_CAST_GIMPLE (gs))[i];
@@ -2498,7 +2565,7 @@ gimple_op (const gimple *gs, unsigned i)
 static inline tree *
 gimple_op_ptr (gimple *gs, unsigned i)
 {
-  if (gimple_has_ops (gs))
+  if (gimple_has_ops_1 (gs))
     {
       gcc_gimple_checking_assert (i < gimple_num_ops (gs));
       return gimple_ops (gs) + i;
@@ -2512,7 +2579,7 @@ gimple_op_ptr (gimple *gs, unsigned i)
 static inline void
 gimple_set_op (gimple *gs, unsigned i, tree op)
 {
-  gcc_gimple_checking_assert (gimple_has_ops (gs) && i < gimple_num_ops (gs));
+  gcc_gimple_checking_assert (gimple_has_ops_1 (gs) && i < gimple_num_ops (gs));
 
   /* Note.  It may be tempting to assert that OP matches
      is_gimple_operand, but that would be wrong.  Different tuples
@@ -6330,6 +6397,77 @@ gimple_omp_continue_set_control_use (gomp_continue *cont_stmt, tree use)
   cont_stmt->control_use = use;
 }
 
+
+static inline tree
+gimple_omp_metadirective_succ_label (const gimple *g)
+{
+  const gomp_metadirective *omp_metadirective
+    = as_a <const gomp_metadirective *> (g);
+  return omp_metadirective->succ_label;
+}
+
+
+static inline void
+gimple_omp_metadirective_set_succ_label (gimple *g, tree succ_label)
+{
+  gomp_metadirective *omp_metadirective = as_a <gomp_metadirective *> (g);
+  omp_metadirective->succ_label = succ_label;
+}
+
+
+static inline gomp_metadirective_body *
+gimple_omp_metadirective_bodies (const gimple *g)
+{
+  const gomp_metadirective *omp_metadirective
+    = as_a <const gomp_metadirective *> (g);
+  return omp_metadirective->bodies;
+}
+
+
+static inline void
+gimple_omp_metadirective_set_bodies (gimple *g,
+				     gomp_metadirective_body *bodies)
+{
+  gomp_metadirective *omp_metadirective = as_a <gomp_metadirective *> (g);
+  omp_metadirective->bodies = bodies;
+}
+
+
+static inline unsigned
+gimple_omp_metadirective_num_clauses (const gimple *g)
+{
+  return gimple_num_ops (g) / 2;
+}
+
+
+static inline tree
+gimple_omp_metadirective_selector (const gimple *g, unsigned n)
+{
+  return gimple_op (g, n * 2);
+}
+
+
+static inline void
+gimple_omp_metadirective_set_selector (gimple *g, unsigned n, tree selector)
+{
+  return gimple_set_op (g, n * 2, selector);
+}
+
+
+static inline tree
+gimple_omp_metadirective_label (const gimple *g, unsigned n)
+{
+  return gimple_op (g, n * 2 + 1);
+}
+
+
+static inline void
+gimple_omp_metadirective_set_label (gimple *g, unsigned n, tree label)
+{
+  return gimple_set_op (g, n * 2 + 1, label);
+}
+
+
 /* Return a pointer to the body for the GIMPLE_TRANSACTION statement
    TRANSACTION_STMT.  */
 
@@ -6478,6 +6616,7 @@ gimple_return_set_retval (greturn *gs, tree retval)
     case GIMPLE_OMP_RETURN:			\
     case GIMPLE_OMP_ATOMIC_LOAD:		\
     case GIMPLE_OMP_ATOMIC_STORE:		\
+    case GIMPLE_OMP_METADIRECTIVE:		\
     case GIMPLE_OMP_CONTINUE
 
 static inline bool
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 2730f225187..49f4bcb01f3 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -5646,6 +5646,7 @@ is_gimple_stmt (tree t)
     case OMP_TASKGROUP:
     case OMP_ORDERED:
     case OMP_CRITICAL:
+    case OMP_METADIRECTIVE:
     case OMP_TASK:
     case OMP_TARGET:
     case OMP_TARGET_DATA:
@@ -13785,6 +13786,49 @@ gimplify_omp_ordered (tree expr, gimple_seq body)
   return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr));
 }
 
+static enum gimplify_status
+gimplify_omp_metadirective (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
+			    bool (*gimple_test_f) (tree), fallback_t fallback)
+{
+  gomp_metadirective_body *first_body = NULL;
+  gomp_metadirective_body *prev_body = NULL;
+  auto_vec<tree> selectors;
+  unsigned i;
+  tree clause = OMP_METADIRECTIVE_CLAUSES (*expr_p);
+
+  while (clause)
+    {
+      tree selector = TREE_PURPOSE (clause);
+      tree directive = TREE_VALUE (clause);
+
+      selectors.safe_push (selector);
+      gomp_metadirective_body *body
+	= gimple_build_omp_metadirective_body (NULL);
+      gimplify_stmt (&directive, gimple_omp_body_ptr (body));
+      if (!first_body)
+	first_body = body;
+      if (prev_body)
+	{
+	  prev_body->next = body;
+	  body->prev = prev_body;
+	}
+      prev_body = body;
+
+      clause = TREE_CHAIN (clause);
+    }
+
+  gomp_metadirective *stmt
+    = gimple_build_omp_metadirective (selectors.length ());
+  gimple_omp_metadirective_set_bodies (stmt, first_body);
+  gimplify_seq_add_stmt (pre_p, stmt);
+
+  tree selector;
+  FOR_EACH_VEC_ELT (selectors, i, selector)
+    gimple_omp_metadirective_set_selector (stmt, i, selector);
+
+  return GS_ALL_DONE;
+}
+
 /* Convert the GENERIC expression tree *EXPR_P to GIMPLE.  If the
    expression produces a value to be used as an operand inside a GIMPLE
    statement, the value will be stored back in *EXPR_P.  This value will
@@ -14680,6 +14724,11 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	  ret = gimplify_omp_atomic (expr_p, pre_p);
 	  break;
 
+	case OMP_METADIRECTIVE:
+	  ret = gimplify_omp_metadirective (expr_p, pre_p, post_p,
+					    gimple_test_f, fallback);
+	  break;
+
 	case TRANSACTION_EXPR:
 	  ret = gimplify_transaction (expr_p, pre_p);
 	  break;
diff --git a/gcc/gsstruct.def b/gcc/gsstruct.def
index 8f777e2bb95..f22ac1f65cb 100644
--- a/gcc/gsstruct.def
+++ b/gcc/gsstruct.def
@@ -50,4 +50,6 @@ DEFGSSTRUCT(GSS_OMP_SINGLE_LAYOUT, gimple_statement_omp_single_layout, false)
 DEFGSSTRUCT(GSS_OMP_CONTINUE, gomp_continue, false)
 DEFGSSTRUCT(GSS_OMP_ATOMIC_LOAD, gomp_atomic_load, false)
 DEFGSSTRUCT(GSS_OMP_ATOMIC_STORE_LAYOUT, gomp_atomic_store, false)
+DEFGSSTRUCT(GSS_OMP_METADIRECTIVE, gomp_metadirective, true)
+DEFGSSTRUCT(GSS_OMP_METADIRECTIVE_BODY, gomp_metadirective_body, false)
 DEFGSSTRUCT(GSS_TRANSACTION, gtransaction, false)
diff --git a/gcc/omp-expand-metadirective.cc b/gcc/omp-expand-metadirective.cc
new file mode 100644
index 00000000000..f4620df577a
--- /dev/null
+++ b/gcc/omp-expand-metadirective.cc
@@ -0,0 +1,125 @@
+/* Expand an OpenMP metadirective.
+
+   Copyright (C) 2021 Free Software Foundation, Inc.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify it under
+the terms of the GNU General Public License as published by the Free
+Software Foundation; either version 3, or (at your option) any later
+version.
+
+GCC is distributed in the hope that it will be useful, but WITHOUT ANY
+WARRANTY; without even the implied warranty of MERCHANTABILITY or
+FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+for more details.
+
+You should have received a copy of the GNU General Public License
+along with GCC; see the file COPYING3.  If not see
+<http://www.gnu.org/licenses/>.  */
+
+#include "config.h"
+#include "system.h"
+#include "coretypes.h"
+#include "backend.h"
+#include "target.h"
+#include "tree.h"
+#include "langhooks.h"
+#include "gimple.h"
+#include "tree-pass.h"
+#include "cgraph.h"
+#include "fold-const.h"
+#include "gimplify.h"
+#include "gimple-iterator.h"
+#include "gimple-walk.h"
+#include "gomp-constants.h"
+#include "omp-general.h"
+#include "diagnostic-core.h"
+#include "tree-cfg.h"
+#include "cfganal.h"
+
+static void
+omp_expand_metadirective (function *fun, basic_block bb)
+{
+  gimple *stmt = last_stmt (bb);
+  tree selected_label = omp_resolve_metadirective (stmt);
+
+  /* This is the last chance for the metadirective to be resolved.  */
+  if (!selected_label)
+    gcc_unreachable ();
+
+  /* Delete all variant BBs except for the selected one.  */
+  calculate_dominance_info (CDI_DOMINATORS);
+  for (unsigned i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+    {
+      tree label = gimple_omp_metadirective_label (stmt, i);
+      edge edge = find_edge (bb, label_to_block (fun, label));
+      if (label == selected_label)
+	edge->flags |= EDGE_FALLTHRU;
+      else
+	remove_edge_and_dominated_blocks (edge);
+    }
+
+  /* Remove the metadirective statement.  */
+  gimple_stmt_iterator gsi = gsi_last_bb (bb);
+  gsi_remove (&gsi, true);
+}
+
+namespace {
+
+const pass_data pass_data_omp_expand_metadirective =
+{
+  GIMPLE_PASS, /* type */
+  "omp_expand_metadirective", /* name */
+  OPTGROUP_OMP, /* optinfo_flags */
+  TV_NONE, /* tv_id */
+  PROP_gimple_lcf, /* properties_required */
+  0, /* properties_provided */
+  0, /* properties_destroyed */
+  0, /* todo_flags_start */
+  0, /* todo_flags_finish */
+};
+
+class pass_omp_expand_metadirective : public gimple_opt_pass
+{
+public:
+  pass_omp_expand_metadirective (gcc::context *ctxt)
+    : gimple_opt_pass (pass_data_omp_expand_metadirective, ctxt)
+  {}
+
+  /* opt_pass methods: */
+  virtual bool gate (function *)
+  {
+    return (flag_openmp);
+  }
+
+  virtual unsigned int execute (function *fun);
+}; // class pass_omp_oacc_kernels_decompose
+
+unsigned int
+pass_omp_expand_metadirective::execute (function *fun)
+{
+  basic_block bb;
+  auto_vec<basic_block> metadirective_bbs;
+
+  FOR_EACH_BB_FN (bb, fun)
+    {
+      gimple *stmt = last_stmt (bb);
+      if (stmt && is_a<gomp_metadirective *> (stmt))
+	metadirective_bbs.safe_push (bb);
+    }
+
+  for (unsigned i = 0; i < metadirective_bbs.length (); i++)
+    omp_expand_metadirective (fun, metadirective_bbs[i]);
+
+  return 0;
+}
+
+} // anon namespace
+
+
+gimple_opt_pass *
+make_pass_omp_expand_metadirective (gcc::context *ctxt)
+{
+  return new pass_omp_expand_metadirective (ctxt);
+}
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 0f843bad79a..2c1affb64f8 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -9866,6 +9866,22 @@ expand_omp_target (struct omp_region *region)
     }
 }
 
+static void
+expand_omp_metadirective (struct omp_region *region)
+{
+  gomp_metadirective *stmt
+    = as_a <gomp_metadirective *> (last_stmt (region->entry));
+  tree succ_label = gimple_omp_metadirective_succ_label (stmt);
+  basic_block succ_bb = label_to_block (cfun, succ_label);
+  gimple_stmt_iterator gsi = gsi_start_bb (succ_bb);
+  while (!gsi_end_p (gsi)
+	 && gimple_code (gsi_stmt (gsi)) != GIMPLE_OMP_RETURN)
+    gsi_next (&gsi);
+
+  gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN);
+  gsi_remove (&gsi, true);
+}
+
 /* Expand the parallel region tree rooted at REGION.  Expansion
    proceeds in depth-first order.  Innermost regions are expanded
    first.  This way, parallel regions that require a new function to
@@ -9952,6 +9968,10 @@ expand_omp (struct omp_region *region)
 	  expand_omp_target (region);
 	  break;
 
+	case GIMPLE_OMP_METADIRECTIVE:
+	  expand_omp_metadirective (region);
+	  break;
+
 	default:
 	  gcc_unreachable ();
 	}
@@ -10388,6 +10408,24 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region,
 	}
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      /* Create an edge to the beginning of the body of each candidate
+	 directive.  */
+      {
+	unsigned i;
+	cur_region = new_omp_region (bb, code, cur_region);
+	gimple *stmt = last_stmt (bb);
+	for (i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+	  {
+	    tree dest = gimple_omp_metadirective_label (stmt, i);
+	    basic_block dest_bb = label_to_block (cfun, dest);
+	    make_edge (bb, dest_bb, 0);
+	  }
+
+	fallthru = false;
+      }
+      break;
+
     default:
       gcc_unreachable ();
     }
diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index a1bb9d8d25d..4839a9849bc 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -44,6 +44,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "tree-iterator.h"
 #include "data-streamer.h"
 #include "streamer-hooks.h"
+#include "tree-pretty-print.h"
 
 enum omp_requires omp_requires_mask;
 
@@ -1100,8 +1101,13 @@ omp_context_name_list_prop (tree prop)
    others need to wait until the whole TU is parsed, others need to wait until
    IPA, others until vectorization.  */
 
+#define DELAY_METADIRECTIVES_AFTER_LTO { \
+  if (metadirective_p && !(cfun->curr_properties & PROP_gimple_lomp_dev))	\
+    return -1;	\
+}
+
 int
-omp_context_selector_matches (tree ctx)
+omp_context_selector_matches (tree ctx, bool metadirective_p)
 {
   int ret = 1;
   for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
@@ -1222,6 +1228,8 @@ omp_context_selector_matches (tree ctx)
 		    const char *arch = omp_context_name_list_prop (t3);
 		    if (arch == NULL)
 		      return 0;
+		    DELAY_METADIRECTIVES_AFTER_LTO;
+
 		    int r = 0;
 		    if (targetm.omp.device_kind_arch_isa != NULL)
 		      r = targetm.omp.device_kind_arch_isa (omp_device_arch,
@@ -1340,6 +1348,8 @@ omp_context_selector_matches (tree ctx)
 			  return 0;
 			continue;
 		      }
+		    DELAY_METADIRECTIVES_AFTER_LTO;
+
 		    int r = 0;
 		    if (targetm.omp.device_kind_arch_isa != NULL)
 		      r = targetm.omp.device_kind_arch_isa (omp_device_kind,
@@ -1379,6 +1389,8 @@ omp_context_selector_matches (tree ctx)
 		    const char *isa = omp_context_name_list_prop (t3);
 		    if (isa == NULL)
 		      return 0;
+		    DELAY_METADIRECTIVES_AFTER_LTO;
+
 		    int r = 0;
 		    if (targetm.omp.device_kind_arch_isa != NULL)
 		      r = targetm.omp.device_kind_arch_isa (omp_device_isa,
@@ -1445,6 +1457,8 @@ omp_context_selector_matches (tree ctx)
   return ret;
 }
 
+#undef DELAY_METADIRECTIVES_AFTER_LTO
+
 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
    in omp_context_selector_set_compare.  */
 
@@ -2459,6 +2473,161 @@ omp_lto_input_declare_variant_alt (lto_input_block *ib, cgraph_node *node,
 						 INSERT) = entryp;
 }
 
+tree
+omp_resolve_metadirective (tree metadirective)
+{
+  auto_vec <tree, 16> clauses;
+  auto_vec <widest_int, 16> scores;
+  tree clause = OMP_METADIRECTIVE_CLAUSES (metadirective);
+  tree default_variant = NULL_TREE;
+
+  while (clause)
+    {
+      tree selector = TREE_PURPOSE (clause);
+      widest_int score;
+
+      if (selector == NULL_TREE)
+	default_variant = TREE_VALUE (clause);
+      else
+	switch (omp_context_selector_matches (selector, true))
+	  {
+	  case -1:
+	    return NULL_TREE;
+	  case 1:
+	    clauses.safe_push (clause);
+	    /* TODO: Handle SIMD score?  */
+	    omp_context_compute_score (selector, &score, false);
+	    scores.safe_push (score);
+	    break;
+	  case 0:
+	    break;
+	  }
+      clause = TREE_CHAIN (clause);
+    }
+
+  /* TODO: Handle case where there is no default.  */
+  if (clauses.is_empty ())
+    {
+      if (dump_file)
+	fprintf (dump_file, "Selecting default directive variant\n");
+      return default_variant;
+    }
+
+  /* A context selector that is a strict subset of another context selector
+     has a score of zero.  */
+  tree clause1, clause2;
+  unsigned int i, j;
+  FOR_EACH_VEC_ELT (clauses, i, clause1)
+    FOR_EACH_VEC_ELT_FROM (clauses, j, clause2, i + 1)
+      {
+	int r = omp_context_selector_compare (TREE_PURPOSE (clause1),
+					      TREE_PURPOSE (clause2));
+	if (r == -1)
+	  {
+	    /* ctx1 is a strict subset of ctx2.  */
+	    scores[i] = 0;
+	    break;
+	  }
+	else if (r == 1)
+	  /* ctx2 is a strict subset of ctx1.  */
+	  scores[j] = 0;
+      }
+
+  widest_int score, highest_score = -1;
+  FOR_EACH_VEC_ELT (scores, i, score)
+    if (score > highest_score)
+      {
+	highest_score = score;
+	clause = clauses[i];
+      }
+
+  if (dump_file)
+    {
+      fprintf (dump_file, "Selecting directive variant with selector:");
+      print_generic_expr (dump_file, TREE_PURPOSE (clause));
+      fprintf (dump_file, "\n");
+    }
+  return TREE_VALUE (clause);
+}
+
+tree
+omp_resolve_metadirective (gimple *gs)
+{
+  auto_vec <tree, 16> labels;
+  auto_vec <tree, 16> selectors;
+  auto_vec <widest_int, 16> scores;
+  tree default_label = gimple_omp_metadirective_succ_label (gs);
+
+  for (unsigned i = 0; i < gimple_omp_metadirective_num_clauses (gs); i++)
+    {
+      tree selector = gimple_omp_metadirective_selector (gs, i);
+      widest_int score;
+      if (selector == NULL_TREE)
+	default_label = gimple_omp_metadirective_label (gs, i);
+      else
+	switch (omp_context_selector_matches (selector, true))
+	  {
+	  case -1:
+	    return NULL;
+	  case 1:
+	    labels.safe_push (gimple_omp_metadirective_label (gs, i));
+	    selectors.safe_push (selector);
+	    /* TODO: Handle SIMD score?.  */
+	    omp_context_compute_score (selector, &score, false);
+	    scores.safe_push (score);
+	    break;
+	  case 0:
+	    break;
+	  }
+    }
+
+  if (scores.is_empty ())
+    {
+      if (dump_file)
+	fprintf (dump_file, "Selecting default directive variant\n");
+      return default_label;
+    }
+
+  /* A context selector that is a strict subset of another context selector
+     has a score of zero.  */
+  tree ctx1, ctx2;
+  unsigned int i, j;
+  FOR_EACH_VEC_ELT (selectors, i, ctx1)
+    FOR_EACH_VEC_ELT_FROM (selectors, j, ctx2, i + 1)
+      {
+	int r = omp_context_selector_compare (ctx1, ctx2);
+	if (r == -1)
+	  {
+	    /* ctx1 is a strict subset of ctx2.  */
+	    scores[i] = 0;
+	    break;
+	  }
+	else if (r == 1)
+	  /* ctx2 is a strict subset of ctx1.  */
+	  scores[j] = 0;
+      }
+
+  unsigned highest_index = 0;
+  widest_int score, highest_score = -1;
+  FOR_EACH_VEC_ELT (scores, i, score)
+  {
+    if (score > highest_score)
+      {
+	highest_score = score;
+	highest_index = i;
+      }
+  }
+
+  if (dump_file)
+    {
+      fprintf (dump_file, "Selecting directive variant with selector:");
+      print_generic_expr (dump_file, selectors[highest_index]);
+      fprintf (dump_file, "\n");
+    }
+
+  return labels[highest_index];
+}
+
 /* Encode an oacc launch argument.  This matches the GOMP_LAUNCH_PACK
    macro on gomp-constants.h.  We do not check for overflow.  */
 
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index aa04895e16d..47cea2eae01 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -104,10 +104,12 @@ extern tree find_combined_omp_for (tree *, int *, void *);
 extern poly_uint64 omp_max_vf (void);
 extern int omp_max_simt_vf (void);
 extern int omp_constructor_traits_to_codes (tree, enum tree_code *);
-extern int omp_context_selector_matches (tree);
+extern int omp_context_selector_matches (tree, bool = false);
 extern int omp_context_selector_set_compare (const char *, tree, tree);
 extern tree omp_get_context_selector (tree, const char *, const char *);
 extern tree omp_resolve_declare_variant (tree);
+extern tree omp_resolve_metadirective (tree);
+extern tree omp_resolve_metadirective (gimple *);
 extern tree oacc_launch_pack (unsigned code, tree device, unsigned op);
 extern tree oacc_replace_fn_attrib_attr (tree attribs, tree dims);
 extern void oacc_replace_fn_attrib (tree fn, tree dims);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index d1136d181b3..c54000cfb9d 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -179,6 +179,10 @@ struct omp_context
   /* Only used for omp target contexts.  True if an OpenMP construct other
      than teams is strictly nested in it.  */
   bool nonteams_nested_p;
+
+  /* Only used for omp metadirectives.  Links to the next shallow
+     clone of this context.  */
+  struct omp_context *next_clone;
 };
 
 static splay_tree all_contexts;
@@ -964,6 +968,7 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
   splay_tree_insert (all_contexts, (splay_tree_key) stmt,
 		     (splay_tree_value) ctx);
   ctx->stmt = stmt;
+  ctx->next_clone = NULL;
 
   if (outer_ctx)
     {
@@ -993,6 +998,17 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
   return ctx;
 }
 
+static omp_context *
+clone_omp_context (omp_context *ctx)
+{
+  omp_context *clone_ctx = XCNEW (omp_context);
+
+  memcpy (clone_ctx, ctx, sizeof (omp_context));
+  ctx->next_clone = clone_ctx;
+
+  return clone_ctx;
+}
+
 static gimple_seq maybe_catch_exception (gimple_seq);
 
 /* Finalize task copyfn.  */
@@ -1039,6 +1055,14 @@ delete_omp_context (splay_tree_value value)
 {
   omp_context *ctx = (omp_context *) value;
 
+  /* Delete clones.  */
+  omp_context *clone = ctx->next_clone;
+  while (clone)
+    {
+      clone = clone->next_clone;
+      XDELETE (clone);
+    }
+
   delete ctx->cb.decl_map;
 
   if (ctx->field_map)
@@ -1073,6 +1097,7 @@ delete_omp_context (splay_tree_value value)
   delete ctx->lastprivate_conditional_map;
   delete ctx->allocate_map;
 
+
   XDELETE (ctx);
 }
 
@@ -3008,6 +3033,23 @@ scan_omp_teams (gomp_teams *stmt, omp_context *outer_ctx)
     ctx->record_type = ctx->receiver_decl = NULL;
 }
 
+/* Scan an OpenMP metadirective.  */
+
+static void
+scan_omp_metadirective (gomp_metadirective *stmt, omp_context *outer_ctx)
+{
+  gomp_metadirective_body *body = gimple_omp_metadirective_bodies (stmt);
+
+  while (body)
+    {
+      gimple_seq *body_p = gimple_omp_body_ptr (body);
+      omp_context *ctx = outer_ctx ? clone_omp_context (outer_ctx) : NULL;
+      scan_omp (body_p, ctx);
+
+      body = (gomp_metadirective_body *) body->next;
+    }
+}
+
 /* Check nesting restrictions.  */
 static bool
 check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
@@ -4045,6 +4087,10 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	scan_omp_teams (as_a <gomp_teams *> (stmt), ctx);
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      scan_omp_metadirective (as_a <gomp_metadirective *> (stmt), ctx);
+      break;
+
     case GIMPLE_BIND:
       {
 	tree var;
@@ -10129,6 +10175,22 @@ lower_omp_for_lastprivate (struct omp_for_data *fd, gimple_seq *body_p,
     }
 }
 
+static void
+lower_omp_metadirective (gimple_stmt_iterator *gsi_p, omp_context *ctx)
+{
+  gimple *stmt = gsi_stmt (*gsi_p);
+  gomp_metadirective_body *body = gimple_omp_metadirective_bodies (stmt);
+  while (body)
+    {
+      gimple_seq *body_p = gimple_omp_body_ptr (body);
+      omp_context *ctx = maybe_lookup_ctx (*body_p);
+      lower_omp (body_p, ctx);
+      body = (gomp_metadirective_body *) (body->next);
+    }
+  gsi_insert_after (gsi_p, gimple_build_omp_return (true),
+		    GSI_CONTINUE_LINKING);
+}
+
 /* Callback for walk_gimple_seq.  Find #pragma omp scan statement.  */
 
 static tree
@@ -13474,6 +13536,9 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       else
 	lower_omp_teams (gsi_p, ctx);
       break;
+    case GIMPLE_OMP_METADIRECTIVE:
+      lower_omp_metadirective (gsi_p, ctx);
+      break;
     case GIMPLE_CALL:
       tree fndecl;
       call_stmt = as_a <gcall *> (stmt);
diff --git a/gcc/passes.def b/gcc/passes.def
index 945d2bc797c..9aad498f266 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -186,6 +186,7 @@ along with GCC; see the file COPYING3.  If not see
   NEXT_PASS (pass_oacc_device_lower);
   NEXT_PASS (pass_omp_device_lower);
   NEXT_PASS (pass_omp_target_link);
+  NEXT_PASS (pass_omp_expand_metadirective);
   NEXT_PASS (pass_adjust_alignment);
   NEXT_PASS (pass_all_optimizations);
   PUSH_INSERT_PASSES_WITHIN (pass_all_optimizations)
diff --git a/gcc/tree-cfg.c b/gcc/tree-cfg.c
index 02256580c98..525e945a2d7 100644
--- a/gcc/tree-cfg.c
+++ b/gcc/tree-cfg.c
@@ -1668,6 +1668,23 @@ cleanup_dead_labels (void)
 	  }
 	  break;
 
+	case GIMPLE_OMP_METADIRECTIVE:
+	  {
+	    int i;
+	    for (i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+	      {
+		label = gimple_omp_metadirective_label (stmt, i);
+		new_label = main_block_label (label, label_for_bb);
+		if (new_label != label)
+		  gimple_omp_metadirective_set_label (stmt, i, new_label);
+	      }
+	    label = gimple_omp_metadirective_succ_label (stmt);
+	    new_label = main_block_label (label, label_for_bb);
+	    if (new_label != label)
+	      gimple_omp_metadirective_set_succ_label (stmt, new_label);
+	  }
+	  break;
+
 	default:
 	  break;
       }
@@ -6078,6 +6095,22 @@ gimple_redirect_edge_and_branch (edge e, basic_block dest)
 				           gimple_block_label (dest));
       break;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      {
+	for (unsigned i = 0; i < gimple_omp_metadirective_num_clauses (stmt); i++)
+	  {
+	    tree label = gimple_omp_metadirective_label (stmt, i);
+	    if (label_to_block (cfun, label) == e->dest)
+	      gimple_omp_metadirective_set_label (stmt, i,
+						  gimple_block_label (dest));
+	  }
+	tree label = gimple_omp_metadirective_succ_label (stmt);
+	if (label_to_block (cfun, label) == e->dest)
+	  gimple_omp_metadirective_set_succ_label (stmt,
+						   gimple_block_label (dest));
+      }
+      break;
+
     default:
       /* Otherwise it must be a fallthru edge, and we don't need to
 	 do anything besides redirecting it.  */
diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c
index 8f945b88c12..3a21268b93b 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -4534,6 +4534,12 @@ estimate_num_insns (gimple *stmt, eni_weights *weights)
       return (weights->omp_cost
               + estimate_num_insns_seq (gimple_omp_body (stmt), weights));
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      /* The actual instruction will disappear eventually, so metadirective
+	 statements have zero cost.  */
+      gcc_assert (gimple_omp_body (stmt) == NULL);
+      return 0;
+
     case GIMPLE_TRANSACTION:
       return (weights->tm_cost
 	      + estimate_num_insns_seq (gimple_transaction_body (
diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h
index 15693fee150..c02dda89f6a 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -418,6 +418,7 @@ extern gimple_opt_pass *make_pass_lower_switch_O0 (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_lower_vector (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_lower_vector_ssa (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_omp_oacc_kernels_decompose (gcc::context *ctxt);
+extern gimple_opt_pass *make_pass_omp_expand_metadirective (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_lower_omp (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_diagnose_omp_blocks (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_expand_omp (gcc::context *ctxt);
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 0a575eb9dad..5b601a303c6 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -3626,6 +3626,34 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
       is_expr = false;
       break;
 
+    case OMP_METADIRECTIVE:
+      {
+	pp_string (pp, "#pragma omp metadirective");
+	newline_and_indent (pp, spc + 2);
+	pp_left_brace (pp);
+
+	tree clause = OMP_METADIRECTIVE_CLAUSES (node);
+	while (clause != NULL_TREE)
+	  {
+	    newline_and_indent (pp, spc + 4);
+	    if (TREE_PURPOSE (clause) == NULL_TREE)
+	      pp_string (pp, "default:");
+	    else
+	      {
+		pp_string (pp, "when (");
+		dump_generic_node (pp, TREE_PURPOSE (clause), spc + 4, flags,
+				   false);
+		pp_string (pp, "):");
+	      }
+	    newline_and_indent (pp, spc + 6);
+	    dump_generic_node (pp, TREE_VALUE (clause), spc + 6, flags, false);
+	    clause = TREE_CHAIN (clause);
+	  }
+	newline_and_indent (pp, spc + 2);
+	pp_right_brace (pp);
+      }
+      break;
+
     case TRANSACTION_EXPR:
       if (TRANSACTION_EXPR_OUTER (node))
 	pp_string (pp, "__transaction_atomic [[outer]]");
diff --git a/gcc/tree-ssa-operands.c b/gcc/tree-ssa-operands.c
index c15575416dd..ee26451d717 100644
--- a/gcc/tree-ssa-operands.c
+++ b/gcc/tree-ssa-operands.c
@@ -978,6 +978,9 @@ operands_scanner::parse_ssa_operands ()
       append_vuse (gimple_vop (fn));
       goto do_default;
 
+    case GIMPLE_OMP_METADIRECTIVE:
+      break;
+
     case GIMPLE_CALL:
       /* Add call-clobbered operands, if needed.  */
       maybe_add_call_vops (as_a <gcall *> (stmt));
diff --git a/gcc/tree.def b/gcc/tree.def
index eda050bdc55..9d50c739539 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1264,6 +1264,11 @@ DEFTREECODE (OMP_TARGET_ENTER_DATA, "omp_target_enter_data", tcc_statement, 1)
    Operand 0: OMP_TARGET_EXIT_DATA_CLAUSES: List of clauses.  */
 DEFTREECODE (OMP_TARGET_EXIT_DATA, "omp_target_exit_data", tcc_statement, 1)
 
+/* OpenMP - #pragma omp metadirective [clause1 ... clauseN]
+   Operand 0: OMP_METADIRECTIVE_CLAUSES: List of selectors and directive
+	variants.  */
+DEFTREECODE (OMP_METADIRECTIVE, "omp_metadirective", tcc_statement, 1)
+
 /* OMP_ATOMIC through OMP_ATOMIC_CAPTURE_NEW must be consecutive,
    or OMP_ATOMIC_SEQ_CST needs adjusting.  */
 
diff --git a/gcc/tree.h b/gcc/tree.h
index 64612cfa368..0d74cc75ce1 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1459,6 +1459,9 @@ class auto_suppress_location_wrappers
 #define OMP_TARGET_EXIT_DATA_CLAUSES(NODE)\
   TREE_OPERAND (OMP_TARGET_EXIT_DATA_CHECK (NODE), 0)
 
+#define OMP_METADIRECTIVE_CLAUSES(NODE) \
+  TREE_OPERAND (OMP_METADIRECTIVE_CHECK (NODE), 0)
+
 #define OMP_SCAN_BODY(NODE)	TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0)
 #define OMP_SCAN_CLAUSES(NODE)	TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1)
 

^ permalink raw reply	[flat|nested] 29+ messages in thread

end of thread, other threads:[~2022-05-30 11:52 UTC | newest]

Thread overview: 29+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-07-09 11:16 [WIP, OpenMP] OpenMP metadirectives support Kwok Cheung Yeung
2021-07-26 11:38 ` Kwok Cheung Yeung
2021-07-26 14:29 ` Jakub Jelinek
2021-07-26 19:28   ` Kwok Cheung Yeung
2021-07-26 19:56     ` Jakub Jelinek
2021-07-26 21:19       ` Kwok Cheung Yeung
2021-07-26 21:23         ` Jakub Jelinek
2021-07-26 21:27           ` Kwok Cheung Yeung
2022-01-28 16:33           ` [PATCH] openmp: Add warning when functions containing metadirectives with 'construct={target}' called directly Kwok Cheung Yeung
2021-12-10 17:27   ` [WIP, OpenMP] OpenMP metadirectives support Kwok Cheung Yeung
2021-12-10 17:29 ` [PATCH 0/7] openmp: " Kwok Cheung Yeung
2021-12-10 17:31   ` [PATCH 1/7] openmp: Add C support for parsing metadirectives Kwok Cheung Yeung
2022-02-18 21:09     ` [PATCH] openmp: Improve handling of nested OpenMP metadirectives in C and C++ (was: Re: [PATCH 1/7] openmp: Add C support for parsing metadirectives) Kwok Cheung Yeung
2022-02-18 21:26       ` [og11][committed] openmp: Improve handling of nested OpenMP metadirectives in C and C++ Kwok Cheung Yeung
2022-05-27 17:44     ` [PATCH 1/7] openmp: Add C support for parsing metadirectives Jakub Jelinek
2021-12-10 17:33   ` [PATCH 2/7] openmp: Add middle-end support for metadirectives Kwok Cheung Yeung
2022-05-30 10:54     ` Jakub Jelinek
2021-12-10 17:35   ` [PATCH 3/7] openmp: Add support for resolving metadirectives during parsing and Gimplification Kwok Cheung Yeung
2022-05-30 11:13     ` Jakub Jelinek
2021-12-10 17:36   ` [PATCH 4/7] openmp: Add support for streaming metadirectives and resolving them after LTO Kwok Cheung Yeung
2022-05-30 11:33     ` Jakub Jelinek
2021-12-10 17:37   ` [PATCH 5/7] openmp: Add C++ support for parsing metadirectives Kwok Cheung Yeung
2022-05-30 11:52     ` Jakub Jelinek
2021-12-10 17:39   ` [PATCH 6/7] openmp, fortran: Add Fortran " Kwok Cheung Yeung
2022-02-14 15:09     ` Kwok Cheung Yeung
2022-02-14 15:17     ` Kwok Cheung Yeung
2021-12-10 17:40   ` [PATCH 7/7] openmp: Add testcases for metadirectives Kwok Cheung Yeung
2022-05-27 13:42     ` Jakub Jelinek
2022-01-24 21:28   ` [PATCH] openmp: Metadirective patch fixes Kwok Cheung Yeung

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).