parser.h (struct cp_omp_declare_simd_data): Add clauses member.

	gcc/cp/
	* parser.h (struct cp_omp_declare_simd_data): Add clauses member.
	(struct cp_parser): Change type the of oacc_routine to
	cp_omp_declare_simd_data.
	* parser.c (cp_ensure_no_oacc_routine): Rework to use
	cp_omp_declare_simd_data.
	(cp_parser_simple_declaration): Remove boolean first.  Update call to
	cp_parser_init_declarator. Don't NULL out oacc_routine.
	(cp_parser_init_declarator): Remove boolean first parameter.  Update
	calls to cp_finalize_oacc_routine.
	(cp_parser_late_return_type_opt): Handle acc routines. 
	(cp_parser_member_declaration): Remove first variable.  Handle
	acc routines like omp declare simd.
	(cp_parser_function_definition_from_specifiers_and_declarator): Update
	call to cp_finalize_oacc_routine.
	(cp_parser_single_declaration): Update call to
	cp_parser_init_declarator.
	(cp_parser_save_member_function_body): Remove first_decl parameter.
	Update call to cp_finalize_oacc_routine.
	(cp_parser_finish_oacc_routine): Delete.
	(cp_parser_oacc_routine): Rework to use cp_omp_declare_simd_data.
	(cp_parser_late_parsing_oacc_routine): New function.
	(cp_finalize_oacc_routine): Remove first argument.  Add more error
	handling and set the acc routine and 'omp declare target' attributes.
	(cp_parser_pragma): Remove unnecessary call to
	cp_ensure_no_oacc_routine.

From-SVN: r230607

gcc/cp/ChangeLog

+2015-11-19  Cesar Philippidis  <cesar@codesourcery.com>
+
+	* parser.h (struct cp_omp_declare_simd_data): Add clauses member.
+	(struct cp_parser): Change type the of oacc_routine to
+	cp_omp_declare_simd_data.
+	* parser.c (cp_ensure_no_oacc_routine): Rework to use
+	cp_omp_declare_simd_data.
+	(cp_parser_simple_declaration): Remove boolean first.  Update call to
+	cp_parser_init_declarator. Don't NULL out oacc_routine.
+	(cp_parser_init_declarator): Remove boolean first parameter.  Update
+	calls to cp_finalize_oacc_routine.
+	(cp_parser_late_return_type_opt): Handle acc routines.
+	(cp_parser_member_declaration): Remove first variable.  Handle
+	acc routines like omp declare simd.
+	(cp_parser_function_definition_from_specifiers_and_declarator): Update
+	call to cp_finalize_oacc_routine.
+	(cp_parser_single_declaration): Update call to
+	cp_parser_init_declarator.
+	(cp_parser_save_member_function_body): Remove first_decl parameter.
+	Update call to cp_finalize_oacc_routine.
+	(cp_parser_finish_oacc_routine): Delete.
+	(cp_parser_oacc_routine): Rework to use cp_omp_declare_simd_data.
+	(cp_parser_late_parsing_oacc_routine): New function.
+	(cp_finalize_oacc_routine): Remove first argument.  Add more error
+	handling and set the acc routine and 'omp declare target' attributes.
+	(cp_parser_pragma): Remove unnecessary call to
+	cp_ensure_no_oacc_routine.
+
 2015-11-17  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* cp-gimplify.c (cp_fold_r): Add support for OACC_LOOP.

gcc/cp/parser.c

@@ -241,7 +241,7 @@ static bool cp_parser_omp_declare_reduction_exprs
 static tree cp_parser_cilk_simd_vectorlength 
   (cp_parser *, tree, bool);
 static void cp_finalize_oacc_routine
-  (cp_parser *, tree, bool, bool);
+  (cp_parser *, tree, bool);
 
 /* Manifest constants.  */
 #define CP_LEXER_BUFFER_SIZE ((256 * 1024) / sizeof (cp_token))
@@ -1318,13 +1318,21 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl)
     }
 }
 
-/* Diagnose if #pragma omp routine isn't followed immediately
-   by function declaration or definition.   */
+/* Diagnose if #pragma acc routine isn't followed immediately by function
+   declaration or definition.  */
 
 static inline void
 cp_ensure_no_oacc_routine (cp_parser *parser)
 {
-  cp_finalize_oacc_routine (parser, NULL_TREE, false, true);
+  if (parser->oacc_routine && !parser->oacc_routine->error_seen)
+    {
+      tree clauses = parser->oacc_routine->clauses;
+      location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses));
+
+      error_at (loc, "%<#pragma oacc routine%> not followed by function "
+		"declaration or definition");
+      parser->oacc_routine = NULL;
+    }
 }
 
 /* Decl-specifiers.  */
@@ -2130,7 +2138,7 @@ static tree cp_parser_decltype
 
 static tree cp_parser_init_declarator
   (cp_parser *, cp_decl_specifier_seq *, vec<deferred_access_check, va_gc> *,
-   bool, bool, int, bool *, tree *, bool, location_t *);
+   bool, bool, int, bool *, tree *, location_t *);
 static cp_declarator *cp_parser_declarator
   (cp_parser *, cp_parser_declarator_kind, int *, bool *, bool, bool);
 static cp_declarator *cp_parser_direct_declarator
@@ -2186,6 +2194,9 @@ static tree cp_parser_late_parsing_omp_declare_simd
 static tree cp_parser_late_parsing_cilk_simd_fn_info
   (cp_parser *, tree);
 
+static tree cp_parser_late_parsing_oacc_routine
+  (cp_parser *, tree);
+
 static tree synthesize_implicit_template_parm
   (cp_parser *, tree);
 static tree finish_fully_implicit_template
@@ -2440,7 +2451,7 @@ static tree cp_parser_single_declaration
 static tree cp_parser_functional_cast
   (cp_parser *, tree);
 static tree cp_parser_save_member_function_body
-  (cp_parser *, cp_decl_specifier_seq *, cp_declarator *, tree, bool);
+  (cp_parser *, cp_decl_specifier_seq *, cp_declarator *, tree);
 static tree cp_parser_save_nsdmi
   (cp_parser *);
 static tree cp_parser_enclosed_template_argument_list
@@ -11870,7 +11881,6 @@ cp_parser_simple_declaration (cp_parser* parser,
   bool saw_declarator;
   location_t comma_loc = UNKNOWN_LOCATION;
   location_t init_loc = UNKNOWN_LOCATION;
-  bool first = true;
 
   if (maybe_range_for_decl)
     *maybe_range_for_decl = NULL_TREE;
@@ -11967,10 +11977,7 @@ cp_parser_simple_declaration (cp_parser* parser,
 					declares_class_or_enum,
 					&function_definition_p,
 					maybe_range_for_decl,
-					first,
 					&init_loc);
-      first = false;
-
       /* If an error occurred while parsing tentatively, exit quickly.
 	 (That usually happens when in the body of a function; each
 	 statement is treated as a declaration-statement until proven
@@ -12069,9 +12076,6 @@ cp_parser_simple_declaration (cp_parser* parser,
 
  done:
   pop_deferring_access_checks ();
-
-  /* Reset any acc routine clauses.  */
-  parser->oacc_routine = NULL;
 }
 
 /* Parse a decl-specifier-seq.
@@ -17811,8 +17815,6 @@ cp_parser_asm_definition (cp_parser* parser)
    if present, will not be consumed.  If returned, this declarator will be
    created with SD_INITIALIZED but will not call cp_finish_decl.
 
-   FIRST indicates if this is the first declarator in a declaration sequence.
-
    If INIT_LOC is not NULL, and *INIT_LOC is equal to UNKNOWN_LOCATION,
    and there is an initializer, the pointed location_t is set to the
    location of the '=' or `(', or '{' in C++11 token introducing the
@@ -17827,7 +17829,6 @@ cp_parser_init_declarator (cp_parser* parser,
 			   int declares_class_or_enum,
 			   bool* function_definition_p,
 			   tree* maybe_range_for_decl,
-			   bool first,
 			   location_t* init_loc)
 {
   cp_token *token = NULL, *asm_spec_start_token = NULL,
@@ -17964,8 +17965,7 @@ cp_parser_init_declarator (cp_parser* parser,
 	    decl = cp_parser_save_member_function_body (parser,
 							decl_specifiers,
 							declarator,
-							prefix_attributes,
-							true);
+							prefix_attributes);
 	  else
 	    decl =
 	      (cp_parser_function_definition_from_specifiers_and_declarator
@@ -18069,7 +18069,7 @@ cp_parser_init_declarator (cp_parser* parser,
 			 range_for_decl_p? SD_INITIALIZED : is_initialized,
 			 attributes, prefix_attributes, &pushed_scope);
       cp_finalize_omp_declare_simd (parser, decl);
-      cp_finalize_oacc_routine (parser, decl, false, first);
+      cp_finalize_oacc_routine (parser, decl, false);
       /* Adjust location of decl if declarator->id_loc is more appropriate:
 	 set, and decl wasn't merged with another decl, in which case its
 	 location would be different from input_location, and more accurate.  */
@@ -18183,7 +18183,7 @@ cp_parser_init_declarator (cp_parser* parser,
       if (decl && TREE_CODE (decl) == FUNCTION_DECL)
 	cp_parser_save_default_args (parser, decl);
       cp_finalize_omp_declare_simd (parser, decl);
-      cp_finalize_oacc_routine (parser, decl, false, first);
+      cp_finalize_oacc_routine (parser, decl, false);
     }
 
   /* Finish processing the declaration.  But, skip member
@@ -19289,13 +19289,17 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator,
 
   bool cilk_simd_fn_vector_p = (parser->cilk_simd_fn_info 
 				&& declarator && declarator->kind == cdk_id);
-  
+
+  bool oacc_routine_p = (parser->oacc_routine
+			 && declarator
+			 && declarator->kind == cdk_id);
+
   /* Peek at the next token.  */
   token = cp_lexer_peek_token (parser->lexer);
   /* A late-specified return type is indicated by an initial '->'. */
   if (token->type != CPP_DEREF
       && token->keyword != RID_REQUIRES
-      && !(declare_simd_p || cilk_simd_fn_vector_p))
+      && !(declare_simd_p || cilk_simd_fn_vector_p || oacc_routine_p))
     return NULL_TREE;
 
   tree save_ccp = current_class_ptr;
@@ -19326,6 +19330,10 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator,
     declarator->std_attributes
       = cp_parser_late_parsing_omp_declare_simd (parser,
 						 declarator->std_attributes);
+  if (oacc_routine_p)
+    declarator->std_attributes
+      = cp_parser_late_parsing_oacc_routine (parser,
+					     declarator->std_attributes);
 
   if (quals >= 0)
     {
@@ -21887,7 +21895,6 @@ cp_parser_member_declaration (cp_parser* parser)
   else
     {
       bool assume_semicolon = false;
-      bool first = true;
 
       /* Clear attributes from the decl_specifiers but keep them
 	 around as prefix attributes that apply them to the entity
@@ -22075,10 +22082,7 @@ cp_parser_member_declaration (cp_parser* parser)
 		  decl = cp_parser_save_member_function_body (parser,
 							      &decl_specifiers,
 							      declarator,
-							      attributes,
-							      first);
-		  first = false;
-
+							      attributes);
 		  if (parser->fully_implicit_function_template_p)
 		    decl = finish_fully_implicit_template (parser, decl);
 		  /* If the member was not a friend, declare it here.  */
@@ -22108,8 +22112,7 @@ cp_parser_member_declaration (cp_parser* parser)
 	    }
 
 	  cp_finalize_omp_declare_simd (parser, decl);
-	  cp_finalize_oacc_routine (parser, decl, false, first);
-	  first = false;
+	  cp_finalize_oacc_routine (parser, decl, false);
 
 	  /* Reset PREFIX_ATTRIBUTES.  */
 	  while (attributes && TREE_CHAIN (attributes) != first_attribute)
@@ -22172,9 +22175,6 @@ cp_parser_member_declaration (cp_parser* parser)
 	  if (assume_semicolon)
 	    goto out;
 	}
-
-      /* Reset any OpenACC routine clauses.  */
-      parser->oacc_routine = NULL;
     }
 
   cp_parser_require (parser, CPP_SEMICOLON, RT_SEMICOLON);
@@ -24716,7 +24716,7 @@ cp_parser_function_definition_from_specifiers_and_declarator
     {
       cp_finalize_omp_declare_simd (parser, current_function_decl);
       parser->omp_declare_simd = NULL;
-      cp_finalize_oacc_routine (parser, current_function_decl, true, true);
+      cp_finalize_oacc_routine (parser, current_function_decl, true);
       parser->oacc_routine = NULL;
     }
 
@@ -25282,7 +25282,7 @@ cp_parser_single_declaration (cp_parser* parser,
 				        member_p,
 				        declares_class_or_enum,
 				        &function_definition_p,
-					NULL, true, NULL);
+					NULL, NULL);
 
     /* 7.1.1-1 [dcl.stc]
 
@@ -25384,15 +25384,14 @@ cp_parser_functional_cast (cp_parser* parser, tree type)
 /* Save the tokens that make up the body of a member function defined
    in a class-specifier.  The DECL_SPECIFIERS and DECLARATOR have
    already been parsed.  The ATTRIBUTES are any GNU "__attribute__"
-   specifiers applied to the declaration. FIRST_DECL indicates if
-   DECLARATOR is the first declarator in a declaration sequence.  Returns
-   the FUNCTION_DECL for the member function.  */
+   specifiers applied to the declaration.  Returns the FUNCTION_DECL
+   for the member function.  */
 
 static tree
 cp_parser_save_member_function_body (cp_parser* parser,
 				     cp_decl_specifier_seq *decl_specifiers,
 				     cp_declarator *declarator,
-				     tree attributes, bool first_decl)
+				     tree attributes)
 {
   cp_token *first;
   cp_token *last;
@@ -25401,7 +25400,7 @@ cp_parser_save_member_function_body (cp_parser* parser,
   /* Create the FUNCTION_DECL.  */
   fn = grokmethod (decl_specifiers, declarator, attributes);
   cp_finalize_omp_declare_simd (parser, fn);
-  cp_finalize_oacc_routine (parser, fn, true, first_decl);
+  cp_finalize_oacc_routine (parser, fn, true);
   /* If something went badly wrong, bail out now.  */
   if (fn == error_mark_node)
     {
@@ -35773,59 +35772,6 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ))
 
-/* Finalize #pragma acc routine clauses after direct declarator has
-   been parsed, and put that into "omp declare target" attribute.  */
-
-static void
-cp_parser_finish_oacc_routine (cp_parser *ARG_UNUSED (parser), tree fndecl,
-			       tree clauses, bool named, bool is_defn,
-			       bool first)
-{
-  location_t loc  = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses));
-
-  if (named && fndecl && is_overloaded_fn (fndecl)
-      && (TREE_CODE (fndecl) != FUNCTION_DECL
-	  || DECL_FUNCTION_TEMPLATE_P  (fndecl)))
-    {
-      error_at (loc, "%<#pragma acc routine%> names a set of overloads");
-      return;
-    }
-
-  if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL
-      || (!named && !first))
-    {
-      error_at (loc, "%<#pragma acc routine%> %s",
-		named ? "does not refer to a function"
-		: "not followed by single function");
-      return;
-    }
-
-  /* Perhaps we should use the same rule as declarations in different
-     namespaces?  */
-  if (named && !DECL_NAMESPACE_SCOPE_P (fndecl))
-    {
-      error_at (loc, "%<#pragma acc routine%> does not refer to a"
-		" namespace scope function");
-      return;
-    }
-
-  if (get_oacc_fn_attrib (fndecl))
-    error_at (loc, "%<#pragma acc routine%> already applied to %D", fndecl);
-
-  if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
-    error_at (OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)),
-	      "%<#pragma acc routine%> must be applied before %s",
-	      TREE_USED (fndecl) ? "use" : "definition");
-
-  /* Process for function attrib  */
-  tree dims = build_oacc_routine_dims (TREE_VALUE (clauses));
-  replace_oacc_fn_attrib (fndecl, dims);
-
-  /* Also attach as a declare.  */
-  DECL_ATTRIBUTES (fndecl)
-    = tree_cons (get_identifier ("omp declare target"),
-		 clauses, DECL_ATTRIBUTES (fndecl));
-}
 
 /* Parse the OpenACC routine pragma.  This has an optional '( name )'
    component, which must resolve to a declared namespace-scope
@@ -35837,16 +35783,50 @@ static void
 cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
 			enum pragma_context context)
 {
+  bool first_p = parser->oacc_routine == NULL;
+  location_t loc = pragma_tok->location;
+  cp_omp_declare_simd_data data;
+  if (first_p)
+    {
+      data.error_seen = false;
+      data.fndecl_seen = false;
+      data.tokens = vNULL;
+      data.clauses = NULL_TREE;
+      parser->oacc_routine = &data;
+    }
+
   tree decl = NULL_TREE;
   /* Create a dummy claue, to record location.  */
   tree c_head = build_omp_clause (pragma_tok->location, OMP_CLAUSE_SEQ);
 
   if (context != pragma_external)
-    cp_parser_error (parser, "%<#pragma acc routine%> not at file scope");
-  
+    {
+      cp_parser_error (parser, "%<#pragma acc routine%> not at file scope");
+      parser->oacc_routine->error_seen = true;
+      parser->oacc_routine = NULL;
+      return;
+    }
+
   /* Look for optional '( name )'.  */
-  if (cp_lexer_next_token_is (parser->lexer,CPP_OPEN_PAREN))
+  if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN))
     {
+      if (!first_p)
+	{
+	  while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)
+		 && cp_lexer_next_token_is_not (parser->lexer, CPP_EOF))
+	    cp_lexer_consume_token (parser->lexer);
+	  if (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+	    parser->oacc_routine->error_seen = true;
+	  cp_parser_require_pragma_eol (parser, pragma_tok);
+
+	  error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses),
+		    "%<#pragma oacc routine%> not followed by a single "
+		    "function declaration or definition");
+
+	  parser->oacc_routine->error_seen = true;
+	  return;
+	}
+
       cp_lexer_consume_token (parser->lexer);
       cp_token *token = cp_lexer_peek_token (parser->lexer);
 
@@ -35868,36 +35848,192 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
 	  || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
 	{
 	  cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+	  parser->oacc_routine = NULL;
 	  return;
 	}
+
+      /* Build a chain of clauses.  */
+      parser->lexer->in_pragma = true;
+      tree clauses = NULL_TREE;
+      clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
+					    "#pragma acc routine",
+					    cp_lexer_peek_token
+					    (parser->lexer));
+
+      /* Force clauses to be non-null, by attaching context to it.  */
+      clauses = tree_cons (c_head, clauses, NULL_TREE);
+
+      if (decl && is_overloaded_fn (decl)
+	  && (TREE_CODE (decl) != FUNCTION_DECL
+	      || DECL_FUNCTION_TEMPLATE_P  (decl)))
+	{
+	  error_at (loc, "%<#pragma acc routine%> names a set of overloads");
+	  parser->oacc_routine = NULL;
+	  return;
+	}
+
+      /* Perhaps we should use the same rule as declarations in different
+	 namespaces?  */
+      if (!DECL_NAMESPACE_SCOPE_P (decl))
+	{
+	  error_at (loc, "%<#pragma acc routine%> does not refer to a "
+		    "namespace scope function");
+	  parser->oacc_routine = NULL;
+	  return;
+	}
+
+      if (!decl || TREE_CODE (decl) != FUNCTION_DECL)
+	{
+	  error_at (loc,
+		    "%<#pragma acc routine%> does not refer to a function");
+	  parser->oacc_routine = NULL;
+	  return;
+	}
+
+      data.clauses = clauses;
+
+      cp_finalize_oacc_routine (parser, decl, false);
+      data.tokens.release ();
+      parser->oacc_routine = NULL;
     }
+  else
+    {
+      while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)
+	     && cp_lexer_next_token_is_not (parser->lexer, CPP_EOF))
+	cp_lexer_consume_token (parser->lexer);
+      if (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+	parser->oacc_routine->error_seen = true;
+      cp_parser_require_pragma_eol (parser, pragma_tok);
+
+      struct cp_token_cache *cp
+	= cp_token_cache_new (pragma_tok, cp_lexer_peek_token (parser->lexer));
+      parser->oacc_routine->tokens.safe_push (cp);
 
-  /* Build a chain of clauses.  */
+      if (first_p)
+	parser->oacc_routine->clauses = c_head;
+
+      while (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA))
+	cp_parser_pragma (parser, context);
+
+      if (first_p)
+	{
+	  /* Create an empty list of clauses.  */
+	  parser->oacc_routine->clauses = tree_cons (c_head, NULL_TREE,
+						     NULL_TREE);
+	  cp_parser_declaration (parser);
+
+	  if (parser->oacc_routine
+	      && !parser->oacc_routine->error_seen
+	      && !parser->oacc_routine->fndecl_seen)
+	    error_at (loc, "%<#pragma acc routine%> not followed by "
+		      "function declaration or definition");
+
+	  data.tokens.release ();
+	  parser->oacc_routine = NULL;
+	}
+    }
+}
+
+/* Finalize #pragma acc routine clauses after direct declarator has
+   been parsed, and put that into "oacc routine" attribute.  */
+
+static tree
+cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
+{
+  struct cp_token_cache *ce;
+  cp_omp_declare_simd_data *data = parser->oacc_routine;
+  tree cl, clauses = parser->oacc_routine->clauses;
+  location_t loc;
+
+  loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses));
+  
+  if ((!data->error_seen && data->fndecl_seen)
+      || data->tokens.length () != 1)
+    {
+      error_at (loc, "%<#pragma oacc routine%> not followed by a single "
+		"function declaration or definition");
+      data->error_seen = true;
+      return attrs;
+    }
+  if (data->error_seen)
+    return attrs;
+
+  ce = data->tokens[0];
+
+  cp_parser_push_lexer_for_tokens (parser, ce);
   parser->lexer->in_pragma = true;
-  tree clauses = NULL_TREE;
-  clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
-					"#pragma acc routine",
-					cp_lexer_peek_token (parser->lexer));
+  gcc_assert (cp_lexer_peek_token (parser->lexer)->type == CPP_PRAGMA);
+
+  cp_token *pragma_tok = cp_lexer_consume_token (parser->lexer);
+  cl = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
+				  "#pragma oacc routine", pragma_tok);
+  cp_parser_pop_lexer (parser);
+
+  tree c_head = build_omp_clause (loc, OMP_CLAUSE_SEQ);
 
   /* Force clauses to be non-null, by attaching context to it.  */
-  clauses = tree_cons (c_head, clauses, NULL_TREE);
+  parser->oacc_routine->clauses = tree_cons (c_head, cl, NULL_TREE);
 
-  if (decl)
-    cp_parser_finish_oacc_routine (parser, decl, clauses, true, false, 0);
-  else
-    parser->oacc_routine = clauses;
+  data->fndecl_seen = true;
+  return attrs;
 }
 
 /* Apply any saved OpenACC routine clauses to a just-parsed
    declaration.  */
 
 static void
-cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn,
-			  bool first)
+cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
 {
-  if (parser->oacc_routine)
-    cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine,
-				   false, is_defn, first);
+  if (__builtin_expect (parser->oacc_routine != NULL, 0))
+    {
+      tree clauses = parser->oacc_routine->clauses;
+      location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses));
+
+      if (parser->oacc_routine->error_seen)
+	return;
+      
+      if (fndecl == error_mark_node)
+	{
+	  parser->oacc_routine = NULL;
+	  return;
+	}
+
+      if (TREE_CODE (fndecl) != FUNCTION_DECL)
+	{
+	  cp_ensure_no_oacc_routine (parser);
+	  return;
+	}
+
+      if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL)
+	{
+	  error_at (loc,
+		    "%<#pragma acc routine%> not followed by single function");
+	  parser->oacc_routine = NULL;
+	}
+	  
+      if (get_oacc_fn_attrib (fndecl))
+	{
+	  error_at (loc, "%<#pragma acc routine%> already applied to %D",
+		    fndecl);
+	  parser->oacc_routine = NULL;
+	}
+
+      if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+	{
+	  error_at (loc, "%<#pragma acc routine%> must be applied before %s",
+		    TREE_USED (fndecl) ? "use" : "definition");
+	  parser->oacc_routine = NULL;
+	}
+
+      /* Process for function attrib  */
+      tree dims = build_oacc_routine_dims (TREE_VALUE (clauses));
+      replace_oacc_fn_attrib (fndecl, dims);
+      
+      /* Add an "omp target" attribute.  */
+      DECL_ATTRIBUTES (fndecl)
+	= tree_cons (get_identifier ("omp declare target"),
+		     NULL_TREE, DECL_ATTRIBUTES (fndecl));
+    }
 }
 
 /* Main entry point to OpenMP statement pragmas.  */
@@ -36381,7 +36517,6 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
   id = pragma_tok->pragma_kind;
   if (id != PRAGMA_OMP_DECLARE_REDUCTION && id != PRAGMA_OACC_ROUTINE)
     cp_ensure_no_omp_declare_simd (parser);
-  cp_ensure_no_oacc_routine (parser);
   switch (id)
     {
     case PRAGMA_GCC_PCH_PREPROCESS:

gcc/cp/parser.h

@@ -203,6 +203,7 @@ struct cp_omp_declare_simd_data {
   bool error_seen; /* Set if error has been reported.  */
   bool fndecl_seen; /* Set if one fn decl/definition has been seen already.  */
   vec<cp_token_cache_ptr> tokens;
+  tree clauses;
 };
 
 
@@ -371,8 +372,8 @@ struct GTY(()) cp_parser {
      necessary.  */
   cp_omp_declare_simd_data * GTY((skip)) cilk_simd_fn_info;
 
-  /* OpenACC routine clauses for subsequent decl/defn.  */
-  tree oacc_routine;
+  /* Parsing information for #pragma acc routine.  */
+  cp_omp_declare_simd_data * GTY((skip)) oacc_routine;
   
   /* Nonzero if parsing a parameter list where 'auto' should trigger an implicit
      template parameter.  */