From patchwork Tue Nov 10 00:31:45 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 542087 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id B6856140D69 for ; Tue, 10 Nov 2015 11:34:29 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=cVVBXOIi; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=MTPKWU5vNBmSMyMs1 6+L3xN3m+VtXlR2hTYH4+VJPIw+ygrDFvL0dkTzpsMLFDM4M9cqZko/AagoGjMAo gDOdeE6Eok4QIOXOlBSGjb9jeqUEVd5b+uQ1k9P18/97uLV2E07nTt64nNR1MrZQ GHgucoAPyzcM02k948CJQcBFgU= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=Tt1BACQYGFr0t9TMpHdigg8 DofI=; b=cVVBXOIi2oZkI0UPEe+v7iKFX59yiCyzQ603+E0LxkQnQx3Wk+akbCr /Uc8BJuYw4R4FlzmUOThkhtG8CCc7drVVqdytXjdevuw2oszWhU53BT89poyAOky xzF9ZynqpLE0ygVdJ4hJkz/pO/5FyJx67GEcPIxfG0cJY6kBd4Hk= Received: (qmail 107062 invoked by alias); 10 Nov 2015 00:32:03 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 106977 invoked by uid 89); 10 Nov 2015 00:32:03 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.9 required=5.0 tests=BAYES_50, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-qg0-f45.google.com Received: from mail-qg0-f45.google.com (HELO mail-qg0-f45.google.com) (209.85.192.45) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Tue, 10 Nov 2015 00:31:50 +0000 Received: by qgeb1 with SMTP id b1so106087856qge.1 for ; Mon, 09 Nov 2015 16:31:47 -0800 (PST) X-Received: by 10.140.181.197 with SMTP id c188mr983245qha.4.1447115507089; Mon, 09 Nov 2015 16:31:47 -0800 (PST) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id 18sm297758qho.24.2015.11.09.16.31.46 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Mon, 09 Nov 2015 16:31:46 -0800 (PST) Subject: Re: [1/2] OpenACC routine support To: Jakub Jelinek References: <5637B1CF.5060408@acm.org> <5637B7C7.70901@acm.org> <20151103153533.GQ478@tucnak.redhat.com> Cc: GCC Patches , Cesar Philippidis From: Nathan Sidwell Message-ID: <56413AF1.8070101@acm.org> Date: Mon, 9 Nov 2015 19:31:45 -0500 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <20151103153533.GQ478@tucnak.redhat.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); >> + bool, bool, tree *, vec, >> + 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 * 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 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 Thomas Schwinge Cesar Philippidis James Norris Julian Brown Nathan Sidwell c-family/ * c-pragma.c (oacc_pragmas): Add "routine". * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE. 2015-11-09 Thomas Schwinge Cesar Philippidis James Norris Julian Brown Nathan Sidwell 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); + 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_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); +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 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 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;