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
This commit is contained in:
Nathan Sidwell
parent
7441dcf4be
commit
3a40d81dcd
|
|
@ -1,3 +1,8 @@
|
|||
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-08 Michael Meissner <meissner@linux.vnet.ibm.com>
|
||||
|
||||
* config/rs6000/constraints.md (wF constraint): New constraints
|
||||
|
|
|
|||
|
|
@ -1,3 +1,12 @@
|
|||
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-pragma.c (oacc_pragmas): Add "routine".
|
||||
* c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE.
|
||||
|
||||
2015-11-08 Eric Botcazou <ebotcazou@adacore.com>
|
||||
|
||||
* c-common.c (c_common_attributes): Add scalar_storage_order.
|
||||
|
|
|
|||
|
|
@ -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 }
|
||||
};
|
||||
|
|
|
|||
|
|
@ -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,
|
||||
|
||||
|
|
|
|||
|
|
@ -1,3 +1,17 @@
|
|||
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 Andreas Arnez <arnez@linux.vnet.ibm.com>
|
||||
|
||||
PR debug/67192
|
||||
|
|
|
|||
141
gcc/c/c-parser.c
141
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<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_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<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 *parser, bool fndef_ok,
|
|||
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 *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
|
||||
*/
|
||||
|
|
|
|||
|
|
@ -1,3 +1,25 @@
|
|||
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>
|
||||
|
||||
* 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 <msebor@redhat.com>
|
||||
|
||||
PR c++/67942
|
||||
|
|
|
|||
170
gcc/cp/parser.c
170
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:
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -12350,6 +12350,50 @@ set_oacc_fn_attrib (tree fn, tree clauses, vec<tree> *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. */
|
||||
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
Loading…
Reference in New Issue