Re: [1/2] OpenACC routine support

[Cesar Philippidis <cesar at codesourcery dot com>]

On 11/09/2015 04:31 PM, Nathan Sidwell wrote:
> On 11/03/15 10:35, Jakub Jelinek wrote:
>> On Mon, Nov 02, 2015 at 02:21:43PM -0500, Nathan Sidwell wrote:
>>> --- gcc/c/c-parser.c    (revision 229667)
>>> +++ gcc/c/c-parser.c    (working copy)
>>> @@ -1160,7 +1160,8 @@ enum c_parser_prec {
>>>   static void c_parser_external_declaration (c_parser *);
>>>   static void c_parser_asm_definition (c_parser *);
>>>   static void c_parser_declaration_or_fndef (c_parser *, bool, bool,
>>> bool,
>>> -                       bool, bool, tree *, vec<c_token>);
>>> +                       bool, bool, tree *, vec<c_token>,
>>> +                       tree);
>>
>> Wonder if this shouldn't be tree = NULL_TREE, then you'd avoid most of
>> the
>> c_parser_declaration_or_fndef caller changes.
>>
>> Otherwise, LGTM.
> 
> This is the patch I've just committed.  It includes c parser adjustments
> to detect the case of two function decls with a single type specifier. 
> Cesar will be applying a patch for the C++ parser for the same  case.

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.

Nathan already approved this patch, so I've applied it to trunk.

Cesar

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

	gcc/cp/
	* parser.c (cp_finalize_oacc_routine): New boolean first argument.
	(cp_ensure_no_oacc_routine): Update call to cp_finalize_oacc_routine.
	(cp_parser_simple_declaration): Maintain a boolean first to keep track
	of each new declarator.  Propagate it to cp_parser_init_declarator.
	(cp_parser_init_declarator): New boolean first argument.  Propagate it
	to cp_parser_save_member_function_body and cp_finalize_oacc_routine.
	(cp_parser_member_declaration): Likewise.
	(cp_parser_single_declaration): Update call to
	cp_parser_init_declarator.
	(cp_parser_save_member_function_body): New boolean first_decl argument.
	Propagate it to cp_finalize_oacc_routine.
	(cp_parser_finish_oacc_routine): New boolean first argument.  Use it to
	determine if multiple declarators follow a routine construct.
	(cp_parser_oacc_routine): Update call to cp_parser_finish_oacc_routine.

	gcc/testsuite/
	* c-c++-common/goacc/routine-5.c: Enable c++ tests.

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 6fc2c6a..f3b4b46 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -246,7 +246,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);
+  (cp_parser *, tree, bool, bool);
 
 /* Manifest constants.  */
 #define CP_LEXER_BUFFER_SIZE ((256 * 1024) / sizeof (cp_token))
@@ -1329,7 +1329,7 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl)
 static inline void
 cp_ensure_no_oacc_routine (cp_parser *parser)
 {
-  cp_finalize_oacc_routine (parser, NULL_TREE, false);
+  cp_finalize_oacc_routine (parser, NULL_TREE, false, true);
 }
 
 /* Decl-specifiers.  */
@@ -2135,7 +2135,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 *, location_t *);
+   bool, bool, int, bool *, tree *, bool, 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
@@ -2445,7 +2445,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);
+  (cp_parser *, cp_decl_specifier_seq *, cp_declarator *, tree, bool);
 static tree cp_parser_save_nsdmi
   (cp_parser *);
 static tree cp_parser_enclosed_template_argument_list
@@ -11909,6 +11909,7 @@ 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;
@@ -12005,7 +12006,10 @@ 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
@@ -12104,6 +12108,9 @@ 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.
@@ -17843,6 +17850,8 @@ 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
@@ -17857,6 +17866,7 @@ 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,
@@ -17993,7 +18003,8 @@ cp_parser_init_declarator (cp_parser* parser,
 	    decl = cp_parser_save_member_function_body (parser,
 							decl_specifiers,
 							declarator,
-							prefix_attributes);
+							prefix_attributes,
+							true);
 	  else
 	    decl =
 	      (cp_parser_function_definition_from_specifiers_and_declarator
@@ -18097,7 +18108,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);
+      cp_finalize_oacc_routine (parser, decl, false, first);
       /* 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.  */
@@ -18211,7 +18222,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);
+      cp_finalize_oacc_routine (parser, decl, false, first);
     }
 
   /* Finish processing the declaration.  But, skip member
@@ -21915,6 +21926,7 @@ 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
@@ -22102,7 +22114,10 @@ cp_parser_member_declaration (cp_parser* parser)
 		  decl = cp_parser_save_member_function_body (parser,
 							      &decl_specifiers,
 							      declarator,
-							      attributes);
+							      attributes,
+							      first);
+		  first = false;
+
 		  if (parser->fully_implicit_function_template_p)
 		    decl = finish_fully_implicit_template (parser, decl);
 		  /* If the member was not a friend, declare it here.  */
@@ -22132,7 +22147,8 @@ cp_parser_member_declaration (cp_parser* parser)
 	    }
 
 	  cp_finalize_omp_declare_simd (parser, decl);
-	  cp_finalize_oacc_routine (parser, decl, false);
+	  cp_finalize_oacc_routine (parser, decl, false, first);
+	  first = false;
 
 	  /* Reset PREFIX_ATTRIBUTES.  */
 	  while (attributes && TREE_CHAIN (attributes) != first_attribute)
@@ -22195,6 +22211,9 @@ 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);
@@ -24736,7 +24755,8 @@ 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);
+      cp_finalize_oacc_routine (parser, current_function_decl, true, true);
+      parser->oacc_routine = NULL;
     }
 
   if (!success_p)
@@ -25301,7 +25321,7 @@ cp_parser_single_declaration (cp_parser* parser,
 				        member_p,
 				        declares_class_or_enum,
 				        &function_definition_p,
-					NULL, NULL);
+					NULL, true, NULL);
 
     /* 7.1.1-1 [dcl.stc]
 
@@ -25403,14 +25423,15 @@ 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.  Returns the FUNCTION_DECL
-   for the member function.  */
+   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.  */
 
 static tree
 cp_parser_save_member_function_body (cp_parser* parser,
 				     cp_decl_specifier_seq *decl_specifiers,
 				     cp_declarator *declarator,
-				     tree attributes)
+				     tree attributes, bool first_decl)
 {
   cp_token *first;
   cp_token *last;
@@ -25419,7 +25440,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);
+  cp_finalize_oacc_routine (parser, fn, true, first_decl);
   /* If something went badly wrong, bail out now.  */
   if (fn == error_mark_node)
     {
@@ -35625,7 +35646,8 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok,
 
 static void
 cp_parser_finish_oacc_routine (cp_parser *ARG_UNUSED (parser), tree fndecl,
-			       tree clauses, bool named, bool is_defn)
+			       tree clauses, bool named, bool is_defn,
+			       bool first)
 {
   location_t loc  = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses));
 
@@ -35637,7 +35659,8 @@ cp_parser_finish_oacc_routine (cp_parser *ARG_UNUSED (parser), tree fndecl,
       return;
     }
 
-  if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL)
+  if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL
+      || (!named && !first))
     {
       error_at (loc, "%<#pragma acc routine%> %s",
 		named ? "does not refer to a function"
@@ -35728,7 +35751,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
   clauses = tree_cons (c_head, clauses, NULL_TREE);
 
   if (decl)
-    cp_parser_finish_oacc_routine (parser, decl, clauses, true, false);
+    cp_parser_finish_oacc_routine (parser, decl, clauses, true, false, 0);
   else
     parser->oacc_routine = clauses;
 }
@@ -35737,14 +35760,12 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
    declaration.  */
 
 static void
-cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
+cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn,
+			  bool first)
 {
   if (parser->oacc_routine)
-    {
-      cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine,
-				     false, is_defn);
-      parser->oacc_routine = NULL_TREE;
-    }
+    cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine,
+				   false, is_defn, first);
 }
 
 /* Main entry point to OpenMP statement pragmas.  */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index c2a8fb2..ccda097 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -3,18 +3,14 @@
 #pragma acc routine /* { dg-error "not followed by" } */
 int a;
 
-#if 0 /* Disable for the moment.  */
-#pragma acc routine /* dg-error "not followed by" */
+#pragma acc routine /* { dg-error "not followed by" } */
 void fn1 (void), fn1b (void);
-#endif
 
 #pragma acc routine /* { dg-error "not followed by" } */
 int b, fn2 (void);
 
-#if 0 /* Disable for the moment.  */
-#pragma acc routine /*  dg-error "not followed by"  */
+#pragma acc routine /* { dg-error "not followed by" } */
 int fn3 (void), b2;
-#endif
 
 #pragma acc routine /* { dg-error "not followed by" } */
 typedef struct c c;