2015-08-17 Nathan Sidwell <nathan@codesourcery.com>
c/
* c-parser.c (c_finish_oacc_routine): Add is_defn arg and adjust
all callers. Detect duplicate, post-use or post-defn application.
cp/
* parser.h (struct cp_parser): Change oacc_routine field
tree. Remove named_oacc_routines field.
* parser.c (cp_ensure_no_omp_declare_simd): Remove oacc routine
checking, move to ...
(cp_ensure_no_oacc_routine): ... here. New function, adjust all
callers.
(cp_parser_late_parsing_oacc_routine): Delete.
(cp_parser_new): Adjust.
(cp_parser_linkage_specification): Call cp_ensure_no_oacc_routine.
(cp_finalize_oacc_routine): Reimplement. Adjust all callers.
(cp_parser_late_return_type_opt): Remove oacc routine handling.
(cp_parser_omp_declare_simd): Remove oacc routine handling. Adjust
callers.
(cp_parser_finish_oacc_routine): New.
(cp_parser_oacc_routine): Reimplement.
testsuite/
* c-c++-common/goacc/routine-2.c: Insert declaration.
libgomp/
* testsuite/libgomp.c-c++-common/routine-2.c: Insert declaration.
===================================================================
@@ -1762,7 +1762,7 @@ finish_oacc_declare (tree fnbody, tree d
static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>);
-static void c_finish_oacc_routine (c_parser *, tree, tree);
+static void c_finish_oacc_routine (c_parser *, tree, tree, bool);
/* Parse a declaration or function definition (C90 6.5, 6.7.1, C99
6.7, 6.9.1). If FNDEF_OK is true, a function definition is
@@ -2020,7 +2020,8 @@ c_parser_declaration_or_fndef (c_parser
c_finish_omp_declare_simd (parser, NULL_TREE, NULL_TREE,
omp_declare_simd_clauses);
if (oacc_routine_clauses)
- c_finish_oacc_routine (parser, NULL_TREE, oacc_routine_clauses);
+ c_finish_oacc_routine (parser, NULL_TREE,
+ oacc_routine_clauses, false);
c_parser_skip_to_end_of_block_or_statement (parser);
return;
}
@@ -2117,9 +2118,6 @@ c_parser_declaration_or_fndef (c_parser
|| !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
c_finish_omp_declare_simd (parser, d, NULL_TREE,
omp_declare_simd_clauses);
-
- if (oacc_routine_clauses)
- c_finish_oacc_routine (parser, d, oacc_routine_clauses);
}
else
{
@@ -2133,14 +2131,14 @@ c_parser_declaration_or_fndef (c_parser
|| !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
c_finish_omp_declare_simd (parser, d, NULL_TREE,
omp_declare_simd_clauses);
- if (oacc_routine_clauses)
- c_finish_oacc_routine (parser, d, oacc_routine_clauses);
-
+
start_init (d, asm_name, global_bindings_p ());
init_loc = c_parser_peek_token (parser)->location;
init = c_parser_initializer (parser);
finish_init ();
}
+ if (oacc_routine_clauses)
+ c_finish_oacc_routine (parser, d, oacc_routine_clauses, false);
if (d != error_mark_node)
{
maybe_warn_string_init (init_loc, TREE_TYPE (d), init);
@@ -2186,8 +2184,8 @@ c_parser_declaration_or_fndef (c_parser
temp_pop_parm_decls ();
}
if (oacc_routine_clauses)
- c_finish_oacc_routine (parser, d, oacc_routine_clauses);
-
+ c_finish_oacc_routine (parser, d, oacc_routine_clauses, false);
+
if (d)
finish_decl (d, UNKNOWN_LOCATION, NULL_TREE,
NULL_TREE, asm_name);
@@ -2298,10 +2296,10 @@ c_parser_declaration_or_fndef (c_parser
|| !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
c_finish_omp_declare_simd (parser, current_function_decl, NULL_TREE,
omp_declare_simd_clauses);
-
if (oacc_routine_clauses)
c_finish_oacc_routine (parser, current_function_decl,
- oacc_routine_clauses);
+ oacc_routine_clauses, true);
+
DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus
= c_parser_peek_token (parser)->location;
@@ -13279,6 +13277,10 @@ c_parser_oacc_parallel (location_t loc,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND))
+/* Parse an OpenACC routine directive. For named directives, we apply
+ immediately to the named function. For unnamed ones we then parse
+ a declaration or definition, which must be for a function. */
+
static void
c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
{
@@ -13325,24 +13327,36 @@ c_parser_oacc_routine (c_parser *parser,
clauses = tree_cons (c_head, clauses, NULL_TREE);
if (decl)
- c_finish_oacc_routine (parser, decl, clauses);
+ c_finish_oacc_routine (parser, decl, clauses, false);
else
c_parser_declaration_or_fndef (parser, true, false, false, false,
true, NULL, vNULL, clauses);
}
+/* Finalize an OpenACC routine pragma, applying it to FNDECL. CLAUSES
+ are the parsed clauses. IS_DEFN is true if we're applying it to
+ the definition (so expect FNDEF to look somewhat defined. */
+
static void
c_finish_oacc_routine (c_parser *ARG_UNUSED (parser),
- tree fndecl, tree clauses)
+ tree fndecl, tree clauses, bool is_defn)
{
+ location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses));
+
if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL)
{
if (fndecl != error_mark_node)
- error_at (OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses)),
- "%<#pragma acc routine%> does not refer to a function");
+ error_at (loc, "%<#pragma acc routine%> does not refer to a 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 (loc, "%<#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);
===================================================================
@@ -251,6 +251,7 @@ static bool cp_parser_omp_declare_reduct
(tree, cp_parser *);
static tree cp_parser_cilk_simd_vectorlength
(cp_parser *, tree, bool);
+static void cp_finalize_oacc_routine (cp_parser *, tree, bool);
/* Manifest constants. */
#define CP_LEXER_BUFFER_SIZE ((256 * 1024) / sizeof (cp_token))
@@ -1292,8 +1293,7 @@ cp_token_cache_new (cp_token *first, cp_
}
/* Diagnose if #pragma omp declare simd isn't followed immediately
- by function declaration or definition. Likewise for
- #pragma acc routine. */
+ by function declaration or definition. */
static inline void
cp_ensure_no_omp_declare_simd (cp_parser *parser)
@@ -1304,13 +1304,6 @@ cp_ensure_no_omp_declare_simd (cp_parser
"function declaration or definition");
parser->omp_declare_simd = NULL;
}
-
- if (parser->oacc_routine && !parser->oacc_routine->error_seen)
- {
- error ("%<#pragma acc routine%> not immediately followed by "
- "function declaration or definition");
- parser->oacc_routine = NULL;
- }
}
/* Finalize #pragma omp declare simd clauses after FNDECL has been parsed,
@@ -1334,56 +1327,13 @@ cp_finalize_omp_declare_simd (cp_parser
}
}
-/* Finalize #pragma acc routine clauses after FNDECL has been parsed,
- and put that into "acc routine" attribute. */
+/* Diagnose if #pragma omp routine isn't followed immediately
+ by function declaration or definition. */
static inline void
-cp_finalize_oacc_routine (cp_parser *parser, tree fndecl)
+cp_ensure_no_oacc_routine (cp_parser *parser)
{
- if (__builtin_expect (parser->omp_declare_simd != NULL, 0))
- {
- if (fndecl == error_mark_node)
- {
- parser->omp_declare_simd = NULL;
- return;
- }
- if (TREE_CODE (fndecl) != FUNCTION_DECL)
- {
- cp_ensure_no_omp_declare_simd (parser);
- return;
- }
- }
- else // Is this fndecl associated with a named routine?
- {
- if (fndecl == NULL_TREE || fndecl == error_mark_node
- || TREE_CODE (fndecl) != FUNCTION_DECL)
- return;
-
- bool found = false;
- int i;
- tree t, clauses = NULL_TREE;
-
- for (i = 0; vec_safe_iterate (parser->named_oacc_routines, i, &t); i++)
- {
- if (!strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)),
- IDENTIFIER_POINTER (t)))
- {
- found = true;
- clauses = TREE_CHAIN (t);
- break;
- }
- }
-
- if (!found)
- return;
-
- if (clauses != NULL_TREE)
- clauses = tree_cons (NULL_TREE, clauses, NULL_TREE);
- clauses = build_tree_list (get_identifier ("omp declare target"),
- clauses);
- TREE_CHAIN (clauses) = DECL_ATTRIBUTES (fndecl);
- DECL_ATTRIBUTES (fndecl) = clauses;
- }
+ cp_finalize_oacc_routine (parser, NULL_TREE, false);
}
/* Decl-specifiers. */
@@ -2244,9 +2194,6 @@ static tree cp_parser_late_parsing_omp_d
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
@@ -3672,9 +3619,8 @@ cp_parser_new (void)
parser->implicit_template_parms = 0;
parser->implicit_template_scope = 0;
- /* The list of OpenACC routines pragmas is unitialized. */
+ /* Active OpenACC routine clauses. */
parser->oacc_routine = NULL;
- parser->named_oacc_routines = NULL;
/* Allow constrained-type-specifiers. */
parser->prevent_constrained_type_specifiers = 0;
@@ -12299,6 +12245,7 @@ cp_parser_linkage_specification (cp_pars
if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_BRACE))
{
cp_ensure_no_omp_declare_simd (parser);
+ cp_ensure_no_oacc_routine (parser);
/* Consume the `{' token. */
cp_lexer_consume_token (parser->lexer);
@@ -16807,6 +16754,7 @@ cp_parser_namespace_definition (cp_parse
bool is_inline;
cp_ensure_no_omp_declare_simd (parser);
+ cp_ensure_no_oacc_routine (parser);
if (cp_lexer_next_token_is_keyword (parser->lexer, RID_INLINE))
{
maybe_warn_cpp0x (CPP0X_INLINE_NAMESPACES);
@@ -17787,7 +17735,7 @@ cp_parser_init_declarator (cp_parser* pa
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);
+ 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. */
@@ -17901,7 +17849,7 @@ cp_parser_init_declarator (cp_parser* pa
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);
+ cp_finalize_oacc_routine (parser, decl, false);
}
/* Finish processing the declaration. But, skip member
@@ -18970,15 +18918,12 @@ cp_parser_late_return_type_opt (cp_parse
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 || oacc_routine_p))
+ && !(declare_simd_p || cilk_simd_fn_vector_p))
return NULL_TREE;
tree save_ccp = current_class_ptr;
@@ -19009,11 +18954,6 @@ cp_parser_late_return_type_opt (cp_parse
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;
@@ -20469,6 +20409,7 @@ cp_parser_class_specifier_1 (cp_parser*
}
cp_ensure_no_omp_declare_simd (parser);
+ cp_ensure_no_oacc_routine (parser);
/* Issue an error message if type-definitions are forbidden here. */
cp_parser_check_type_definition (parser);
@@ -21782,7 +21723,7 @@ cp_parser_member_declaration (cp_parser*
}
cp_finalize_omp_declare_simd (parser, decl);
- cp_finalize_oacc_routine (parser, decl);
+ cp_finalize_oacc_routine (parser, decl, false);
/* Reset PREFIX_ATTRIBUTES. */
while (attributes && TREE_CHAIN (attributes) != first_attribute)
@@ -24390,8 +24331,7 @@ cp_parser_function_definition_from_speci
cp_finalize_omp_declare_simd (parser, current_function_decl);
parser->omp_declare_simd = NULL;
- cp_finalize_oacc_routine (parser, current_function_decl);
- parser->oacc_routine = NULL;
+ cp_finalize_oacc_routine (parser, current_function_decl, true);
}
if (!success_p)
@@ -25074,7 +25014,7 @@ cp_parser_save_member_function_body (cp_
/* Create the FUNCTION_DECL. */
fn = grokmethod (decl_specifiers, declarator, attributes);
cp_finalize_omp_declare_simd (parser, fn);
- cp_finalize_oacc_routine (parser, fn);
+ cp_finalize_oacc_routine (parser, fn, true);
/* If something went badly wrong, bail out now. */
if (fn == error_mark_node)
{
@@ -33693,41 +33633,29 @@ cp_parser_oacc_wait (cp_parser *parser,
static void
cp_parser_omp_declare_simd (cp_parser *parser, cp_token *pragma_tok,
- enum pragma_context context, bool is_omp)
+ enum pragma_context context)
{
- bool first_p = is_omp ? parser->omp_declare_simd == NULL
- : parser->oacc_routine == NULL;
+ bool first_p = parser->omp_declare_simd == NULL;
cp_omp_declare_simd_data data;
if (first_p)
{
data.error_seen = false;
data.fndecl_seen = false;
data.tokens = vNULL;
- if (is_omp)
- parser->omp_declare_simd = &data;
- else
- parser->oacc_routine = &data;
+ parser->omp_declare_simd = &data;
}
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))
- {
- if (is_omp)
- parser->omp_declare_simd->error_seen = true;
- else
- parser->oacc_routine->error_seen = true;
- }
+ parser->omp_declare_simd->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));
- if (is_omp)
- parser->omp_declare_simd->tokens.safe_push (cp);
- else
- parser->oacc_routine->tokens.safe_push (cp);
+ parser->omp_declare_simd->tokens.safe_push (cp);
if (first_p)
{
@@ -33748,23 +33676,14 @@ cp_parser_omp_declare_simd (cp_parser *p
cp_parser_declaration_statement (parser);
break;
}
- if (is_omp && parser->omp_declare_simd
+ if (parser->omp_declare_simd
&& !parser->omp_declare_simd->error_seen
&& !parser->omp_declare_simd->fndecl_seen)
error_at (pragma_tok->location,
"%<#pragma omp declare simd%> not immediately followed by "
"function declaration or definition");
- else if (!is_omp && parser->oacc_routine
- && !parser->oacc_routine->error_seen
- && !parser->oacc_routine->fndecl_seen)
- error_at (pragma_tok->location,
- "%<#pragma acc routine%> not immediately followed by "
- "function declaration or definition");
data.tokens.release ();
- if (is_omp)
- parser->omp_declare_simd = NULL;
- else
- parser->oacc_routine = NULL;
+ parser->omp_declare_simd = NULL;
}
}
@@ -34339,8 +34258,7 @@ cp_parser_omp_declare (cp_parser *parser
if (strcmp (p, "simd") == 0)
{
cp_lexer_consume_token (parser->lexer);
- cp_parser_omp_declare_simd (parser, pragma_tok,
- context, true);
+ cp_parser_omp_declare_simd (parser, pragma_tok, context);
return;
}
cp_ensure_no_omp_declare_simd (parser);
@@ -34391,42 +34309,103 @@ cp_parser_omp_declare (cp_parser *parser
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND))
+/* 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)
+{
+ location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses));
+
+ if (!fndecl)
+ {
+ error ("%<#pragma oacc routine%> not immediately followed by "
+ "function declaration or definition");
+ return;
+ }
+
+ if (named && is_overloaded_fn (fndecl)
+ && (TREE_CODE (fndecl) != FUNCTION_DECL
+ || DECL_FUNCTION_TEMPLATE_P (fndecl)))
+ {
+ error_at (loc, "%D names a set of overloads", OVL_CURRENT (fndecl));
+ return;
+ }
+
+ if (TREE_CODE (fndecl) != FUNCTION_DECL)
+ {
+ error_at (loc, "%D is not a function", fndecl);
+ return;
+ }
+
+ /* Perhaps we should use the same rule as declarations in different
+ namespaces? */
+ if (named && !DECL_NAMESPACE_SCOPE_P (fndecl))
+ {
+ error_at (loc, "%D is not at namespace scope", fndecl);
+ 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
+ function. The clauses are either processed directly (for a named
+ function), or defered until the immediatley following declaration
+ is parsed.
+*/
+
static void
cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
enum pragma_context context)
{
- tree name = NULL_TREE;
-
- //cp_lexer_consume_token (parser->lexer);
+ tree decl = NULL_TREE;
+ /* Create a dummy claue, to record location. */
+ tree c_head = build_omp_clause (pragma_tok->location, OMP_CLAUSE_SEQ);
- /* Scan for optional '( name )'. */
- if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN))
+ if (context != pragma_external)
+ cp_parser_error (parser, "%<#pragma acc routine%> not at file scope");
+
+ /* Look for optional '( name )'. */
+ if (cp_lexer_next_token_is (parser->lexer,CPP_OPEN_PAREN))
{
cp_lexer_consume_token (parser->lexer);
- name = cp_parser_id_expression (parser, /*template_p=*/false,
- /*check_dependency_p=*/true,
- /*template_p=*/NULL,
- /*declarator_p=*/false,
- /*optional_p=*/false);
- if (name == error_mark_node)
- return;
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
- if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_PAREN))
+ /* We parse the name as an id-expression. If it resolves to
+ anything other than a non-overloaded function at namespace
+ scope, it's an error. */
+ tree id = cp_parser_id_expression (parser,
+ /*template_keyword_p=*/false,
+ /*check_dependency_p=*/false,
+ /*template_p=*/NULL,
+ /*declarator_p=*/false,
+ /*optional_p=*/false);
+ decl = cp_parser_lookup_name_simple (parser, id, token->location);
+
+ if (decl == error_mark_node
+ || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
{
- error_at (cp_lexer_peek_token (parser->lexer)->location,
- "expected %<)%>");
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
return;
}
- cp_lexer_consume_token (parser->lexer);
- }
-
- /* If this routine construct doesn't explicitly have an optional 'name',
- then handle it the same way as an omp declare simd. */
- if (!name)
- {
- cp_parser_omp_declare_simd (parser, pragma_tok, context, false);
- cp_ensure_no_omp_declare_simd (parser);
- return;
}
/* Build a chain of clauses. */
@@ -34437,64 +34416,27 @@ cp_parser_oacc_routine (cp_parser *parse
cp_lexer_peek_token (parser->lexer),
OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK);
- TREE_CHAIN (name) = clauses;
- vec_safe_push (parser->named_oacc_routines, name);
+ /* Force clauses to be non-null, by attaching context to it. */
+ clauses = tree_cons (c_head, clauses, NULL_TREE);
+
+ if (decl)
+ cp_parser_finish_oacc_routine (parser, decl, clauses, true, false);
+ else
+ parser->oacc_routine = clauses;
}
-/* Finalize #pragma acc routine clauses after direct declarator has
- been parsed, and put that into "omp declare target" attribute. */
+/* Apply any saved OpenACC routine clauses to a just-parsed
+ declaration. */
-static tree
-cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
+static void
+cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
{
- struct cp_token_cache *ce;
- cp_omp_declare_simd_data *data = parser->oacc_routine;
- int i;
-
- if (!data->error_seen && data->fndecl_seen)
- {
- error ("%<#pragma acc routine%> not immediately followed by "
- "a single function declaration or definition");
- data->error_seen = true;
- return attrs;
- }
- if (data->error_seen)
- return attrs;
-
- tree c, cl = NULL_TREE;
-
- FOR_EACH_VEC_ELT (data->tokens, i, ce)
+ if (parser->oacc_routine)
{
- cp_parser_push_lexer_for_tokens (parser, ce);
- parser->lexer->in_pragma = true;
- gcc_assert (cp_lexer_peek_token (parser->lexer)->type == CPP_PRAGMA);
- cp_token *pragma_tok = cp_lexer_consume_token (parser->lexer);
- c = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
- "#pragma acc routine", pragma_tok);
- cp_parser_pop_lexer (parser);
-
- if (cl == NULL_TREE)
- cl = c;
- else if (c != NULL_TREE)
- {
- OMP_CLAUSE_CHAIN (c) = cl;
- cl = c;
- TREE_CHAIN (c) = attrs;
- if (processing_template_decl)
- ATTR_IS_DEPENDENT (c) = 1;
- attrs = c;
- }
+ cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine,
+ false, is_defn);
+ parser->oacc_routine = NULL_TREE;
}
-
- tree dims = build_oacc_routine_dims (cl);
- attrs = tree_cons (get_identifier ("oacc function"), dims, attrs);
-
- if (cl != NULL_TREE)
- cl = tree_cons (NULL_TREE, cl, NULL_TREE);
-
- attrs = tree_cons (get_identifier ("omp declare target"), cl, attrs);
- data->fndecl_seen = true;
- return attrs;
}
/* Main entry point to OpenMP statement pragmas. */
@@ -34979,6 +34921,7 @@ cp_parser_pragma (cp_parser *parser, enu
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:
===================================================================
@@ -373,10 +373,9 @@ typedef struct GTY(()) cp_parser {
necessary. */
cp_omp_declare_simd_data * GTY((skip)) cilk_simd_fn_info;
- /* OpenACC specific parser information. */
- cp_omp_declare_simd_data * GTY((skip)) oacc_routine;
- vec <tree, va_gc> *named_oacc_routines;
-
+ /* OpenACC routine clauses for subsequent decl/defn. */
+ tree oacc_routine;
+
/* Nonzero if parsing a parameter list where 'auto' should trigger an implicit
template parameter. */
bool auto_is_implicit_function_template_parm_p;
===================================================================
@@ -1,6 +1,8 @@
void *malloc (__SIZE_TYPE__);
void free (void *);
+int fact (int n);
+
#pragma acc routine (fact)
int
===================================================================
@@ -6,8 +6,9 @@
#include <stdio.h>
#include <stdlib.h>
-#pragma acc routine (fact)
+int fact (int);
+#pragma acc routine (fact)
int fact (int n)
{