Rework C/C++ OpenACC routine parsing (was: C/C++: Simplify handling of location information for OpenACC routine directives)
Thomas Schwinge
thomas@codesourcery.com
Wed Jul 13 14:11:00 GMT 2016
Hi!
On Wed, 13 Jul 2016 11:25:46 +0200, I wrote:
> Working on something else regarding the C/C++ OpenACC routine directive,
> I couldn't but untangle [...]
> (Another C/C++ OpenACC routine
> cleanup patch is emerging, depending on this one.)
Here it is; likewise, OK for trunk? (Further cleanup especially of C++
OpenACC routine handling seems to be possible, but I want to synchronize
my work at this point.)
commit 0bd30acaf4dd634499b1c695ddee555e7675aa18
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu Jun 23 13:28:09 2016 +0200
Rework C/C++ OpenACC routine parsing
gcc/c/
* c-parser.c (struct oacc_routine_data): Add error_seen and
fndecl_seen members.
(c_finish_oacc_routine): Use these.
(c_parser_declaration_or_fndef): Adjust.
(c_parser_oacc_routine): Likewise. Support more C language
constructs, and improve diagnostics. Move pragma context
checking...
(c_parser_pragma): ... here.
gcc/cp/
* parser.c (cp_ensure_no_oacc_routine): Improve diagnostics.
(cp_parser_late_parsing_cilk_simd_fn_info): Fix diagnostics.
(cp_parser_late_parsing_oacc_routine, cp_finalize_oacc_routine):
Simplify code, and improve diagnostics.
(cp_parser_oacc_routine): Likewise. Move pragma context
checking...
(cp_parser_pragma): ... here.
gcc/testsuite/
* c-c++-common/goacc/routine-5.c: Update.
---
gcc/c/c-parser.c | 161 +++++++++++++++-------
gcc/cp/parser.c | 182 +++++++++++-------------
gcc/testsuite/c-c++-common/goacc/routine-5.c | 199 +++++++++++++++++++++++----
3 files changed, 369 insertions(+), 173 deletions(-)
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 7f84ce9..809118a 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1273,6 +1273,8 @@ enum c_parser_prec {
/* Helper data structure for parsing #pragma acc routine. */
struct oacc_routine_data {
+ bool error_seen; /* Set if error has been reported. */
+ bool fndecl_seen; /* Set if one fn decl/definition has been seen already. */
tree clauses;
location_t loc;
};
@@ -1565,8 +1567,7 @@ c_parser_external_declaration (c_parser *parser)
}
static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>);
-static void c_finish_oacc_routine (struct oacc_routine_data *, tree, bool,
- bool, bool);
+static void c_finish_oacc_routine (struct oacc_routine_data *, 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
@@ -1751,8 +1752,7 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
}
c_parser_consume_token (parser);
if (oacc_routine_data)
- c_finish_oacc_routine (oacc_routine_data, NULL_TREE, false, true,
- false);
+ c_finish_oacc_routine (oacc_routine_data, NULL_TREE, false);
return;
}
@@ -1850,7 +1850,7 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
prefix_attrs = specs->attrs;
all_prefix_attrs = prefix_attrs;
specs->attrs = NULL_TREE;
- for (bool first = true;; first = false)
+ while (true)
{
struct c_declarator *declarator;
bool dummy = false;
@@ -1870,8 +1870,7 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
c_finish_omp_declare_simd (parser, NULL_TREE, NULL_TREE,
omp_declare_simd_clauses);
if (oacc_routine_data)
- c_finish_oacc_routine (oacc_routine_data, NULL_TREE,
- false, first, false);
+ c_finish_oacc_routine (oacc_routine_data, NULL_TREE, false);
c_parser_skip_to_end_of_block_or_statement (parser);
return;
}
@@ -1987,8 +1986,7 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
finish_init ();
}
if (oacc_routine_data)
- c_finish_oacc_routine (oacc_routine_data, d,
- false, first, false);
+ c_finish_oacc_routine (oacc_routine_data, d, false);
if (d != error_mark_node)
{
maybe_warn_string_init (init_loc, TREE_TYPE (d), init);
@@ -2033,8 +2031,7 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
temp_pop_parm_decls ();
}
if (oacc_routine_data)
- c_finish_oacc_routine (oacc_routine_data, d,
- false, first, false);
+ c_finish_oacc_routine (oacc_routine_data, d, false);
if (d)
finish_decl (d, UNKNOWN_LOCATION, NULL_TREE,
NULL_TREE, asm_name);
@@ -2146,8 +2143,7 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
c_finish_omp_declare_simd (parser, current_function_decl, NULL_TREE,
omp_declare_simd_clauses);
if (oacc_routine_data)
- c_finish_oacc_routine (oacc_routine_data, current_function_decl,
- false, first, true);
+ c_finish_oacc_routine (oacc_routine_data, current_function_decl, true);
DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus
= c_parser_peek_token (parser)->location;
fnbody = c_parser_compound_statement (parser);
@@ -10117,6 +10113,13 @@ c_parser_pragma (c_parser *parser, enum pragma_context context, bool *if_p)
return false;
case PRAGMA_OACC_ROUTINE:
+ if (context != pragma_external)
+ {
+ error_at (c_parser_peek_token (parser)->location,
+ "%<#pragma acc routine%> must be at file scope");
+ c_parser_skip_until_found (parser, CPP_PRAGMA_EOL, NULL);
+ return false;
+ }
c_parser_oacc_routine (parser, context);
return false;
@@ -14029,29 +14032,32 @@ c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser,
static void
c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
{
- tree decl = NULL_TREE;
+ gcc_checking_assert (context == pragma_external);
+
oacc_routine_data data;
+ data.error_seen = false;
+ data.fndecl_seen = false;
data.clauses = NULL_TREE;
data.loc = c_parser_peek_token (parser)->location;
-
- if (context != pragma_external)
- c_parser_error (parser, "%<#pragma acc routine%> not at file scope");
c_parser_consume_pragma (parser);
- /* Scan for optional '( name )'. */
+ /* Look for optional '( name )'. */
if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
{
- c_parser_consume_token (parser);
+ c_parser_consume_token (parser); /* '(' */
- c_token *token = c_parser_peek_token (parser);
- if (token->type == CPP_NAME && (token->id_kind == C_ID_ID
- || token->id_kind == C_ID_TYPENAME))
+ tree decl = NULL_TREE;
+ c_token *name_token = c_parser_peek_token (parser);
+ location_t name_loc = name_token->location;
+ if (name_token->type == CPP_NAME
+ && (name_token->id_kind == C_ID_ID
+ || name_token->id_kind == C_ID_TYPENAME))
{
- decl = lookup_name (token->value);
+ decl = lookup_name (name_token->value);
if (!decl)
- error_at (token->location, "%qE has not been declared",
- token->value);
+ error_at (name_loc,
+ "%qE has not been declared", name_token->value);
c_parser_consume_token (parser);
}
else
@@ -14063,22 +14069,56 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
c_parser_skip_to_pragma_eol (parser, false);
return;
}
+
+ data.clauses
+ = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
+ "#pragma acc routine");
+
+ if (TREE_CODE (decl) != FUNCTION_DECL)
+ {
+ error_at (name_loc, "%qD does not refer to a function", decl);
+ return;
+ }
+
+ c_finish_oacc_routine (&data, decl, false);
}
+ else /* No optional '( name )'. */
+ {
+ data.clauses
+ = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
+ "#pragma acc routine");
- /* Build a chain of clauses. */
- parser->in_pragma = true;
- data.clauses
- = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
- "#pragma acc routine");
+ /* Emit a helpful diagnostic if there's another pragma following this
+ one. Also don't allow a static assertion declaration, as in the
+ following we'll just parse a *single* "declaration or function
+ definition", and the static assertion counts an one. */
+ if (c_parser_next_token_is (parser, CPP_PRAGMA)
+ || c_parser_next_token_is_keyword (parser, RID_STATIC_ASSERT))
+ {
+ error_at (data.loc,
+ "%<#pragma acc routine%> not immediately followed by"
+ " function declaration or definition");
+ /* ..., and then just keep going. */
+ return;
+ }
- if (decl)
- c_finish_oacc_routine (&data, decl, true, true, false);
- else if (c_parser_peek_token (parser)->type == CPP_PRAGMA)
- /* This will emit an error. */
- c_finish_oacc_routine (&data, NULL_TREE, false, true, false);
- else
- c_parser_declaration_or_fndef (parser, true, false, false, false,
- true, NULL, vNULL, &data);
+ /* We only have to consider the pragma_external case here. */
+ if (c_parser_next_token_is (parser, CPP_KEYWORD)
+ && c_parser_peek_token (parser)->keyword == RID_EXTENSION)
+ {
+ int ext = disable_extension_diagnostics ();
+ do
+ c_parser_consume_token (parser);
+ while (c_parser_next_token_is (parser, CPP_KEYWORD)
+ && c_parser_peek_token (parser)->keyword == RID_EXTENSION);
+ c_parser_declaration_or_fndef (parser, true, true, true, false, true,
+ NULL, vNULL, &data);
+ restore_extension_diagnostics (ext);
+ }
+ else
+ c_parser_declaration_or_fndef (parser, true, true, true, false, true,
+ NULL, vNULL, &data);
+ }
}
/* Finalize an OpenACC routine pragma, applying it to FNDECL.
@@ -14086,24 +14126,46 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
static void
c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
- bool named, bool first, bool is_defn)
+ bool is_defn)
{
- if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL || !first)
+ /* Keep going if we're in error reporting mode. */
+ if (data->error_seen
+ || fndecl == error_mark_node)
+ return;
+
+ if (data->fndecl_seen)
{
- if (fndecl != error_mark_node)
- error_at (data->loc, "%<#pragma acc routine%> %s",
- named ? "does not refer to a function"
- : "not followed by single function");
+ error_at (data->loc,
+ "%<#pragma acc routine%> not immediately followed by"
+ " a single function declaration or definition");
+ data->error_seen = true;
+ return;
+ }
+ if (fndecl == NULL_TREE || TREE_CODE (fndecl) != FUNCTION_DECL)
+ {
+ error_at (data->loc,
+ "%<#pragma acc routine%> not immediately followed by"
+ " function declaration or definition");
+ data->error_seen = true;
return;
}
if (get_oacc_fn_attrib (fndecl))
- error_at (data->loc,
- "%<#pragma acc routine%> already applied to %D", fndecl);
+ {
+ error_at (data->loc,
+ "%<#pragma acc routine%> already applied to %qD", fndecl);
+ data->error_seen = true;
+ return;
+ }
if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
- error_at (data->loc, "%<#pragma acc routine%> must be applied before %s",
- TREE_USED (fndecl) ? "use" : "definition");
+ {
+ error_at (data->loc,
+ "%<#pragma acc routine%> must be applied before %s",
+ TREE_USED (fndecl) ? "use" : "definition");
+ data->error_seen = true;
+ return;
+ }
/* Process the routine's dimension clauses. */
tree dims = build_oacc_routine_dims (data->clauses);
@@ -14113,6 +14175,9 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
DECL_ATTRIBUTES (fndecl)
= tree_cons (get_identifier ("omp declare target"),
NULL_TREE, DECL_ATTRIBUTES (fndecl));
+
+ /* Remember that we've used this "#pragma acc routine". */
+ data->fndecl_seen = true;
}
/* OpenACC 2.0:
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 28417d8..43bed89 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -1383,8 +1383,8 @@ cp_ensure_no_oacc_routine (cp_parser *parser)
if (parser->oacc_routine && !parser->oacc_routine->error_seen)
{
error_at (parser->oacc_routine->loc,
- "%<#pragma acc routine%> not followed by a function "
- "declaration or definition");
+ "%<#pragma acc routine%> not immediately followed by "
+ "function declaration or definition");
parser->oacc_routine = NULL;
}
}
@@ -35660,6 +35660,8 @@ cp_parser_omp_declare_simd (cp_parser *parser, cp_token *pragma_tok,
used while this scope is live. */
parser->omp_declare_simd = &data;
}
+
+ /* Store away all pragma tokens. */
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);
@@ -35669,6 +35671,7 @@ cp_parser_omp_declare_simd (cp_parser *parser, cp_token *pragma_tok,
struct cp_token_cache *cp
= cp_token_cache_new (pragma_tok, cp_lexer_peek_token (parser->lexer));
parser->omp_declare_simd->tokens.safe_push (cp);
+
if (first_p)
{
while (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA))
@@ -35713,9 +35716,9 @@ cp_parser_late_parsing_cilk_simd_fn_info (cp_parser *parser, tree attrs)
if (parser->omp_declare_simd != NULL
|| lookup_attribute ("simd", attrs))
{
- error ("%<#pragma omp declare simd%> of %<simd%> attribute cannot be "
+ error ("%<#pragma omp declare simd%> or %<simd%> attribute cannot be "
"used in the same function marked as a Cilk Plus SIMD-enabled "
- " function");
+ "function");
parser->cilk_simd_fn_info->tokens.release ();
XDELETE (parser->cilk_simd_fn_info);
parser->cilk_simd_fn_info = NULL;
@@ -36487,64 +36490,39 @@ static void
cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
enum pragma_context context)
{
- bool first_p = parser->oacc_routine == NULL;
+ gcc_checking_assert (context == pragma_external);
+ /* The checking for "another pragma following this one" in the "no optional
+ '( name )'" case makes sure that we dont re-enter. */
+ gcc_checking_assert (parser->oacc_routine == NULL);
+
cp_oacc_routine_data data;
- if (first_p)
- {
- data.error_seen = false;
- data.fndecl_seen = false;
- data.tokens = vNULL;
- data.clauses = NULL_TREE;
- data.loc = pragma_tok->location;
- /* It is safe to take the address of a local variable; it will only be
- used while this scope is live. */
- parser->oacc_routine = &data;
- }
+ data.error_seen = false;
+ data.fndecl_seen = false;
+ data.tokens = vNULL;
+ data.clauses = NULL_TREE;
+ data.loc = pragma_tok->location;
+ /* It is safe to take the address of a local variable; it will only be
+ used while this scope is live. */
+ parser->oacc_routine = &data;
- if (context != pragma_external)
- {
- 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 )'. */
+ /* Do we have an optional '( name )'? */
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 (parser->oacc_routine->loc,
- "%<#pragma acc routine%> not followed by a "
- "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);
+ cp_lexer_consume_token (parser->lexer); /* '(' */
/* 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);
- tree decl = cp_parser_lookup_name_simple (parser, id, token->location);
- if (id != error_mark_node && decl == error_mark_node)
- cp_parser_name_lookup_error (parser, id, decl, NLE_NULL,
- token->location);
+ location_t name_loc = cp_lexer_peek_token (parser->lexer)->location;
+ tree name = cp_parser_id_expression (parser,
+ /*template_keyword_p=*/false,
+ /*check_dependency_p=*/false,
+ /*template_p=*/NULL,
+ /*declarator_p=*/false,
+ /*optional_p=*/false);
+ tree decl = cp_parser_lookup_name_simple (parser, name, name_loc);
+ if (name != error_mark_node && decl == error_mark_node)
+ cp_parser_name_lookup_error (parser, name, decl, NLE_NULL, name_loc);
if (decl == error_mark_node
|| !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
@@ -36554,8 +36532,6 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
return;
}
- /* Build a chain of clauses. */
- parser->lexer->in_pragma = true;
data.clauses
= cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
"#pragma acc routine",
@@ -36565,7 +36541,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
&& (TREE_CODE (decl) != FUNCTION_DECL
|| DECL_FUNCTION_TEMPLATE_P (decl)))
{
- error_at (data.loc,
+ error_at (name_loc,
"%<#pragma acc routine%> names a set of overloads");
parser->oacc_routine = NULL;
return;
@@ -36575,60 +36551,58 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
namespaces? */
if (!DECL_NAMESPACE_SCOPE_P (decl))
{
- error_at (data.loc,
- "%<#pragma acc routine%> does not refer to a "
- "namespace scope function");
+ error_at (name_loc,
+ "%qD does not refer to a namespace scope function", decl);
parser->oacc_routine = NULL;
return;
}
- if (!decl || TREE_CODE (decl) != FUNCTION_DECL)
+ if (TREE_CODE (decl) != FUNCTION_DECL)
{
- error_at (data.loc,
- "%<#pragma acc routine%> does not refer to a function");
+ error_at (name_loc, "%qD does not refer to a function", decl);
parser->oacc_routine = NULL;
return;
}
cp_finalize_oacc_routine (parser, decl, false);
- data.tokens.release ();
parser->oacc_routine = NULL;
}
- else
+ else /* No optional '( name )'. */
{
+ /* Store away all pragma tokens. */
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);
- while (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA))
- cp_parser_pragma (parser, context, NULL);
-
- if (first_p)
+ /* Emit a helpful diagnostic if there's another pragma following this
+ one. */
+ if (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA))
{
- cp_parser_declaration (parser);
-
- if (parser->oacc_routine
- && !parser->oacc_routine->error_seen
- && !parser->oacc_routine->fndecl_seen)
- error_at (data.loc,
- "%<#pragma acc routine%> not followed by a "
- "function declaration or definition");
-
+ cp_ensure_no_oacc_routine (parser);
data.tokens.release ();
- parser->oacc_routine = NULL;
+ /* ..., and then just keep going. */
+ return;
}
+
+ /* We only have to consider the pragma_external case here. */
+ cp_parser_declaration (parser);
+ if (parser->oacc_routine
+ && !parser->oacc_routine->fndecl_seen)
+ cp_ensure_no_oacc_routine (parser);
+ else
+ parser->oacc_routine = NULL;
+ data.tokens.release ();
}
}
/* Finalize #pragma acc routine clauses after direct declarator has
- been parsed, and put that into "oacc function" attribute. */
+ been parsed. */
static tree
cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
@@ -36636,17 +36610,17 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
struct cp_token_cache *ce;
cp_oacc_routine_data *data = parser->oacc_routine;
- if ((!data->error_seen && data->fndecl_seen)
- || data->tokens.length () != 1)
+ if (!data->error_seen && data->fndecl_seen)
{
error_at (data->loc,
- "%<#pragma acc routine%> not followed by a "
- "function declaration or definition");
+ "%<#pragma acc routine%> not immediately followed by "
+ "a single function declaration or definition");
data->error_seen = true;
}
if (data->error_seen)
return attrs;
+ gcc_checking_assert (data->tokens.length () == 1);
ce = data->tokens[0];
cp_parser_push_lexer_for_tokens (parser, ce);
@@ -36654,12 +36628,14 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
gcc_assert (cp_lexer_peek_token (parser->lexer)->type == CPP_PRAGMA);
cp_token *pragma_tok = cp_lexer_consume_token (parser->lexer);
+ gcc_checking_assert (parser->oacc_routine->clauses == NULL_TREE);
parser->oacc_routine->clauses
= cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
"#pragma acc routine", pragma_tok);
cp_parser_pop_lexer (parser);
+ /* Later, cp_finalize_oacc_routine will process the clauses, and then set
+ fndecl_seen. */
- data->fndecl_seen = true;
return attrs;
}
@@ -36671,34 +36647,29 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
{
if (__builtin_expect (parser->oacc_routine != NULL, 0))
{
- if (parser->oacc_routine->error_seen)
+ /* Keep going if we're in error reporting mode. */
+ if (parser->oacc_routine->error_seen
+ || fndecl == error_mark_node)
return;
-
- if (fndecl == error_mark_node)
+
+ if (parser->oacc_routine->fndecl_seen)
{
+ error_at (parser->oacc_routine->loc,
+ "%<#pragma acc routine%> not immediately followed by"
+ " a single function declaration or definition");
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 (parser->oacc_routine->loc,
- "%<#pragma acc routine%> not followed by a function "
- "declaration or definition");
- parser->oacc_routine = NULL;
- return;
- }
-
if (get_oacc_fn_attrib (fndecl))
{
error_at (parser->oacc_routine->loc,
- "%<#pragma acc routine%> already applied to %D", fndecl);
+ "%<#pragma acc routine%> already applied to %qD", fndecl);
parser->oacc_routine = NULL;
return;
}
@@ -36720,6 +36691,11 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
DECL_ATTRIBUTES (fndecl)
= tree_cons (get_identifier ("omp declare target"),
NULL_TREE, DECL_ATTRIBUTES (fndecl));
+
+ /* Don't unset parser->oacc_routine here: we may still need it to
+ diagnose wrong usage. But, remember that we've used this "#pragma acc
+ routine". */
+ parser->oacc_routine->fndecl_seen = true;
}
}
@@ -37319,6 +37295,12 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p)
return false;
case PRAGMA_OACC_ROUTINE:
+ if (context != pragma_external)
+ {
+ error_at (pragma_tok->location,
+ "%<#pragma acc routine%> must be at file scope");
+ break;
+ }
cp_parser_oacc_routine (parser, pragma_tok, context);
return false;
diff --git gcc/testsuite/c-c++-common/goacc/routine-5.c gcc/testsuite/c-c++-common/goacc/routine-5.c
index 1efd154..17fe67c 100644
--- gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -1,64 +1,211 @@
-/* { dg-do compile } */
+/* Miscellaneous OpenACC routine front end checking. */
-#pragma acc routine /* { dg-error "not followed by" } */
+/* Pragma context. */
+
+struct PC
+{
+#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+};
+
+void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c++ } } */
+#pragma acc routine
+ /* { dg-error ".#pragma acc routine. must be at file scope" "" { target c } 11 }
+ { dg-error ".#pragma. is not allowed here" "" { target c++ } 11 } */
+) /* { dg-bogus "expected declaration specifiers or .\\.\\.\\.. before .\\). token" "TODO" { xfail c } } */
+{
+}
+
+void PC2()
+{
+ if (0)
+#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+ ;
+}
+
+void PC3()
+{
+#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+}
+
+
+/* "( name )" syntax. */
+
+#pragma acc routine ( /* { dg-error "expected (function name|unqualified-id) before end of line" } */
+#pragma acc routine () /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */
+#pragma acc routine (+) /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */
+#pragma acc routine (?) /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */
+#pragma acc routine (:) /* { dg-error "expected (function name|unqualified-id) before .:. token" } */
+#pragma acc routine (4) /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */
+#pragma acc routine ('4') /* { dg-error "expected (function name|unqualified-id) before .4." } */
+#pragma acc routine ("4") /* { dg-error "expected (function name|unqualified-id) before string constant" } */
+extern void R1(void);
+extern void R2(void);
+#pragma acc routine (R1, R2, R3) worker /* { dg-error "expected .\\). before .,. token" } */
+#pragma acc routine (R1 R2 R3) worker /* { dg-error "expected .\\). before .R2." } */
+#pragma acc routine (R1) worker
+#pragma acc routine (R2) worker
+
+
+/* "#pragma acc routine" not immediately followed by (a single) function
+ declaration or definition. */
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
int a;
-#pragma acc routine /* { dg-error "not followed by" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
void fn1 (void), fn1b (void);
-#pragma acc routine /* { dg-error "not followed by" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
int b, fn2 (void);
-#pragma acc routine /* { dg-error "not followed by" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+int b_, fn2_ (void), B_;
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
int fn3 (void), b2;
-#pragma acc routine /* { dg-error "not followed by" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
typedef struct c c;
-#pragma acc routine /* { dg-error "not followed by" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
struct d {} d;
-#pragma acc routine /* { dg-error "not followed by" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+void fn1_2 (void), fn1b_2 (void);
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+int b_2, fn2_2 (void);
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+int b_2_, fn2_2_ (void), B_2_;
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+int fn3_2 (void), b2_2;
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+typedef struct c_2 c_2;
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+struct d_2 {} d_2;
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
#pragma acc routine
int fn4 (void);
int fn5a (void);
-
-#pragma acc routine /* { dg-error "not followed by" } */
+int fn5b (void);
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
#pragma acc routine (fn5a)
+#pragma acc routine (fn5b)
int fn5 (void);
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine (fn6a) /* { dg-error ".fn6a. has not been declared" } */
+#pragma acc routine (fn6b) /* { dg-error ".fn6b. has not been declared" } */
+int fn6 (void);
+
#ifdef __cplusplus
-#pragma acc routine /* { dg-error "not followed by" "" { target c++ } } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
namespace f {}
namespace g {}
-#pragma acc routine /* { dg-error "not followed by" "" { target c++ } } */
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
using namespace g;
-#pragma acc routine (g) /* { dg-error "does not refer to a function" "" { target c++ } } */
+#pragma acc routine (g) /* { dg-error ".g. does not refer to a function" "" { target c++ } } */
#endif /* __cplusplus */
-#pragma acc routine (a) /* { dg-error "does not refer to a function" } */
+#pragma acc routine (a) /* { dg-error ".a. does not refer to a function" } */
-#pragma acc routine (c) /* { dg-error "does not refer to a function" } */
+#pragma acc routine (c) /* { dg-error ".c. does not refer to a function" } */
-#pragma acc routine () vector /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */
+/* Static assert. */
-#pragma acc routine (+) /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */
+#pragma acc routine /* { dg-bogus ".#pragma acc routine. not immediately followed by function declaration or definition" "TODO" { xfail *-*-* } } */
+#ifndef __cplusplus /* C */
+_Static_assert(0, ""); /* { dg-error "static assertion failed" "" { target c } } */
+#elif __cplusplus < 201103L /* C++98 */
+/* C++98 doesn't support static_assert, so fake an error in combination, and as
+ expected with the "#pragma acc routine" above. */
+int dummy_instead_of_static_assert;
+#else /* C++ */
+static_assert(0, ""); /* { dg-error "static assertion failed" "" { target c++11 } } */
+#endif
+void f_static_assert();
+/* Check that we already recognized "f_static_assert" as an OpenACC routine. */
+#pragma acc routine (f_static_assert) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */
-extern void R1(void);
-extern void R2(void);
-#pragma acc routine (R1, R2, R3) worker /* { dg-error "expected .\\). before .,. token" } */
-#pragma acc routine (R1 R2 R3) worker /* { dg-error "expected .\\). before .R2." } */
-#pragma acc routine (R1) worker
-#pragma acc routine (R2) worker
+/* __extension__ usage. */
+#pragma acc routine
+__extension__ extern void ex1();
+#pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */
+
+#pragma acc routine
+__extension__ __extension__ __extension__ __extension__ __extension__ void ex2()
+{
+}
+#pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */
+
+#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+__extension__ int ex3;
+#pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */
+
+
+/* "#pragma acc routine" already applied. */
+
+extern void fungsi_1();
+#pragma acc routine(fungsi_1) gang
+#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
+#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
+#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
+
+#pragma acc routine seq
+extern void fungsi_2();
+#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
+#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
+#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
+
+#pragma acc routine vector
+extern void fungsi_3();
+#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */
+void fungsi_3()
+{
+}
+
+extern void fungsi_4();
+#pragma acc routine (fungsi_4) worker
+#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */
+void fungsi_4()
+{
+}
+
+#pragma acc routine gang
+void fungsi_5()
+{
+}
+#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */
+
+#pragma acc routine seq
+void fungsi_6()
+{
+}
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */
+extern void fungsi_6();
+
+
+/* "#pragma acc routine" must be applied before. */
void Bar ();
@@ -67,13 +214,15 @@ void Foo ()
Bar ();
}
-#pragma acc routine (Bar) // { dg-error "must be applied before use" }
+#pragma acc routine (Bar) // { dg-error ".#pragma acc routine. must be applied before use" }
-#pragma acc routine (Foo) gang // { dg-error "must be applied before definition" }
+#pragma acc routine (Foo) gang // { dg-error ".#pragma acc routine. must be applied before definition" }
#pragma acc routine (Baz) // { dg-error "not been declared" }
+/* OpenACC declare. */
+
int vb1; /* { dg-error "directive for use" } */
extern int vb2; /* { dg-error "directive for use" } */
static int vb3; /* { dg-error "directive for use" } */
Grüße
Thomas
More information about the Gcc-patches
mailing list