From patchwork Mon Nov 2 19:21:43 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 539083 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 34511140D98 for ; Tue, 3 Nov 2015 06:21:57 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=CYsA4PSL; 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:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=JxnpTrEoscUWyYceb UGzcx+0cGIL1AgJRzXeP0uhot817+qEzaRcX5TlH/G1QscrHv8feXc7Jigghr0+h YfjSNVO43YDwiKg4vvdf6r5tNjj/yWv0KtTEsF3kTJ8CNVLr1wt88DGjdN7Za2ow 6apikWAzNdDE9w5wnnu2cDO5Yo= 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:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=57Lo00D+Zm/NrEgPuwyp5cK YLFA=; b=CYsA4PSLTlo6Ur9EXqSysd6wj8OpLz5kIiX/vuWqVcQ2XPM/uy1RI5c lSUWEYeS2d2su7Z8vtvXKyXO76AuwpZqsMR6EAUoi0vzhPn+Lp1riz2ADQfnV0Qn DlXfvJUMPzr2ctdYbsY92X2rmoeaptqaG5aKOZjg/8F+QQrGflSg= Received: (qmail 96677 invoked by alias); 2 Nov 2015 19:21:50 -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 96667 invoked by uid 89); 2 Nov 2015 19:21:49 -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-f52.google.com Received: from mail-qg0-f52.google.com (HELO mail-qg0-f52.google.com) (209.85.192.52) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Mon, 02 Nov 2015 19:21:46 +0000 Received: by qgad10 with SMTP id d10so124412910qga.3 for ; Mon, 02 Nov 2015 11:21:44 -0800 (PST) X-Received: by 10.140.93.104 with SMTP id c95mr31187780qge.101.1446492104372; Mon, 02 Nov 2015 11:21:44 -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 22sm8357911qhq.39.2015.11.02.11.21.43 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Mon, 02 Nov 2015 11:21:43 -0800 (PST) Subject: Re: [1/2] OpenACC routine support To: Jakub Jelinek , GCC Patches References: <5637B1CF.5060408@acm.org> From: Nathan Sidwell Message-ID: <5637B7C7.70901@acm.org> Date: Mon, 2 Nov 2015 14:21:43 -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: <5637B1CF.5060408@acm.org> This is the core changes, an C & C++ FE parsing pieces. Parsing only deals with the gang, worker, vector & seq clauses. The nohost and bind clauses will be a later patch to port. The parsing is very similar to the omp declare simd parsing, in the way that it's hooked into the rest of the parser. conversion of the gang, worker, vecto & seq clauses to the internal 'oacc function' attribute format is done by a new routine, build_oacc_routine_dims, in omp-low.c. This sets the dimensions over which a routine might partition a loop to integer_zero_node, and sets TREE_PURPOSE of such dimensions to zero. while those dimensions over which it must not partition a loop have TREE_PURPOSE set to non-zero. the handling for validating this in the lower_oacc_device pass, and in the PTX backend is already there. ok? nathan 2015-11-02 Nathan Sidwell * 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 Cesar Philippidis James Norris Julian Brown Nathan Sidwell 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 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-02 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/c/c-parser.c =================================================================== --- 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); 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); +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 omp_declare_simd_clauses) + vec 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 " Index: gcc/c-family/c-pragma.c =================================================================== --- gcc/c-family/c-pragma.c (revision 229667) +++ gcc/c-family/c-pragma.c (working copy) @@ -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 } }; Index: gcc/c-family/c-pragma.h =================================================================== --- gcc/c-family/c-pragma.h (revision 229667) +++ gcc/c-family/c-pragma.h (working copy) @@ -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, Index: gcc/cp/parser.c =================================================================== --- gcc/cp/parser.c (revision 229667) +++ gcc/cp/parser.c (working copy) @@ -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: Index: gcc/cp/parser.h =================================================================== --- gcc/cp/parser.h (revision 229667) +++ 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; Index: gcc/omp-low.c =================================================================== --- gcc/omp-low.c (revision 229667) +++ gcc/omp-low.c (working copy) @@ -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. */ Index: gcc/omp-low.h =================================================================== --- gcc/omp-low.h (revision 229667) +++ 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 GTY(()) vec *offload_funcs;