diff mbox

[1/2] OpenACC routine support

Message ID 564CCB29.30909@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Nov. 18, 2015, 7:02 p.m. UTC
On 11/10/2015 12:16 AM, Jakub Jelinek wrote:
> On Mon, Nov 09, 2015 at 09:28:47PM -0800, Cesar Philippidis wrote:
>> Here's the patch that Nathan was referring to. I ended up introducing a
>> boolean variable named first in the various functions which call
>> finalize_oacc_routines. The problem the original approach was having was
>> that the routine clauses is only applied to the first function
>> declarator in a declaration list. By using 'first', which is set to true
>> if the current declarator is the first in a sequence of declarators, I
>> was able to defer setting parser->oacc_routine to NULL.
> 
> The #pragma omp declare simd has identical restrictions, but doesn't need
> to add any of the first parameters to the C++ parser.
> So, what are you doing differently that you need it?  Handling both
> differently is a consistency issue, and unnecessary additional complexity to
> the parser.

I reworked how acc routines are handed in this patch to be more similar
to #pragma omp declare simd. Things get kind of messy though. For
starters, I had to add a new tree clauses member to
cp_omp_declare_simd_data. This serves two purposes:

  * It allows the c++ FE to record the location of the first
    #pragma acc routine, which is nice because it allows test cases to
    be shared with the c FE.

  * Unlike omp declare simd, only one acc routine may be associated with
    a function decl. This meant that I had to defer attaching the acc
    geometry and 'omp target' attributes to cp_finalize_oacc_routine
    instead of in cp_parser_late_parsing_oacc_routine like in omp. So
    what happens is, cp_parser_late_parsing_oacc_routine ends up
    creating a function geometry clause.

I don't really like this approach. I did try to postpone parsing the
clauses till cp_finalize_oacc_routine, but that got messy. Plus, while
I'd be able to remove the clauses field from cp_omp_declare_simd_data,
we'd still need a location_t field for cp_ensure_no_oacc_routine.

Is this OK for trunk?

Cesar

Comments

Jakub Jelinek Nov. 19, 2015, 11:23 a.m. UTC | #1
On Wed, Nov 18, 2015 at 11:02:01AM -0800, Cesar Philippidis wrote:
>  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));

Formatting, missing space before (clauses));

> @@ -19326,7 +19330,11 @@ 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);
> +  

Trailing whitespace at the end of line.

Otherwise LGTM.

	Jakub
Thomas Schwinge Dec. 1, 2015, 2:40 p.m. UTC | #2
Hi Cesar!

I noticed while working on other test cases:

On Wed, 18 Nov 2015 11:02:01 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote:
> --- a/gcc/cp/parser.c
> +++ b/gcc/cp/parser.c

> @@ -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;
> +    }
>  }

"#pragma acc routine", not "oacc".  Also in a few other places.

Next, in the function quoted above, you use "not followed by function
declaration or definition", but you use "not followed by a single
function declaration or definition" in a lot of (but not all) other
places -- is that intentional?

For example:

>  cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
>  			enum pragma_context context)
>  {
> [...]
> +	  error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses),
> +		    "%<#pragma oacc routine%> not followed by a single "
> +		    "function declaration or definition");

"a single".

> [...]
> +	  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");

Not "a single".

> +
> +	  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.  */

There is no "oacc routine" attribute (anymore)?

> +static tree
> +cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
> +{
> [...]
> +  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");

"a single".

(I have not verified all of the parser(s) source code.)


Grüße
 Thomas
Cesar Philippidis Dec. 1, 2015, 2:49 p.m. UTC | #3
On 12/01/2015 06:40 AM, Thomas Schwinge wrote:

> I noticed while working on other test cases:
> 
> On Wed, 18 Nov 2015 11:02:01 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote:
>> --- a/gcc/cp/parser.c
>> +++ b/gcc/cp/parser.c
> 
>> @@ -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;
>> +    }
>>  }
> 
> "#pragma acc routine", not "oacc".  Also in a few other places.

Good eyes. Thanks for catching that.

> Next, in the function quoted above, you use "not followed by function
> declaration or definition", but you use "not followed by a single
> function declaration or definition" in a lot of (but not all) other
> places -- is that intentional?

I probably wasn't being consistent. Which error message do you prefer?
I'll take a look at what the c front end does.

> For example:
> 
>>  cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
>>  			enum pragma_context context)
>>  {
>> [...]
>> +	  error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses),
>> +		    "%<#pragma oacc routine%> not followed by a single "
>> +		    "function declaration or definition");
> 
> "a single".
> 
>> [...]
>> +	  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");
> 
> Not "a single".
> 
>> +
>> +	  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.  */
> 
> There is no "oacc routine" attribute (anymore)?

You're right, it was renamed to 'oacc function'.

>> +static tree
>> +cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
>> +{
>> [...]
>> +  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");
> 
> "a single".
> 
> (I have not verified all of the parser(s) source code.)

Thanks. I'll go through and update the comments and error messages.

Cesar
Cesar Philippidis Dec. 2, 2015, 11:37 p.m. UTC | #4
On 12/01/2015 06:49 AM, Cesar Philippidis wrote:
> On 12/01/2015 06:40 AM, Thomas Schwinge wrote:
> 
>> I noticed while working on other test cases:
>>
>> On Wed, 18 Nov 2015 11:02:01 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote:
>>> --- a/gcc/cp/parser.c
>>> +++ b/gcc/cp/parser.c
>>
>>> @@ -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;
>>> +    }
>>>  }
>>
>> "#pragma acc routine", not "oacc".  Also in a few other places.
> 
> Good eyes. Thanks for catching that.
> 
>> Next, in the function quoted above, you use "not followed by function
>> declaration or definition", but you use "not followed by a single
>> function declaration or definition" in a lot of (but not all) other
>> places -- is that intentional?
> 
> I probably wasn't being consistent. Which error message do you prefer?
> I'll take a look at what the c front end does.
> 
>> For example:
>>
>>>  cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
>>>  			enum pragma_context context)
>>>  {
>>> [...]
>>> +	  error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses),
>>> +		    "%<#pragma oacc routine%> not followed by a single "
>>> +		    "function declaration or definition");
>>
>> "a single".
>>
>>> [...]
>>> +	  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");
>>
>> Not "a single".
>>
>>> +
>>> +	  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.  */
>>
>> There is no "oacc routine" attribute (anymore)?
> 
> You're right, it was renamed to 'oacc function'.
> 
>>> +static tree
>>> +cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
>>> +{
>>> [...]
>>> +  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");
>>
>> "a single".
>>
>> (I have not verified all of the parser(s) source code.)
> 
> Thanks. I'll go through and update the comments and error messages.

Here's the updated patch. The test cases were written in a way such that
none of them needed to be updated with these changes.

I'm tempted to commit this as obvious, but I want to make sure you're ok
with these new messages. The major change is to report these errors as
"pragma acc routine not followed by a function declaration or
definition". I think that's more descriptive then "not followed by a
single function". That said, it looks like the c front end uses the
latter error message.

Is this OK or do you prefer the "not followed by a single function" message?

Cesar
Thomas Schwinge Dec. 3, 2015, 8:36 a.m. UTC | #5
Hi Cesar!

On Wed, 2 Dec 2015 15:37:17 -0800, Cesar Philippidis <cesar_philippidis@mentor.com> wrote:
> On 12/01/2015 06:49 AM, Cesar Philippidis wrote:
> > On 12/01/2015 06:40 AM, Thomas Schwinge wrote:
> > 
> >> I noticed while working on other test cases:
> >>
> >> On Wed, 18 Nov 2015 11:02:01 -0800, Cesar Philippidis <cesar@codesourcery.com> wrote:
> >>> --- a/gcc/cp/parser.c
> >>> +++ b/gcc/cp/parser.c
> >>
> >>> @@ -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;
> >>> +    }
> >>>  }

> >> Next, in the function quoted above, you use "not followed by function
> >> declaration or definition", but you use "not followed by a single
> >> function declaration or definition" in a lot of (but not all) other
> >> places -- is that intentional?
> > 
> > I probably wasn't being consistent. Which error message do you prefer?
> > I'll take a look at what the c front end does.
> > 
> >> For example: [...]

> >> (I have not verified all of the parser(s) source code.)
> > 
> > Thanks. I'll go through and update the comments and error messages.
> 
> Here's the updated patch.

ENOPATCH.

> The test cases were written in a way such that
> none of them needed to be updated with these changes.

... which potentially means they'd match for all kinds of "random"
diagnostics.  ;-)

> I'm tempted to commit this as obvious, but I want to make sure you're ok
> with these new messages.

I don't care very much, as long as it's understandable for a user.  I
just tripped over this because of mismatches between C and C++ as well as
different C++ diagnostic variants.

> The major change is to report these errors as
> "pragma acc routine not followed by a function declaration or
> definition". I think that's more descriptive then "not followed by a
> single function". That said, it looks like the c front end uses the
> latter error message.

(In the C front end, the "a" is missing: "not followed by single
function"; that should be fixed up as well.)

> Is this OK or do you prefer the "not followed by a single function" message?

"not followed by a function declaration or definition" sounds good to me.


Grüße
 Thomas
diff mbox

Patch

2015-11-17  Cesar Philippidis  <cesar@codesourcery.com>

	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.

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 0e1116b..8de3bce 100644
--- a/gcc/cp/parser.c
+++ b/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,7 +19330,11 @@  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)
     {
       current_class_ptr = save_ccp;
@@ -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);
+
+      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];
 
-  /* Build a chain of clauses.  */
+  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:
diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h
index 022d037..a6b8e74 100644
--- a/gcc/cp/parser.h
+++ b/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.  */