[5/8] OpenMP: C++ front-end support for metadirectives

Message ID 20240106185257.126445-6-sandra@codesourcery.com
State Unresolved
Headers
Series OpenMP: Implement metadirective support |

Checks

Context Check Description
snail/gcc-patch-check warning Git am fail log

Commit Message

Sandra Loosemore Jan. 6, 2024, 6:52 p.m. UTC
  From: Kwok Cheung Yeung <kcy@codesourcery.com>

	This patch adds C++ support for metadirectives.  It uses the
	c-family support committed with the corresponding C front end patch
	to do early parse-time metadirective resolution when possible.

	Additional C/C++ common testcases are provided in a subsequent
	patch in the series.

gcc/cp/ChangeLog
	* parser.cc (cp_parser_skip_to_end_of_block_or_statement): Add
	metadirective_p parameter, use it to control brace/parentheses
	behavior for metadirectives.
	(mangle_metadirective_region_label): New.
	(cp_parser_label_for_labeled_statement): Use it.
	(cp_parser_jump_statement): Likewise.
	(cp_parser_omp_context_selector): Add metadirective_p
	parameter, use it to control error behavior for non-constant exprs
	properties.
	(cp_parser_omp_context_selector_specification): Add metadirective_p
	parameter, use it for cp_parser_omp_context_selector call.
	(cp_finish_omp_declare_variant): Adjust call to
	cp_parser_omp_context_selector_specification.
	(analyze_metadirective_body): New.
	(cp_parser_omp_metadirective): New.
	(cp_parser_pragma): Handle PRAGMA_OMP_METADIRECTIVE.
	* parser.h (struct cp_parser): Add fields for metadirective parsing
	state.
	* pt.cc (tsubst_omp_context_selector): New.
	(tsubst_stmt): Handle OMP_METADIRECTIVE.

gcc/testsuite/ChangeLog
	* g++.dg/gomp/attrs-metadirective-1.C: New.
	* g++.dg/gomp/attrs-metadirective-2.C: New.
	* g++.dg/gomp/attrs-metadirective-3.C: New.
	* g++.dg/gomp/attrs-metadirective-4.C: New.
	* g++.dg/gomp/attrs-metadirective-5.C: New.
	* g++.dg/gomp/attrs-metadirective-6.C: New.
	* g++.dg/gomp/attrs-metadirective-7.C: New.
	* g++.dg/gomp/attrs-metadirective-8.C: New.

libgomp/ChangeLog
	* testsuite/libgomp.c++/metadirective-template-1.C: New.
	* testsuite/libgomp.c++/metadirective-template-2.C: New.
	* testsuite/libgomp.c++/metadirective-template-3.C: New.

Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com>
---
 gcc/cp/parser.cc                              | 524 +++++++++++++++++-
 gcc/cp/parser.h                               |   7 +
 gcc/cp/pt.cc                                  | 118 ++++
 .../g++.dg/gomp/attrs-metadirective-1.C       |  40 ++
 .../g++.dg/gomp/attrs-metadirective-2.C       |  74 +++
 .../g++.dg/gomp/attrs-metadirective-3.C       |  31 ++
 .../g++.dg/gomp/attrs-metadirective-4.C       |  41 ++
 .../g++.dg/gomp/attrs-metadirective-5.C       |  24 +
 .../g++.dg/gomp/attrs-metadirective-6.C       |  31 ++
 .../g++.dg/gomp/attrs-metadirective-7.C       |  31 ++
 .../g++.dg/gomp/attrs-metadirective-8.C       |  16 +
 .../libgomp.c++/metadirective-template-1.C    |  37 ++
 .../libgomp.c++/metadirective-template-2.C    |  41 ++
 .../libgomp.c++/metadirective-template-3.C    |  41 ++
 14 files changed, 1043 insertions(+), 13 deletions(-)
 create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C
 create mode 100644 libgomp/testsuite/libgomp.c++/metadirective-template-1.C
 create mode 100644 libgomp/testsuite/libgomp.c++/metadirective-template-2.C
 create mode 100644 libgomp/testsuite/libgomp.c++/metadirective-template-3.C
  

Patch

diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index b864a9b5275..4ec4ed8a01c 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -2998,7 +2998,7 @@  static void cp_parser_skip_to_end_of_statement
 static void cp_parser_consume_semicolon_at_end_of_statement
   (cp_parser *);
 static void cp_parser_skip_to_end_of_block_or_statement
-  (cp_parser *);
+  (cp_parser *, bool = false);
 static bool cp_parser_skip_to_closing_brace
   (cp_parser *);
 static bool cp_parser_skip_entire_template_parameter_list
@@ -4174,9 +4174,11 @@  cp_parser_consume_semicolon_at_end_of_statement (cp_parser *parser)
    have consumed a non-nested `;'.  */
 
 static void
-cp_parser_skip_to_end_of_block_or_statement (cp_parser* parser)
+cp_parser_skip_to_end_of_block_or_statement (cp_parser* parser,
+					     bool metadirective_p)
 {
   int nesting_depth = 0;
+  int bracket_depth = 0;
 
   /* Unwind generic function template scope if necessary.  */
   if (parser->fully_implicit_function_template_p)
@@ -4198,7 +4200,7 @@  cp_parser_skip_to_end_of_block_or_statement (cp_parser* parser)
 
 	case CPP_SEMICOLON:
 	  /* Stop if this is an unnested ';'. */
-	  if (!nesting_depth)
+	  if (!nesting_depth && (!metadirective_p || bracket_depth <= 0))
 	    nesting_depth = -1;
 	  break;
 
@@ -4217,6 +4219,19 @@  cp_parser_skip_to_end_of_block_or_statement (cp_parser* parser)
 	  nesting_depth++;
 	  break;
 
+	case CPP_OPEN_PAREN:
+	  /* Track parentheses in case the statement is a standalone 'for'
+	     statement - we want to skip over the semicolons separating the
+	     operands.  */
+	  if (metadirective_p && nesting_depth == 0)
+	    bracket_depth++;
+	  break;
+
+	case CPP_CLOSE_PAREN:
+	  if (metadirective_p && nesting_depth == 0)
+	    bracket_depth--;
+	  break;
+
 	case CPP_KEYWORD:
 	  if (!cp_token_is_module_directive (token))
 	    break;
@@ -12834,6 +12849,18 @@  attr_chainon (tree attrs, tree attr)
   return chainon (attrs, attr);
 }
 
+
+/* Helper function for cp_parser_label: mangle a metadirective region
+   label NAME.  */
+static tree
+mangle_metadirective_region_label (cp_parser *parser, tree name)
+{
+  const char *old_name = IDENTIFIER_POINTER (name);
+  char *new_name = (char *) alloca (strlen (old_name) + 32);
+  sprintf (new_name, "%s_MDR%u", old_name, parser->metadirective_region_num);
+  return get_identifier (new_name);
+}
+
 /* Parse the label for a labeled-statement, i.e.
 
    label:
@@ -12936,7 +12963,12 @@  cp_parser_label_for_labeled_statement (cp_parser* parser, tree attributes)
 
     default:
       /* Anything else must be an ordinary label.  */
-      label = finish_label_stmt (cp_parser_identifier (parser));
+      cp_expr identifier = cp_parser_identifier (parser);
+      if (identifier != error_mark_node
+	  && parser->in_metadirective_body
+	  && parser->metadirective_body_labels->contains (*identifier))
+	*identifier = mangle_metadirective_region_label (parser, *identifier);
+      label = finish_label_stmt (identifier);
       if (label && TREE_CODE (label) == LABEL_DECL)
 	FALLTHROUGH_LABEL_P (label) = fallthrough_p;
       break;
@@ -14741,7 +14773,15 @@  cp_parser_jump_statement (cp_parser* parser)
 	  finish_goto_stmt (cp_parser_expression (parser));
 	}
       else
-	finish_goto_stmt (cp_parser_identifier (parser));
+	{
+	  cp_expr identifier = cp_parser_identifier (parser);
+	  if (identifier != error_mark_node
+	      && parser->in_metadirective_body
+	      && parser->metadirective_body_labels->contains (*identifier))
+	    *identifier = mangle_metadirective_region_label (parser,
+							     *identifier);
+	  finish_goto_stmt (identifier);
+	}
       /* Look for the final `;'.  */
       cp_parser_require (parser, CPP_SEMICOLON, RT_SEMICOLON);
       break;
@@ -47431,7 +47471,7 @@  cp_parser_omp_declare_simd (cp_parser *parser, cp_token *pragma_tok,
 
 static tree
 cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set,
-				bool has_parms_p)
+				bool has_parms_p, bool metadirective_p)
 {
   tree ret = NULL_TREE;
   do
@@ -47591,17 +47631,25 @@  cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set,
 	      while (1);
 	      break;
 	    case OMP_TRAIT_PROPERTY_EXPR:
-	      /* FIXME: this is bogus, the expression need
-		 not be constant.  */
-	      t = cp_parser_constant_expression (parser);
+	      /* FIXME: I believe it is an unimplemented feature rather
+		 than a user error to have non-constant expressions
+		 inside "declare variant".  */
+	      t = metadirective_p
+		? cp_parser_expression (parser)
+		: cp_parser_constant_expression (parser);
 	      if (t != error_mark_node)
 		{
 		  t = fold_non_dependent_expr (t);
-		  if (!value_dependent_expression_p (t)
+		  if (!metadirective_p
+		      && !value_dependent_expression_p (t)
 		      && (!INTEGRAL_TYPE_P (TREE_TYPE (t))
 			  || !tree_fits_shwi_p (t)))
 		    error_at (token->location, "property must be "
 			      "constant integer expression");
+		  if (metadirective_p
+		      && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+		    error_at (token->location,
+			      "property must be integer expression");
 		  else
 		    properties = make_trait_property (NULL_TREE, t,
 						      properties);
@@ -47680,7 +47728,8 @@  cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set,
 
 static tree
 cp_parser_omp_context_selector_specification (cp_parser *parser,
-					      bool has_parms_p)
+					      bool has_parms_p,
+					      bool metadirective_p)
 {
   tree ret = NULL_TREE;
   do
@@ -47707,7 +47756,8 @@  cp_parser_omp_context_selector_specification (cp_parser *parser,
 	return error_mark_node;
 
       tree selectors
-	= cp_parser_omp_context_selector (parser, set, has_parms_p);
+	= cp_parser_omp_context_selector (parser, set, has_parms_p,
+					  metadirective_p);
       if (selectors == error_mark_node)
 	{
 	  cp_parser_skip_to_closing_brace (parser);
@@ -48017,7 +48067,8 @@  cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok,
   if (!parens.require_open (parser))
     goto fail;
 
-  tree ctx = cp_parser_omp_context_selector_specification (parser, true);
+  tree ctx = cp_parser_omp_context_selector_specification (parser, true,
+							   false);
   if (ctx == error_mark_node)
     goto fail;
   ctx = omp_check_context_selector (match_loc, ctx, false);
@@ -48650,6 +48701,449 @@  cp_parser_omp_end (cp_parser *parser, cp_token *pragma_tok)
     }
 }
 
+
+/* Helper function for cp_parser_omp_metadirective.  */
+
+static void
+analyze_metadirective_body (cp_parser *parser,
+			    vec<cp_token> &tokens,
+			    vec<tree> &labels)
+{
+  int nesting_depth = 0;
+  int bracket_depth = 0;
+  bool in_case = false;
+  bool in_label_decl = false;
+  cp_token *pragma_tok = NULL;
+
+  while (1)
+    {
+      cp_token *token = cp_lexer_peek_token (parser->lexer);
+      bool stop = false;
+
+      if (cp_lexer_next_token_is_keyword (parser->lexer, RID_CASE))
+	in_case = true;
+      else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_LABEL))
+	in_label_decl = true;
+
+      switch (token->type)
+	{
+	case CPP_EOF:
+	  break;
+	case CPP_NAME:
+	  if ((!in_case
+	       && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+	      || in_label_decl)
+	    labels.safe_push (token->u.value);
+	  goto add;
+	case CPP_OPEN_BRACE:
+	  ++nesting_depth;
+	  goto add;
+	case CPP_CLOSE_BRACE:
+	  if (--nesting_depth == 0 && bracket_depth == 0)
+	    stop = true;
+	  goto add;
+	case CPP_OPEN_PAREN:
+	  ++bracket_depth;
+	  goto add;
+	case CPP_CLOSE_PAREN:
+	  --bracket_depth;
+	  goto add;
+	case CPP_COLON:
+	  in_case = false;
+	  goto add;
+	case CPP_SEMICOLON:
+	  if (nesting_depth == 0 && bracket_depth == 0)
+	    stop = true;
+	  /* Local label declarations are terminated by a semicolon.  */
+	  in_label_decl = false;
+	  goto add;
+	case CPP_PRAGMA:
+	  parser->lexer->in_pragma = true;
+	  pragma_tok = token;
+	  goto add;
+	case CPP_PRAGMA_EOL:
+	  /* C++ attribute syntax for OMP directives lexes as a pragma,
+	     but we must reset the associated lexer state when we reach
+	     the end in order to get the tokens for the statement that
+	     come after it.  */
+	  tokens.safe_push (*token);
+	  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+	  pragma_tok = NULL;
+	  continue;
+	default:
+	add:
+	  tokens.safe_push (*token);
+	  cp_lexer_consume_token (parser->lexer);
+	  if (stop)
+	    break;
+	  continue;
+	}
+      break;
+    }
+}
+
+/* OpenMP 5.0:
+
+  # pragma omp metadirective [clause[, clause]]
+*/
+
+static void
+cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok,
+			     bool *if_p)
+{
+  static unsigned int metadirective_region_count = 0;
+
+  auto_vec<cp_token> directive_tokens;
+  auto_vec<cp_token> body_tokens;
+  auto_vec<tree> body_labels;
+  auto_vec<const struct c_omp_directive *> directives;
+  auto_vec<tree> ctxs;
+  bool default_seen = false;
+  int directive_token_idx = 0;
+  location_t pragma_loc = pragma_tok->location;
+  tree standalone_body = NULL_TREE;
+  vec<struct omp_metadirective_variant> candidates;
+  bool requires_body = false;
+
+  tree ret = make_node (OMP_METADIRECTIVE);
+  SET_EXPR_LOCATION (ret, pragma_loc);
+  TREE_TYPE (ret) = void_type_node;
+  OMP_METADIRECTIVE_VARIANTS (ret) = NULL_TREE;
+
+  while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+    {
+      if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+	cp_lexer_consume_token (parser->lexer);
+      if (cp_lexer_next_token_is_not (parser->lexer, CPP_NAME)
+	  && cp_lexer_next_token_is_not (parser->lexer, CPP_KEYWORD))
+	{
+	  cp_parser_error (parser, "expected %<when%>, "
+			   "%<otherwise%>, or %<default%> clause");
+	  goto fail;
+	}
+
+      location_t match_loc = cp_lexer_peek_token (parser->lexer)->location;
+      const char *p
+	= IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value);
+      cp_lexer_consume_token (parser->lexer);
+      bool default_p
+	= strcmp (p, "default") == 0 || strcmp (p, "otherwise") == 0;
+      if (default_p)
+	{
+	  if (default_seen)
+	    {
+	      error_at (match_loc, "too many %<otherwise%> or %<default%> "
+			"clauses in %<metadirective%>");
+	      cp_parser_skip_to_end_of_block_or_statement (parser, true);
+	      goto fail;
+	    }
+	  else
+	    default_seen = true;
+	}
+      if (!strcmp (p, "when") == 0 && !default_p)
+	{
+	  error_at (match_loc, "%qs is not valid for %qs",
+		    p, "metadirective");
+	  cp_parser_skip_to_end_of_block_or_statement (parser, true);
+	  goto fail;
+	}
+
+      matching_parens parens;
+      tree ctx = NULL_TREE;
+      bool skip = false;
+
+      if (!parens.require_open (parser))
+	goto fail;
+
+      if (!default_p)
+	{
+	  ctx = cp_parser_omp_context_selector_specification (parser, false,
+							      true);
+	  if (ctx == error_mark_node)
+	    goto fail;
+	  ctx = omp_check_context_selector (match_loc, ctx, true);
+	  if (ctx == error_mark_node)
+	    goto fail;
+
+	  /* Remove the selector from further consideration if it can be
+	     evaluated as a non-match at this point.  */
+	  /* FIXME: we could still do this if the context selector
+	     doesn't have any dependent subexpressions.  */
+	  skip = (!processing_template_decl
+		  && omp_context_selector_matches (ctx, true, true) == 0);
+
+	  if (cp_lexer_next_token_is_not (parser->lexer, CPP_COLON))
+	    {
+	      cp_parser_require (parser, CPP_COLON, RT_COLON);
+	      goto fail;
+	    }
+	  cp_lexer_consume_token (parser->lexer);
+	}
+
+      /* Read in the directive type and create a dummy pragma token for
+	 it.  */
+      location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+
+      const char *directive[3] = {};
+      int i;
+      for (i = 0; i < 3; i++)
+	{
+	  tree id;
+	  if (cp_lexer_peek_nth_token (parser->lexer, i + 1)->type
+	      == CPP_CLOSE_PAREN)
+	    {
+	      if (i == 0)
+		directive[i++] = "nothing";
+	      break;
+	    }
+	  else if (cp_lexer_peek_nth_token (parser->lexer, i + 1)->type
+		   == CPP_NAME)
+	    id = cp_lexer_peek_nth_token (parser->lexer, i + 1)->u.value;
+	  else if (cp_lexer_peek_nth_token (parser->lexer, i + 1)->keyword
+		   != RID_MAX)
+	    {
+	      enum rid rid
+		= cp_lexer_peek_nth_token (parser->lexer, i + 1)->keyword;
+	      id = ridpointers[rid];
+	    }
+	  else
+	    break;
+
+	  directive[i] = IDENTIFIER_POINTER (id);
+	}
+      if (i == 0)
+	{
+	  error_at (loc, "expected directive name");
+	  cp_parser_skip_to_end_of_block_or_statement (parser, true);
+	  goto fail;
+	}
+
+      const struct c_omp_directive *omp_directive
+	= c_omp_categorize_directive (directive[0],
+				      directive[1],
+				      directive[2]);
+
+      if (omp_directive == NULL)
+	{
+	  for (int j = 0; j < i; j++)
+	    cp_lexer_consume_token (parser->lexer);
+	  cp_parser_error (parser, "unknown directive name");
+	  goto fail;
+	}
+      else
+	{
+	  int token_count = 0;
+	  if (omp_directive->first) token_count++;
+	  if (omp_directive->second) token_count++;
+	  if (omp_directive->third) token_count++;
+	  for (int j = 0; j < token_count; j++)
+	    cp_lexer_consume_token (parser->lexer);
+	}
+      if (p == NULL)
+	{
+	  cp_parser_error (parser, "expected directive name");
+	  goto fail;
+	}
+      if (omp_directive->id == PRAGMA_OMP_METADIRECTIVE)
+	{
+	  cp_parser_error (parser,
+			   "metadirectives cannot be used as variants of a "
+			   "%<metadirective%>");
+	  goto fail;
+	}
+      if (omp_directive->kind == C_OMP_DIR_DECLARATIVE)
+	{
+	  sorry_at (loc, "declarative directive variants of a "
+			 "%<metadirective%> are not supported");
+	  goto fail;
+	}
+      if (omp_directive->kind == C_OMP_DIR_CONSTRUCT)
+	requires_body = true;
+
+      if (!skip)
+	{
+	  cp_token pragma_token;
+	  pragma_token.type = CPP_PRAGMA;
+	  pragma_token.location = loc;
+	  pragma_token.u.value = build_int_cst (NULL, omp_directive->id);
+
+	  directives.safe_push (omp_directive);
+	  directive_tokens.safe_push (pragma_token);
+	  ctxs.safe_push (ctx);
+	}
+
+      /* Read in tokens for the directive clauses.  */
+      int nesting_depth = 0;
+      while (1)
+	{
+	  cp_token *token = cp_lexer_peek_token (parser->lexer);
+	  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)
+		directive_tokens.safe_push (*token);
+	      cp_lexer_consume_token (parser->lexer);
+	      continue;
+	    }
+	  break;
+	}
+
+      cp_lexer_consume_token (parser->lexer);
+
+      if (!skip)
+	{
+	  cp_token eol_token = {};
+	  eol_token.type = CPP_PRAGMA_EOL;
+	  eol_token.keyword = RID_MAX;
+	  directive_tokens.safe_push (eol_token);
+	}
+    }
+  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+
+  if (!default_seen)
+    {
+      /* Add a default clause that evaluates to 'omp nothing'.  */
+      const struct c_omp_directive *omp_directive
+	= c_omp_categorize_directive ("nothing", NULL, NULL);
+
+      cp_token pragma_token = {};
+      pragma_token.type = CPP_PRAGMA;
+      pragma_token.keyword = RID_MAX;
+      pragma_token.location = UNKNOWN_LOCATION;
+      pragma_token.u.value = build_int_cst (NULL, PRAGMA_OMP_NOTHING);
+
+      directives.safe_push (omp_directive);
+      directive_tokens.safe_push (pragma_token);
+      ctxs.safe_push (NULL_TREE);
+
+      cp_token eol_token = {};
+      eol_token.type = CPP_PRAGMA_EOL;
+      eol_token.keyword = RID_MAX;
+      directive_tokens.safe_push (eol_token);
+    }
+
+  if (requires_body)
+    analyze_metadirective_body (parser, body_tokens, body_labels);
+
+  /* Process each candidate directive.  */
+  unsigned i;
+  tree ctx;
+  cp_lexer *lexer;
+
+  lexer = cp_lexer_alloc ();
+  lexer->debugging_p = parser->lexer->debugging_p;
+  vec_safe_reserve (lexer->buffer,
+		    directive_tokens.length () + body_tokens.length () + 2);
+
+  FOR_EACH_VEC_ELT (ctxs, i, ctx)
+    {
+      lexer->buffer->truncate (0);
+
+      /* Add the directive tokens.  */
+      do
+	lexer->buffer->quick_push (directive_tokens [directive_token_idx++]);
+      while (lexer->buffer->last ().type != CPP_PRAGMA_EOL);
+
+      /* Add the body tokens.  */
+      gcc_assert (requires_body || body_tokens.is_empty ());
+      for (unsigned j = 0; j < body_tokens.length (); j++)
+	lexer->buffer->quick_push (body_tokens[j]);
+
+      /* Make sure nothing tries to read past the end of the tokens.  */
+      cp_token eof_token = {};
+      eof_token.type = CPP_EOF;
+      eof_token.keyword = RID_MAX;
+      lexer->buffer->quick_push (eof_token);
+      lexer->buffer->quick_push (eof_token);
+
+      lexer->next_token = lexer->buffer->address();
+      lexer->last_token = lexer->next_token + lexer->buffer->length () - 1;
+
+      cp_lexer *old_lexer = parser->lexer;
+      bool old_in_metadirective_body = parser->in_metadirective_body;
+      vec<tree> *old_metadirective_body_labels
+	= parser->metadirective_body_labels;
+      unsigned int old_metadirective_region_num
+	= parser->metadirective_region_num;
+      parser->lexer = lexer;
+      cp_lexer_set_source_position_from_token (lexer->next_token);
+      parser->in_metadirective_body = true;
+      parser->metadirective_body_labels = &body_labels;
+      parser->metadirective_region_num = ++metadirective_region_count;
+
+      int prev_errorcount = errorcount;
+      tree directive = push_stmt_list ();
+      tree directive_stmt = begin_compound_stmt (0);
+
+      cp_parser_pragma (parser, pragma_compound, if_p);
+      finish_compound_stmt (directive_stmt);
+      directive = pop_stmt_list (directive);
+
+      bool standalone_p
+	= directives[i]->kind == C_OMP_DIR_STANDALONE
+	  || directives[i]->kind == C_OMP_DIR_UTILITY;
+      if (standalone_p && requires_body)
+	{
+	  /* Parsing standalone directives will not consume the body
+	     tokens, so do that here.  */
+	  if (standalone_body == NULL_TREE)
+	    {
+	      standalone_body = push_stmt_list ();
+	      cp_parser_statement (parser, NULL_TREE, false, if_p);
+	      standalone_body = pop_stmt_list (standalone_body);
+	    }
+	  else
+	    cp_parser_skip_to_end_of_block_or_statement (parser, true);
+	}
+
+      tree body = standalone_p ? standalone_body : NULL_TREE;
+      tree variant = make_omp_metadirective_variant (ctx, directive, body);
+      OMP_METADIRECTIVE_VARIANTS (ret)
+	= chainon (OMP_METADIRECTIVE_VARIANTS (ret), variant);
+
+      /* Check that all valid tokens have been consumed if no parse errors
+	 encountered.  */
+      gcc_assert (errorcount != prev_errorcount
+		  || cp_lexer_next_token_is (parser->lexer, CPP_EOF));
+
+      parser->lexer = old_lexer;
+      cp_lexer_set_source_position_from_token (old_lexer->next_token);
+      parser->in_metadirective_body = old_in_metadirective_body;
+      parser->metadirective_body_labels = old_metadirective_body_labels;
+      parser->metadirective_region_num = old_metadirective_region_num;
+    }
+
+  /* Try to resolve the metadirective early.  */
+  if (!processing_template_decl)
+    {
+      candidates = omp_resolve_metadirective (ret);
+      if (!candidates.is_empty ())
+	ret = c_omp_expand_metadirective (candidates);
+    }
+
+  add_stmt (ret);
+  return;
+
+fail:
+  /* Skip the metadirective pragma.  */
+  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+
+  /* Skip the metadirective body.  */
+  cp_parser_skip_to_end_of_block_or_statement (parser, true);
+}
+
+
 /* 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
@@ -50582,6 +51076,10 @@  cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p)
       cp_parser_omp_nothing (parser, pragma_tok);
       return false;
 
+    case PRAGMA_OMP_METADIRECTIVE:
+      cp_parser_omp_metadirective (parser, pragma_tok, if_p);
+      return true;
+
     case PRAGMA_OMP_ERROR:
       return cp_parser_omp_error (parser, pragma_tok, context);
 
diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h
index 8a40d3870ca..1e4ea425b72 100644
--- a/gcc/cp/parser.h
+++ b/gcc/cp/parser.h
@@ -443,6 +443,13 @@  struct GTY(()) cp_parser {
   /* Pointer to state for parsing omp_loops.  Managed by
      cp_parser_omp_for_loop in parser.cc and not used outside that file.  */
   struct omp_for_parse_data * GTY((skip)) omp_for_parse_state;
+
+  /* Set if we are processing a statement body associated with a
+     metadirective variant.  */
+  bool in_metadirective_body;
+
+  vec<tree> * GTY((skip)) metadirective_body_labels;
+  unsigned metadirective_region_num;
 };
 
 /* In parser.cc  */
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index e38e7a773f0..e7a0186fb03 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -17694,6 +17694,78 @@  tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
   return new_clauses;
 }
 
+/* Like tsubst_copy, but specifically for OpenMP context selectors.  */
+static tree
+tsubst_omp_context_selector (tree ctx, tree args, tsubst_flags_t complain,
+			     tree in_decl)
+{
+  tree new_ctx = NULL_TREE;
+  for (tree set = ctx; set; set = TREE_CHAIN (set))
+    {
+      tree selectors = NULL_TREE;
+      for (tree sel = OMP_TSS_TRAIT_SELECTORS (set); sel;
+	   sel = TREE_CHAIN (sel))
+	{
+	  enum omp_ts_code code = OMP_TS_CODE (sel);
+	  tree properties = NULL_TREE;
+	  tree score = OMP_TS_SCORE (sel);
+	  tree t;
+
+	  if (score)
+	    {
+	      score = tsubst_expr (score, args, complain, in_decl);
+	      score = fold_non_dependent_expr (score);
+	      if (!INTEGRAL_TYPE_P (TREE_TYPE (score))
+		  || TREE_CODE (score) != INTEGER_CST)
+		{
+		  error_at (cp_expr_loc_or_input_loc (score),
+			    "%<score%> argument must "
+			    "be constant integer expression");
+		  score = NULL_TREE;
+		}
+	      else if (tree_int_cst_sgn (score) < 0)
+		{
+		  error_at (cp_expr_loc_or_input_loc (score),
+			    "%<score%> argument must be non-negative");
+		  score = NULL_TREE;
+		}
+	    }
+
+	  switch (omp_ts_map[OMP_TS_CODE (sel)].tp_type)
+	      {
+	      case OMP_TRAIT_PROPERTY_EXPR:
+		t = tsubst_expr (OMP_TP_VALUE (OMP_TS_PROPERTIES (sel)),
+				 args, complain, in_decl);
+		t = fold_non_dependent_expr (t);
+		if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+		  error_at (cp_expr_loc_or_input_loc (t),
+			    "property must be integer expression");
+		properties = make_trait_property (NULL_TREE, t, NULL_TREE);
+		break;
+	      case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
+		if (OMP_TS_CODE (sel) == OMP_TRAIT_CONSTRUCT_SIMD)
+		  properties = tsubst_omp_clauses (OMP_TS_PROPERTIES (sel),
+						   C_ORT_OMP_DECLARE_SIMD,
+						   args, complain, in_decl);
+		break;
+	      default:
+		/* Nothing to do here, just copy.  */
+		for (tree prop = OMP_TS_PROPERTIES (sel);
+		     prop; prop = TREE_CHAIN (prop))
+		  properties = make_trait_property (OMP_TP_NAME (prop),
+						    OMP_TP_VALUE (prop),
+						    properties);
+	      }
+	  selectors = make_trait_selector (code, score, properties, selectors);
+	}
+      new_ctx = make_trait_set_selector (OMP_TSS_CODE (set),
+					 nreverse (selectors),
+					 new_ctx);
+    }
+  return nreverse (new_ctx);
+}
+
+
 /* Like tsubst_expr, but unshare TREE_LIST nodes.  */
 
 static tree
@@ -19243,6 +19315,52 @@  tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl)
 	}
       break;
 
+    case OMP_METADIRECTIVE:
+      {
+	tree variants = NULL_TREE;
+	for (tree v = OMP_METADIRECTIVE_VARIANTS (t); v; v = TREE_CHAIN (v))
+	  {
+	    tree ctx = OMP_METADIRECTIVE_VARIANT_SELECTOR (v);
+	    tree directive = OMP_METADIRECTIVE_VARIANT_DIRECTIVE (v);
+	    tree body = OMP_METADIRECTIVE_VARIANT_BODY (v);
+	    tree s;
+
+	    /* CTX is null if this is the default variant.  */
+	    if (ctx)
+	      {
+		ctx = tsubst_omp_context_selector (ctx, args, complain,
+						   in_decl);
+		/* Remove the selector from further consideration if it can be
+		   evaluated as a non-match at this point.  */
+		if (omp_context_selector_matches (ctx, true, true) == 0)
+		  continue;
+	      }
+	    s = push_stmt_list ();
+	    RECUR (directive);
+	    directive = pop_stmt_list (s);
+	    if (body)
+	      {
+		s = push_stmt_list ();
+		RECUR (body);
+		body = pop_stmt_list (s);
+	      }
+	    variants
+	      = chainon (variants,
+			 make_omp_metadirective_variant (ctx, directive,
+							 body));
+	  }
+	t = copy_node (t);
+	OMP_METADIRECTIVE_VARIANTS (t) = variants;
+
+	/* Try to resolve the metadirective early.  */
+	vec<struct omp_metadirective_variant> candidates
+	  = omp_resolve_metadirective (t);
+	if (!candidates.is_empty ())
+	  t = c_omp_expand_metadirective (candidates);
+	add_stmt (t);
+	break;
+      }
+
     case TRANSACTION_EXPR:
       {
 	int flags = 0;
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C
new file mode 100644
index 00000000000..a0b09088d3a
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C
@@ -0,0 +1,40 @@ 
+// { dg-do compile { target c++11 } }
+
+#define N 100
+
+void f (int a[], int b[], int c[])
+{
+  int i;
+
+  [[omp::directive (metadirective
+      default (teams loop)
+      default (parallel loop))]] /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */
+    for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+  [[omp::directive (metadirective
+      default (bad_directive))]] /* { dg-error "unknown directive name before '\\)' token" } */
+    for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+  [[omp::directive (metadirective
+      default (teams loop)
+		    where (device={arch("nvptx")}: parallel loop))]] /* { dg-error "'where' is not valid for 'metadirective'" } */
+    for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+  [[omp::directive (metadirective
+      default (teams loop)
+      when (device={arch("nvptx")} parallel loop))]] /* { dg-error "expected ':' before 'parallel'" } */
+    for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+  [[omp::directive (metadirective
+      default (metadirective default (flush)))]]	/* { dg-error "metadirectives cannot be used as variants of a 'metadirective' before 'default'" } */
+    for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+  /* Test improperly nested metadirectives - even though the second
+     metadirective resolves to 'omp nothing', that is not the same as there
+     being literally nothing there.  */
+  [[omp::directive (metadirective
+      when (implementation={vendor("gnu")}: parallel for))]]
+  [[omp::directive (metadirective      /* { dg-error "'#pragma' is not allowed here" } */
+      when (implementation={vendor("cray")}: parallel for))]]
+      for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+}
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C
new file mode 100644
index 00000000000..44c87df1776
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C
@@ -0,0 +1,74 @@ 
+// { dg-do compile { target c++11 } }
+
+#define N 100
+
+int main (void)
+{
+  int x = 0;
+  int y = 0;
+
+  /* Test implicit default (nothing).  */
+  [[omp::directive (metadirective,
+      when (device={arch("nvptx")}: barrier))]]
+    x = 1;
+
+  /* Test with multiple standalone directives.  */
+  [[omp::directive (metadirective,
+      when (device={arch("nvptx")}: barrier),
+      default (flush))]]
+    x = 1;
+
+  /* Test combining a standalone directive with one that takes a statement
+     body.  */
+  [[omp::directive (metadirective,
+      when (device={arch("nvptx")}: parallel),
+      default (barrier))]]
+    x = 1;
+
+  /* Test combining a standalone directive with one that takes a for loop.  */
+  [[omp::directive (metadirective,
+      when (device={arch("nvptx")}: parallel for),
+      default (barrier))]]
+    for (int i = 0; i < N; i++)
+      x += i;
+
+  /* Test combining a directive that takes a for loop with one that takes
+     a regular statement body.  */
+  [[omp::directive (metadirective,
+      when (device={arch("nvptx")}: parallel for),
+      default (parallel))]]
+    for (int i = 0; i < N; i++)
+      x += i;
+
+  /* Test labels inside statement body.  */
+  [[omp::directive (metadirective,
+      when (device={arch("nvptx")}: teams num_teams(512)),
+      when (device={arch("gcn")}: teams num_teams(256)),
+    default (teams num_teams(4)))]]
+  {
+    if (x)
+      goto l1;
+    else
+      goto l2;
+  l1: ;
+  l2: ;
+  }
+
+  /* Test local labels inside statement body.  */
+  [[omp::directive (metadirective,
+      when (device={arch("nvptx")}: teams num_teams(512)),
+      when (device={arch("gcn")}: teams num_teams(256)),
+      default (teams num_teams(4)))]]
+  {
+    //__label__ l1, l2;
+
+    if (x)
+      goto l1;
+    else
+      goto l2;
+  l1: ;
+  l2: ;
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C
new file mode 100644
index 00000000000..0c2bbdd2f10
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C
@@ -0,0 +1,31 @@ 
+// { dg-do compile { target c++11 } }
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+/* { dg-additional-options "-fdump-tree-optimized" } */
+
+#define N 100
+
+void f (int x[], int y[], int z[])
+{
+  int i;
+
+  [[omp::sequence (directive (target map(to: x, y) map(from: z)),
+		   directive (metadirective
+			      when (device={arch("nvptx")}: teams loop)
+			      default (parallel loop)))]]
+   for (i = 0; i < N; i++)
+     z[i] = x[i] * y[i];
+}
+
+/* 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/g++.dg/gomp/attrs-metadirective-4.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C
new file mode 100644
index 00000000000..43b939be43b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C
@@ -0,0 +1,41 @@ 
+// { dg-do compile { target c++11 } }
+
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 100
+
+#pragma omp declare target
+void f(double a[], double x) {
+  int i;
+
+  [[omp::directive (metadirective
+	when (construct={target}: distribute parallel for)
+	default (parallel for simd))]]
+   for (i = 0; i < N; i++)
+     a[i] = x * i;
+}
+#pragma omp end declare target
+
+ int main()
+{
+  double a[N];
+
+  #pragma omp target teams map(from: a[0:N])
+    f (a, 3.14159);
+
+  /* TODO: This does not execute a version of f with the default clause
+     active as might be expected.  */
+  f (a, 2.71828); /* { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } */
+
+  return 0;
+ }
+
+ /* 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/g++.dg/gomp/attrs-metadirective-5.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C
new file mode 100644
index 00000000000..1a9cee15be3
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C
@@ -0,0 +1,24 @@ 
+// { dg-do compile { target c++11 } }
+/* { dg-additional-options "-fdump-tree-original" } */
+
+#define N 100
+
+void f (int a[], int flag)
+{
+  int i;
+  [[omp::directive (metadirective
+	when (user={condition(flag)}:
+		target teams distribute parallel for map(from: a[0:N]))
+	default (parallel for))]]
+  for (i = 0; i < N; i++)
+    a[i] = i;
+}
+
+/* The metadirective should be resolved at parse time.  */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times  "#pragma omp distribute" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp for" 2 "original" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C
new file mode 100644
index 00000000000..8a104ff2ebe
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C
@@ -0,0 +1,31 @@ 
+// { dg-do compile { target c++11 } }
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 100
+
+void bar (int a[], int run_parallel, int run_guided)
+{
+  [[omp::directive (metadirective
+		    when (user={condition(run_parallel)}: parallel))]]
+  {
+    int i;
+  [[omp::directive (metadirective
+	when (construct={parallel}, user={condition(run_guided)}:
+	      for schedule(guided))
+	when (construct={parallel}: for schedule(static)))]]
+      for (i = 0; i < N; i++)
+	a[i] = i;
+   }
+ }
+
+/* The outer metadirective should be resolved at parse time.  */
+/* The inner metadirective should be resolved during Gimplificiation.  */
+
+/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp for" 4 "original" } } */
+/* { dg-final { scan-tree-dump-times "when \\(construct = .parallel" 4 "original" } } */
+/* { dg-final { scan-tree-dump-times "default:" 2 "original" } } */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C
new file mode 100644
index 00000000000..161bbf86bd5
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C
@@ -0,0 +1,31 @@ 
+// { dg-do compile { target c++11 } }
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 256
+
+void f (int a[], int num)
+{
+  int i;
+
+  [[omp::directive (metadirective
+      when (target_device={device_num(num), kind("gpu"), arch("nvptx")}:
+	    target parallel for map(tofrom: a[0:N]))
+      when (target_device={device_num(num), kind("gpu"),
+			   arch("amdgcn"), isa("gfx906")}:
+	    target parallel for)
+      when (target_device={device_num(num), kind("cpu"), arch("x86_64")}:
+	    parallel for))]]
+    for (i = 0; i < N; i++)
+      a[i] += i;
+
+  [[omp::directive (metadirective
+      when (target_device={kind("gpu"), arch("nvptx")}:
+	    target parallel for map(tofrom: a[0:N])))]]
+    for (i = 0; i < N; i++)
+      a[i] += i;
+}
+
+/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"gpu\"\\\[0\\\], &\"amdgcn\"\\\[0\\\], &\"gfx906\"\\\[0\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"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/g++.dg/gomp/attrs-metadirective-8.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C
new file mode 100644
index 00000000000..dcb3c365b80
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C
@@ -0,0 +1,16 @@ 
+// { dg-do compile { target c++11 } }
+
+#define N 256
+
+void f ()
+{
+  int i;
+  int a[N];
+
+  [[omp::directive (metadirective
+      when( device={kind(nohost)}: nothing )
+      when( device={arch("nvptx")}: nothing)
+      default( parallel for))]]
+    for (i = 0; i < N; i++)
+      a[i] = i;
+}
diff --git a/libgomp/testsuite/libgomp.c++/metadirective-template-1.C b/libgomp/testsuite/libgomp.c++/metadirective-template-1.C
new file mode 100644
index 00000000000..2d826574861
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/metadirective-template-1.C
@@ -0,0 +1,37 @@ 
+# include <stdio.h>
+
+template <bool tasking> int
+fib (int n)
+{
+  int i, j;
+  if (n < 2)
+    return n;
+  else if ( tasking && n < 8 )  // serial/taskless cutoff for n<8
+    return fib<false> (n);
+  else
+    {
+#pragma omp metadirective				\
+  when (user = {condition (tasking)}: task shared(i))
+      i = fib<tasking> (n - 1);
+#pragma omp metadirective					\
+  when (user = {condition (score(10): tasking)}: task shared(j))
+      j = fib<tasking> (n - 2);
+#pragma omp metadirective			\
+  when (user = {condition (tasking)}: taskwait)
+      return i + j;
+    }
+}
+
+int main ()
+{
+  int n = 15, o = 610;
+#pragma omp parallel
+#pragma omp single
+  {
+    if (fib<true> (n) != o)
+      __builtin_abort ();
+    if (fib<false> (n) != o)
+      __builtin_abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/metadirective-template-2.C b/libgomp/testsuite/libgomp.c++/metadirective-template-2.C
new file mode 100644
index 00000000000..9134fefc7ac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/metadirective-template-2.C
@@ -0,0 +1,41 @@ 
+# include <stdio.h>
+
+template <bool tasking> int
+fib (int n, bool flag)
+{
+  int i, j;
+  if (n < 2)
+    return n;
+  else if ( tasking && flag && n < 8 )  // serial/taskless cutoff for n<8
+    return fib<false> (n, false);
+  else
+    {
+#pragma omp metadirective				\
+  when (user = {condition (tasking && flag)}: task shared(i))
+      i = fib<tasking> (n - 1, flag);
+#pragma omp metadirective					\
+  when (user = {condition (score(10): tasking && flag)}: task shared(j))
+      j = fib<tasking> (n - 2, flag);
+#pragma omp metadirective			\
+  when (user = {condition (tasking && flag)}: taskwait)
+      return i + j;
+    }
+}
+
+int main ()
+{
+  int n = 15, o = 610;
+#pragma omp parallel
+#pragma omp single
+  {
+    if (fib<true> (n, true) != o)
+      __builtin_abort ();
+    if (fib<true> (n, false) != o)
+      __builtin_abort ();
+    if (fib<false> (n, false) != o)
+      __builtin_abort ();
+    if (fib<false> (n, true) != o)
+      __builtin_abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/metadirective-template-3.C b/libgomp/testsuite/libgomp.c++/metadirective-template-3.C
new file mode 100644
index 00000000000..cc48fde0fda
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/metadirective-template-3.C
@@ -0,0 +1,41 @@ 
+# include <stdio.h>
+
+template <bool tasking> int
+fib (int n, bool flag)
+{
+  int i, j;
+  if (n < 2)
+    return n;
+  else if ( tasking && flag && n < 8 )  // serial/taskless cutoff for n<8
+    return fib<false> (n, false);
+  else
+    {
+#pragma omp metadirective				\
+  when (user = {condition (tasking && flag)}: task shared(i))	\
+  when (user = {condition (!tasking && !flag)}: nothing) \
+  otherwise (error at(execution) message("oops 1"))
+      i = fib<tasking> (n - 1, flag);
+#pragma omp metadirective					\
+  when (user = {condition (score(10): tasking && flag)}: task shared(j)) \
+  when (user = {condition (tasking || flag)} : \
+	  error at(execution) message ("oops 2"))
+      j = fib<tasking> (n - 2, flag);
+#pragma omp metadirective			\
+  when (user = {condition (tasking && flag)}: taskwait)
+      return i + j;
+    }
+}
+
+int main ()
+{
+  int n = 15, o = 610;
+#pragma omp parallel
+#pragma omp single
+  {
+    if (fib<true> (n, true) != o)
+      __builtin_abort ();
+    if (fib<false> (n, false) != o)
+      __builtin_abort ();
+  }
+  return 0;
+}