From: Nathan Sidwell Date: Tue, 10 Nov 2015 00:27:26 +0000 (+0000) Subject: omp-low.h (replace_oacc_fn_attrib, [...]): Declare. X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=3a40d81dcd1bae2c8dea064cdd0f28c9269813da;p=gcc.git omp-low.h (replace_oacc_fn_attrib, [...]): Declare. * omp-low.h (replace_oacc_fn_attrib, build_oacc_routine_dims): Declare. * omp-low.c (build_oacc_routine_dims): New. 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. c-family/ * c-pragma.c (oacc_pragmas): Add "routine". * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE. 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. From-SVN: r230072 --- diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 08e176026ac..bf785384ec8 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2015-11-09 Nathan Sidwell + + * omp-low.h (replace_oacc_fn_attrib, build_oacc_routine_dims): Declare. + * omp-low.c (build_oacc_routine_dims): New. + 2015-11-08 Michael Meissner * config/rs6000/constraints.md (wF constraint): New constraints diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index f52b3e354c5..4188535df78 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,12 @@ +2015-11-09 Thomas Schwinge + Cesar Philippidis + James Norris + Julian Brown + Nathan Sidwell + + * c-pragma.c (oacc_pragmas): Add "routine". + * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE. + 2015-11-08 Eric Botcazou * c-common.c (c_common_attributes): Add scalar_storage_order. diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c index ffc5b33a4cd..50324b82498 100644 --- a/gcc/c-family/c-pragma.c +++ b/gcc/c-family/c-pragma.c @@ -1257,6 +1257,7 @@ static const struct omp_pragma_def oacc_pragmas[] = { { "kernels", PRAGMA_OACC_KERNELS }, { "loop", PRAGMA_OACC_LOOP }, { "parallel", PRAGMA_OACC_PARALLEL }, + { "routine", PRAGMA_OACC_ROUTINE }, { "update", PRAGMA_OACC_UPDATE }, { "wait", PRAGMA_OACC_WAIT } }; diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 59231ce030e..afeceff1edb 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -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, diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index a666a3e32a3..0124dcc578d 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,17 @@ +2015-11-09 Thomas Schwinge + Cesar Philippidis + James Norris + Julian Brown + Nathan Sidwell + + 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 Andreas Arnez PR debug/67192 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 66eb2dd188a..dd4434dc7ae 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -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); + bool, bool, tree *, vec, + 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_parser *, enum pragma_context); 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 *parser) } static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec); +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 *parser, bool fndef_ok, bool static_assert_ok, bool empty_ok, bool nested, bool start_attr_ok, tree *objc_foreach_object_declaration, - vec omp_declare_simd_clauses) + vec 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 *parser, bool fndef_ok, 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 *parser, bool fndef_ok, 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 *parser, bool fndef_ok, || !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 *parser, bool fndef_ok, 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 *parser, bool fndef_ok, 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 *parser, bool fndef_ok, || !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 pragma_context context) 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) { @@ -13399,6 +13423,117 @@ c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, return c_finish_omp_construct (loc, code, block, clauses); } +/* 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 */ diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 11ca9ab584a..91cbf067dba 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,25 @@ +2015-11-09 Thomas Schwinge + Cesar Philippidis + James Norris + Julian Brown + Nathan Sidwell + + * 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. + 2015-11-08 Martin Sebor PR c++/67942 diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index d1f4970f1eb..6fc2c6af389 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -245,6 +245,8 @@ static bool cp_parser_omp_declare_reduction_exprs (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 *parser, tree fndecl) } } } + +/* 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_parser* parser) 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_parser* parser) 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* parser, range_for_decl_p? SD_INITIALIZED : is_initialized, attributes, prefix_attributes, &pushed_scope); cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl, false); /* 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* parser, if (decl && TREE_CODE (decl) == FUNCTION_DECL) cp_parser_save_default_args (parser, decl); cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl, false); } /* Finish processing the declaration. But, skip member @@ -20800,6 +20818,7 @@ cp_parser_class_specifier_1 (cp_parser* 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* 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_specifiers_and_declarator { cp_finalize_omp_declare_simd (parser, current_function_decl); parser->omp_declare_simd = NULL; + cp_finalize_oacc_routine (parser, current_function_decl, true); } if (!success_p) @@ -25398,6 +25419,7 @@ cp_parser_save_member_function_body (cp_parser* parser, /* Create the FUNCTION_DECL. */ fn = grokmethod (decl_specifiers, declarator, attributes); cp_finalize_omp_declare_simd (parser, fn); + cp_finalize_oacc_routine (parser, fn, true); /* If something went badly wrong, bail out now. */ if (fn == error_mark_node) { @@ -35584,6 +35606,147 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok, 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, enum pragma_context context) 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, enum pragma_context context) 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: diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h index 760467c2eb0..fdbff6664ef 100644 --- a/gcc/cp/parser.h +++ b/gcc/cp/parser.h @@ -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; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 159f78483cf..2a552da1e5b 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12350,6 +12350,50 @@ set_oacc_fn_attrib (tree fn, tree clauses, vec *args) } } +/* 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. */ diff --git a/gcc/omp-low.h b/gcc/omp-low.h index ee0f8ac1ad5..194b3d14d6b 100644 --- a/gcc/omp-low.h +++ b/gcc/omp-low.h @@ -30,6 +30,8 @@ extern tree omp_reduction_init (tree, tree); 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);