2015-11-02 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-02 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. Adjust all callers.
(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-02 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-02 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.
===================================================================
@@ -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);
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,
@@ -1249,6 +1250,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. */
@@ -1428,12 +1430,13 @@ c_parser_external_declaration (c_parser
only tell which after parsing the declaration specifiers, if
any, and the first declarator. */
c_parser_declaration_or_fndef (parser, true, true, true, false, true,
- NULL, vNULL);
+ NULL, vNULL, NULL_TREE);
break;
}
}
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);
/* 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
@@ -1511,7 +1514,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;
@@ -1581,6 +1585,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, false);
return;
}
@@ -1697,6 +1704,9 @@ 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, false);
c_parser_skip_to_end_of_block_or_statement (parser);
return;
}
@@ -1811,6 +1821,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, false);
if (d != error_mark_node)
{
maybe_warn_string_init (init_loc, TREE_TYPE (d), init);
@@ -1854,6 +1867,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, false);
if (d)
finish_decl (d, UNKNOWN_LOCATION, NULL_TREE,
NULL_TREE, asm_name);
@@ -1958,12 +1974,15 @@ c_parser_declaration_or_fndef (c_parser
while (c_parser_next_token_is_not (parser, CPP_EOF)
&& c_parser_next_token_is_not (parser, CPP_OPEN_BRACE))
c_parser_declaration_or_fndef (parser, false, false, false,
- true, false, NULL, vNULL);
+ true, false, NULL, vNULL, NULL_TREE);
store_parm_decls ();
if (omp_declare_simd_clauses.exists ()
|| !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, true);
DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus
= c_parser_peek_token (parser)->location;
fnbody = c_parser_compound_statement (parser);
@@ -4634,7 +4653,7 @@ c_parser_compound_statement_nostart (c_p
last_label = false;
mark_valid_location_for_stdc_pragma (false);
c_parser_declaration_or_fndef (parser, true, true, true, true,
- true, NULL, vNULL);
+ true, NULL, vNULL, NULL_TREE);
if (last_stmt)
pedwarn_c90 (loc, OPT_Wdeclaration_after_statement,
"ISO C90 forbids mixed declarations and code");
@@ -4659,7 +4678,7 @@ c_parser_compound_statement_nostart (c_p
last_label = false;
mark_valid_location_for_stdc_pragma (false);
c_parser_declaration_or_fndef (parser, true, true, true, true,
- true, NULL, vNULL);
+ true, NULL, vNULL, NULL_TREE);
/* Following the old parser, __extension__ does not
disable this diagnostic. */
restore_extension_diagnostics (ext);
@@ -4808,7 +4827,7 @@ c_parser_label (c_parser *parser)
/*static_assert_ok*/ true,
/*empty_ok*/ true, /*nested*/ true,
/*start_attr_ok*/ true, NULL,
- vNULL);
+ vNULL, NULL_TREE);
}
}
}
@@ -5580,7 +5599,7 @@ c_parser_for_statement (c_parser *parser
else if (c_parser_next_tokens_start_declaration (parser))
{
c_parser_declaration_or_fndef (parser, true, true, true, true, true,
- &object_expression, vNULL);
+ &object_expression, vNULL, NULL_TREE);
parser->objc_could_be_foreach_context = false;
if (c_parser_next_token_is_keyword (parser, RID_IN))
@@ -5609,7 +5628,8 @@ c_parser_for_statement (c_parser *parser
ext = disable_extension_diagnostics ();
c_parser_consume_token (parser);
c_parser_declaration_or_fndef (parser, true, true, true, true,
- true, &object_expression, vNULL);
+ true, &object_expression, vNULL,
+ NULL_TREE);
parser->objc_could_be_foreach_context = false;
restore_extension_diagnostics (ext);
@@ -8745,8 +8765,8 @@ c_parser_objc_methodprotolist (c_parser
c_parser_consume_token (parser);
}
else
- c_parser_declaration_or_fndef (parser, false, false, true,
- false, true, NULL, vNULL);
+ c_parser_declaration_or_fndef (parser, false, false, true, false,
+ true, NULL, vNULL, NULL_TREE);
break;
}
}
@@ -9703,6 +9723,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)
{
@@ -13265,6 +13289,111 @@ 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)
+ {
+ decl = lookup_name (token->value);
+ if (!decl)
+ {
+ error_at (token->location, "%qE has not been declared",
+ token->value);
+ decl = error_mark_node;
+ }
+ c_parser_consume_token (parser);
+ }
+ else
+ c_parser_error (parser, "expected function name");
+
+ 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, 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 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 (loc, "%<#pragma acc routine%> %s",
+ named ? "does not refer to a function"
+ : "not followed by 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
*/
@@ -13929,7 +14058,7 @@ c_parser_omp_for_loop (location_t loc, c
vec_safe_push (for_block, c_begin_compound_stmt (true));
this_pre_body = push_stmt_list ();
c_parser_declaration_or_fndef (parser, true, true, true, true, true,
- NULL, vNULL);
+ NULL, vNULL, NULL_TREE);
if (this_pre_body)
{
this_pre_body = pop_stmt_list (this_pre_body);
@@ -15506,12 +15635,12 @@ c_parser_omp_declare_simd (c_parser *par
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, clauses);
+ NULL, clauses, NULL_TREE);
restore_extension_diagnostics (ext);
}
else
c_parser_declaration_or_fndef (parser, true, true, true, false, true,
- NULL, clauses);
+ NULL, clauses, NULL_TREE);
break;
case pragma_struct:
case pragma_param:
@@ -15531,7 +15660,7 @@ c_parser_omp_declare_simd (c_parser *par
if (c_parser_next_tokens_start_declaration (parser))
{
c_parser_declaration_or_fndef (parser, true, true, true, true,
- true, NULL, clauses);
+ true, NULL, clauses, NULL_TREE);
restore_extension_diagnostics (ext);
break;
}
@@ -15540,7 +15669,7 @@ c_parser_omp_declare_simd (c_parser *par
else if (c_parser_next_tokens_start_declaration (parser))
{
c_parser_declaration_or_fndef (parser, true, true, true, true, true,
- NULL, clauses);
+ NULL, clauses, NULL_TREE);
break;
}
c_parser_error (parser, "%<#pragma omp declare simd%> must be followed by "
===================================================================
@@ -1211,6 +1211,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 }
};
===================================================================
@@ -34,6 +34,7 @@ enum pragma_kind {
PRAGMA_OACC_KERNELS,
PRAGMA_OACC_LOOP,
PRAGMA_OACC_PARALLEL,
+ PRAGMA_OACC_ROUTINE,
PRAGMA_OACC_UPDATE,
PRAGMA_OACC_WAIT,
===================================================================
@@ -244,6 +244,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))
@@ -1319,6 +1321,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. */
@@ -3619,6 +3630,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;
@@ -12535,6 +12549,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);
@@ -17062,6 +17077,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);
@@ -18085,6 +18101,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. */
@@ -18198,6 +18215,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
@@ -20804,6 +20822,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);
@@ -22117,6 +22136,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)
@@ -24720,6 +24740,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)
@@ -25402,6 +25423,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)
{
@@ -35453,6 +35475,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 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
@@ -35929,8 +36092,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:
@@ -36040,6 +36204,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_CACHE:
case PRAGMA_OACC_DATA:
case PRAGMA_OACC_ENTER_DATA:
===================================================================
@@ -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;
===================================================================
@@ -12083,6 +12083,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. */
===================================================================
@@ -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 GTY(()) vec<tree, va_gc> *offload_funcs;