This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [1/2] OpenACC routine support
- From: Nathan Sidwell <nathan at acm dot org>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: GCC Patches <gcc-patches at gcc dot gnu dot org>, Cesar Philippidis <cesar at codesourcery dot com>
- Date: Mon, 9 Nov 2015 19:31:45 -0500
- Subject: Re: [1/2] OpenACC routine support
- Authentication-results: sourceware.org; auth=none
- References: <5637B1CF dot 5060408 at acm dot org> <5637B7C7 dot 70901 at acm dot org> <20151103153533 dot GQ478 at tucnak dot redhat dot com>
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.
nathan
2015-11-09 Nathan Sidwell <nathan@codesourcery.com>
* omp-low.h (replace_oacc_fn_attrib, build_oacc_routine_dims): Declare.
* omp-low.c (build_oacc_routine_dims): New.
2015-11-09 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
James Norris <jnorris@codesourcery.com>
Julian Brown <julian@codesourcery.com>
Nathan Sidwell <nathan@codesourcery.com>
c/
* c-parser.c (c_parser_declaration_or_fndef): Add OpenACC
routine arg.
(c_parser_declaration_or_fndef): Call c_finish_oacc_routine.
(c_parser_pragma): Parse 'acc routine'.
(OACC_ROUTINE_CLAUSE_MARK): Define.
(c_parser_oacc_routine, (c_finish_oacc_routine): New.
2015-11-09 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
James Norris <jnorris@codesourcery.com>
Julian Brown <julian@codesourcery.com>
Nathan Sidwell <nathan@codesourcery.com>
c-family/
* c-pragma.c (oacc_pragmas): Add "routine".
* c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE.
2015-11-09 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
James Norris <jnorris@codesourcery.com>
Julian Brown <julian@codesourcery.com>
Nathan Sidwell <nathan@codesourcery.com>
cp/
* parser.h (struct cp_parser): Add oacc_routine field.
* parser.c (cp_ensure_no_oacc_routine): New.
(cp_parser_new): Initialize oacc_routine field.
(cp_parser_linkage_specification): Call cp_ensure_no_oacc_routine.
(cp_parser_namespace_definition,
cp_parser_class_specifier_1): Likewise.
(cp_parser_init_declarator): Call cp_finalize_oacc_routine.
(cp_parser_function_definition,
cp_parser_save_member_function_body): Likewise.
(OACC_ROUTINE_CLAUSE_MASK): New.
(cp_parser_finish_oacc_routine, cp_parser_oacc_routine,
cp_finalize_oacc_routine): New.
(cp_parser_pragma): Adjust omp_declare_simd checking. Call
cp_ensure_no_oacc_routine.
(cp_parser_pragma): Add OpenACC routine handling.
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c (revision 230040)
+++ gcc/omp-low.c (working copy)
@@ -12361,6 +12361,50 @@ set_oacc_fn_attrib (tree fn, tree clause
}
}
+/* Process the routine's dimension clauess to generate an attribute
+ value. Issue diagnostics as appropriate. We default to SEQ
+ (OpenACC 2.5 clarifies this). All dimensions have a size of zero
+ (dynamic). TREE_PURPOSE is set to indicate whether that dimension
+ can have a loop partitioned on it. non-zero indicates
+ yes, zero indicates no. By construction once a non-zero has been
+ reached, further inner dimensions must also be non-zero. We set
+ TREE_VALUE to zero for the dimensions that may be partitioned and
+ 1 for the other ones -- if a loop is (erroneously) spawned at
+ an outer level, we don't want to try and partition it. */
+
+tree
+build_oacc_routine_dims (tree clauses)
+{
+ /* Must match GOMP_DIM ordering. */
+ static const omp_clause_code ids[] =
+ {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
+ int ix;
+ int level = -1;
+
+ for (; clauses; clauses = OMP_CLAUSE_CHAIN (clauses))
+ for (ix = GOMP_DIM_MAX + 1; ix--;)
+ if (OMP_CLAUSE_CODE (clauses) == ids[ix])
+ {
+ if (level >= 0)
+ error_at (OMP_CLAUSE_LOCATION (clauses),
+ "multiple loop axes specified for routine");
+ level = ix;
+ break;
+ }
+
+ /* Default to SEQ. */
+ if (level < 0)
+ level = GOMP_DIM_MAX;
+
+ tree dims = NULL_TREE;
+
+ for (ix = GOMP_DIM_MAX; ix--;)
+ dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
+ build_int_cst (integer_type_node, ix < level), dims);
+
+ return dims;
+}
+
/* Retrieve the oacc function attrib and return it. Non-oacc
functions will return NULL. */
Index: gcc/omp-low.h
===================================================================
--- gcc/omp-low.h (revision 230040)
+++ gcc/omp-low.h (working copy)
@@ -30,6 +30,8 @@ extern tree omp_reduction_init (tree, tr
extern bool make_gimple_omp_edges (basic_block, struct omp_region **, int *);
extern void omp_finish_file (void);
extern tree omp_member_access_dummy_var (tree);
+extern void replace_oacc_fn_attrib (tree, tree);
+extern tree build_oacc_routine_dims (tree);
extern tree get_oacc_fn_attrib (tree);
extern int get_oacc_ifn_dim_arg (const gimple *);
extern int get_oacc_fn_dim_size (tree, int);
Index: gcc/c-family/c-pragma.c
===================================================================
--- gcc/c-family/c-pragma.c (revision 230040)
+++ gcc/c-family/c-pragma.c (working copy)
@@ -1257,6 +1257,7 @@ static const struct omp_pragma_def oacc_
{ "kernels", PRAGMA_OACC_KERNELS },
{ "loop", PRAGMA_OACC_LOOP },
{ "parallel", PRAGMA_OACC_PARALLEL },
+ { "routine", PRAGMA_OACC_ROUTINE },
{ "update", PRAGMA_OACC_UPDATE },
{ "wait", PRAGMA_OACC_WAIT }
};
Index: gcc/c-family/c-pragma.h
===================================================================
--- gcc/c-family/c-pragma.h (revision 230040)
+++ gcc/c-family/c-pragma.h (working copy)
@@ -35,6 +35,7 @@ enum pragma_kind {
PRAGMA_OACC_KERNELS,
PRAGMA_OACC_LOOP,
PRAGMA_OACC_PARALLEL,
+ PRAGMA_OACC_ROUTINE,
PRAGMA_OACC_UPDATE,
PRAGMA_OACC_WAIT,
Index: gcc/c/c-parser.c
===================================================================
--- gcc/c/c-parser.c (revision 230040)
+++ gcc/c/c-parser.c (working copy)
@@ -1162,7 +1162,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 = NULL_TREE);
static void c_parser_static_assert_declaration_no_semi (c_parser *);
static void c_parser_static_assert_declaration (c_parser *);
static void c_parser_declspecs (c_parser *, struct c_declspecs *, bool, bool,
@@ -1251,6 +1252,7 @@ static bool c_parser_omp_target (c_parse
static void c_parser_omp_end_declare_target (c_parser *);
static void c_parser_omp_declare (c_parser *, enum pragma_context);
static bool c_parser_omp_ordered (c_parser *, enum pragma_context);
+static void c_parser_oacc_routine (c_parser *parser, enum pragma_context);
/* These Objective-C parser functions are only ever called when
compiling Objective-C. */
@@ -1436,6 +1438,7 @@ c_parser_external_declaration (c_parser
}
static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>);
+static void c_finish_oacc_routine (c_parser *, tree, tree, bool, bool, 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
@@ -1513,7 +1516,8 @@ c_parser_declaration_or_fndef (c_parser
bool static_assert_ok, bool empty_ok,
bool nested, bool start_attr_ok,
tree *objc_foreach_object_declaration,
- vec<c_token> omp_declare_simd_clauses)
+ vec<c_token> omp_declare_simd_clauses,
+ tree oacc_routine_clauses)
{
struct c_declspecs *specs;
tree prefix_attrs;
@@ -1583,6 +1587,9 @@ c_parser_declaration_or_fndef (c_parser
pedwarn (here, 0, "empty declaration");
}
c_parser_consume_token (parser);
+ if (oacc_routine_clauses)
+ c_finish_oacc_routine (parser, NULL_TREE,
+ oacc_routine_clauses, false, true, false);
return;
}
@@ -1680,7 +1687,7 @@ c_parser_declaration_or_fndef (c_parser
prefix_attrs = specs->attrs;
all_prefix_attrs = prefix_attrs;
specs->attrs = NULL_TREE;
- while (true)
+ for (bool first = true;; first = false)
{
struct c_declarator *declarator;
bool dummy = false;
@@ -1699,6 +1706,10 @@ c_parser_declaration_or_fndef (c_parser
|| !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
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,
+ false, first, false);
c_parser_skip_to_end_of_block_or_statement (parser);
return;
}
@@ -1813,6 +1824,9 @@ c_parser_declaration_or_fndef (c_parser
init = c_parser_initializer (parser);
finish_init ();
}
+ if (oacc_routine_clauses)
+ c_finish_oacc_routine (parser, d, oacc_routine_clauses,
+ false, first, false);
if (d != error_mark_node)
{
maybe_warn_string_init (init_loc, TREE_TYPE (d), init);
@@ -1856,6 +1870,9 @@ c_parser_declaration_or_fndef (c_parser
if (parms)
temp_pop_parm_decls ();
}
+ if (oacc_routine_clauses)
+ c_finish_oacc_routine (parser, d, oacc_routine_clauses,
+ false, first, false);
if (d)
finish_decl (d, UNKNOWN_LOCATION, NULL_TREE,
NULL_TREE, asm_name);
@@ -1966,6 +1983,9 @@ 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, false, first, true);
DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus
= c_parser_peek_token (parser)->location;
fnbody = c_parser_compound_statement (parser);
@@ -9706,6 +9726,10 @@ c_parser_pragma (c_parser *parser, enum
c_parser_oacc_enter_exit_data (parser, false);
return false;
+ case PRAGMA_OACC_ROUTINE:
+ c_parser_oacc_routine (parser, context);
+ return false;
+
case PRAGMA_OACC_UPDATE:
if (context != pragma_compound)
{
@@ -13400,6 +13424,117 @@ c_parser_oacc_kernels_parallel (location
}
/* OpenACC 2.0:
+ # pragma acc routine oacc-routine-clause[optseq] new-line
+ function-definition
+
+ # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line
+*/
+
+#define OACC_ROUTINE_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) )
+
+/* 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)
+{
+ tree decl = NULL_TREE;
+ /* Create a dummy claue, to record location. */
+ tree c_head = build_omp_clause (c_parser_peek_token (parser)->location,
+ OMP_CLAUSE_SEQ);
+
+ if (context != pragma_external)
+ c_parser_error (parser, "%<#pragma acc routine%> not at file scope");
+
+ c_parser_consume_pragma (parser);
+
+ /* Scan for optional '( name )'. */
+ if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+ {
+ 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))
+ {
+ decl = lookup_name (token->value);
+ if (!decl)
+ {
+ error_at (token->location, "%qE has not been declared",
+ token->value);
+ decl = error_mark_node;
+ }
+ }
+ else
+ c_parser_error (parser, "expected function name");
+
+ if (token->type != CPP_CLOSE_PAREN)
+ c_parser_consume_token (parser);
+
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
+ }
+
+ /* Build a chain of clauses. */
+ parser->in_pragma = true;
+ tree clauses = c_parser_oacc_all_clauses
+ (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine");
+
+ /* Force clauses to be non-null, by attaching context to it. */
+ clauses = tree_cons (c_head, clauses, NULL_TREE);
+
+ if (decl)
+ c_finish_oacc_routine (parser, decl, clauses, true, true, false);
+ else if (c_parser_peek_token (parser)->type == CPP_PRAGMA)
+ /* This will emit an error. */
+ c_finish_oacc_routine (parser, NULL_TREE, clauses, false, true, 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, bool named, bool first, bool is_defn)
+{
+ location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses));
+
+ if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL || !first)
+ {
+ if (fndecl != error_mark_node)
+ error_at (loc, "%<#pragma acc routine%> %s",
+ named ? "does not refer to a function"
+ : "not followed by single 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);
+
+ /* Also attach as a declare. */
+ DECL_ATTRIBUTES (fndecl)
+ = tree_cons (get_identifier ("omp declare target"),
+ clauses, DECL_ATTRIBUTES (fndecl));
+}
+
+/* OpenACC 2.0:
# pragma acc update oacc-update-clause[optseq] new-line
*/
Index: gcc/cp/parser.c
===================================================================
--- gcc/cp/parser.c (revision 230040)
+++ gcc/cp/parser.c (working copy)
@@ -245,6 +245,8 @@ 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))
@@ -1320,6 +1322,15 @@ cp_finalize_omp_declare_simd (cp_parser
}
}
}
+
+/* Diagnose if #pragma omp 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);
+}
/* Decl-specifiers. */
@@ -3620,6 +3631,9 @@ cp_parser_new (void)
parser->implicit_template_parms = 0;
parser->implicit_template_scope = 0;
+ /* Active OpenACC routine clauses. */
+ parser->oacc_routine = NULL;
+
/* Allow constrained-type-specifiers. */
parser->prevent_constrained_type_specifiers = 0;
@@ -12541,6 +12555,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);
@@ -17058,6 +17073,7 @@ cp_parser_namespace_definition (cp_parse
int nested_definition_count = 0;
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);
@@ -18081,6 +18097,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, 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. */
@@ -18194,6 +18211,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, false);
}
/* Finish processing the declaration. But, skip member
@@ -20800,6 +20818,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);
@@ -22113,6 +22132,7 @@ cp_parser_member_declaration (cp_parser*
}
cp_finalize_omp_declare_simd (parser, decl);
+ cp_finalize_oacc_routine (parser, decl, false);
/* Reset PREFIX_ATTRIBUTES. */
while (attributes && TREE_CHAIN (attributes) != first_attribute)
@@ -24716,6 +24736,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, true);
}
if (!success_p)
@@ -25398,6 +25419,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, true);
/* If something went badly wrong, bail out now. */
if (fn == error_mark_node)
{
@@ -35584,6 +35606,147 @@ cp_parser_omp_taskloop (cp_parser *parse
return ret;
}
+
+/* OpenACC 2.0:
+ # pragma acc routine oacc-routine-clause[optseq] new-line
+ function-definition
+
+ # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line
+*/
+
+#define OACC_ROUTINE_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \
+ | (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)
+{
+ 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)
+ {
+ 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
+ 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 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");
+
+ /* Look for optional '( name )'. */
+ if (cp_lexer_next_token_is (parser->lexer,CPP_OPEN_PAREN))
+ {
+ cp_lexer_consume_token (parser->lexer);
+ cp_token *token = cp_lexer_peek_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);
+ 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);
+
+ if (decl == error_mark_node
+ || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ {
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ 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)
+ cp_parser_finish_oacc_routine (parser, decl, clauses, true, false);
+ else
+ parser->oacc_routine = clauses;
+}
+
+/* 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)
+{
+ if (parser->oacc_routine)
+ {
+ cp_parser_finish_oacc_routine (parser, fndecl, parser->oacc_routine,
+ false, is_defn);
+ parser->oacc_routine = NULL_TREE;
+ }
+}
+
/* Main entry point to OpenMP statement pragmas. */
static void
@@ -36063,8 +36226,9 @@ cp_parser_pragma (cp_parser *parser, enu
parser->lexer->in_pragma = true;
id = pragma_tok->pragma_kind;
- if (id != PRAGMA_OMP_DECLARE_REDUCTION)
+ 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:
@@ -36174,6 +36338,10 @@ cp_parser_pragma (cp_parser *parser, enu
cp_parser_omp_declare (parser, pragma_tok, context);
return false;
+ case PRAGMA_OACC_ROUTINE:
+ cp_parser_oacc_routine (parser, pragma_tok, context);
+ return false;
+
case PRAGMA_OACC_ATOMIC:
case PRAGMA_OACC_CACHE:
case PRAGMA_OACC_DATA:
Index: gcc/cp/parser.h
===================================================================
--- gcc/cp/parser.h (revision 230040)
+++ gcc/cp/parser.h (working copy)
@@ -371,6 +371,9 @@ 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;
+
/* Nonzero if parsing a parameter list where 'auto' should trigger an implicit
template parameter. */
bool auto_is_implicit_function_template_parm_p;