From patchwork Sat Nov 12 16:51:00 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 81927 Delivered-To: patch@linaro.org Received: by 10.140.97.165 with SMTP id m34csp310809qge; Sat, 12 Nov 2016 08:51:36 -0800 (PST) X-Received: by 10.98.4.134 with SMTP id 128mr1102577pfe.156.1478969496136; Sat, 12 Nov 2016 08:51:36 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id l11si15722098pgc.328.2016.11.12.08.51.35 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Sat, 12 Nov 2016 08:51:36 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-441231-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) client-ip=209.132.180.131; Authentication-Results: mx.google.com; dkim=pass header.i=@gcc.gnu.org; spf=pass (google.com: domain of gcc-patches-return-441231-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-441231-patch=linaro.org@gcc.gnu.org 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=bDE1kHF+p2c7Dh5cz 1f+3Vowv2egKxGtfNA10Z7EWn3Hc8/n+5T+a/Yl7uARXFipuJq0qtjq3/1jKTTPm bi8+biJI6gHRGsqubm64HDpg4vtwyVVuXkET8InEoGBmqetzrCfYl7A6KsF+Q6S3 /4Vq5jqq/BD/Zt9ccVFi8/blyM= 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=NMJZ+0PdhoyfD39ldmkmHBJ PRRI=; b=L8P/XvW027i1EJ2j8dO+MujNmA8yUTvzOINZ82mN2GzK3VO3dLtPiM7 rBo2KkkL2iL/0S0yLioD0DIcHINdF2Wzq+E2Uc47XZZTt6y9hgNKLXEvzg/oKs2j 3lAoFAsJ1s+pkk5XsgtRUScmpCLh5glDHnX3ewtnQDv0wrxGzSbI= Received: (qmail 70202 invoked by alias); 12 Nov 2016 16:51:20 -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 70177 invoked by uid 89); 12 Nov 2016 16:51:18 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=flag_cilkplus, sk:decl_in, sk:DECL_FU, sk:DECL_IN X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Nov 2016 16:51:08 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1c5bWC-0001qY-5l from Cesar_Philippidis@mentor.com ; Sat, 12 Nov 2016 08:51:04 -0800 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Sat, 12 Nov 2016 08:51:01 -0800 Subject: Re: [Patch 4/5] OpenACC tile clause support, Fortran front-end parts To: Jakub Jelinek , Chung-Lin Tang References: <87a13495-a378-b5e1-8223-0f5679f98356@mentor.com> <20161111103441.GY3541@tucnak.redhat.com> CC: gcc-patches , Nathan Sidwell From: Cesar Philippidis Message-ID: Date: Sat, 12 Nov 2016 08:51:00 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.3.0 MIME-Version: 1.0 In-Reply-To: <20161111103441.GY3541@tucnak.redhat.com> X-ClientProxiedBy: svr-orw-mbx-04.mgc.mentorg.com (147.34.90.204) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) On 11/11/2016 02:34 AM, Jakub Jelinek wrote: > On Thu, Nov 10, 2016 at 06:46:46PM +0800, Chung-Lin Tang wrote: And here's the patch. Cesar >> 2016-XX-XX Cesar Philippidis >> >> fortran/ >> * openmp.c (resolve_oacc_positive_int_expr): Promote the warning >> to an error. >> (resolve_oacc_loop_blocks): Use integer zero to represent the '*' >> tile argument. >> (resolve_omp_clauses): Error on directives containing both tile >> and collapse clauses. >> * trans-openmp.c (gfc_trans_omp_do): Lower tiled loops like >> collapsed loops. >> >> >> gcc/testsuite/ >> * gfortran.dg/goacc/loop-2.f95: Change expected tile clause >> warnings to errors. >> * gfortran.dg/goacc/loop-5.f95: Likewise. >> * gfortran.dg/goacc/sie.f95: Likewise. >> * gfortran.dg/goacc/tile-1.f90: New test. >> * gfortran.dg/goacc/tile-2.f90: New test >> * gfortran.dg/goacc/tile-lowering.f95: New test. > > Again, 8 spaces in ChangeLog. Missing full stop after New test > >> --- fortran/openmp.c (revision 241809) >> +++ fortran/openmp.c (working copy) >> @@ -3024,8 +3024,8 @@ resolve_oacc_positive_int_expr (gfc_expr *expr, co >> resolve_oacc_scalar_int_expr (expr, clause); >> if (expr->expr_type == EXPR_CONSTANT && expr->ts.type == BT_INTEGER >> && mpz_sgn(expr->value.integer) <= 0) >> - gfc_warning (0, "INTEGER expression of %s clause at %L must be positive", >> - clause, &expr->where); >> + gfc_error ("INTEGER expression of %s clause at %L must be positive", >> + clause, &expr->where); >> } > > This can't be against current trunk. The current routine is shared with > OpenMP and gfc_error is undesirable there. > >> @@ -3859,6 +3859,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus >> if (omp_clauses->wait_list) >> for (el = omp_clauses->wait_list; el; el = el->next) >> resolve_oacc_scalar_int_expr (el->expr, "WAIT"); >> + if (omp_clauses->collapse && omp_clauses->tile_list) >> + gfc_error ("Incompatible use of TILE and COLLAPSE at %L", &code->loc); > > Shouldn't you in that case for error recovery clear collapse (or tile_list)? > > Jakub > 2016-11-11 Cesar Philippidis Thomas Schwinge gcc/cp/ * cp-tree.h (bind_decls_match): Declare. * decl.c (bind_decls_match): New function. * parser.c (cp_parser_omp_clause_name): (cp_parser_oacc_shape_clause): New location_t loc argument. Use it to report more accurate diagnostics. (cp_parser_oacc_clause_bind): New function. (cp_parser_oacc_all_clauses): Handle OpenACC bind and nohost clauses. Update calls to c_parser_oacc_{simple,shape}_clause. (OACC_ROUTINE_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{BIND,NOHOST}. (cp_parser_oacc_routine): Update diagnostics. (cp_parser_late_parsing_oacc_routine): Likewise. (cp_finalize_oacc_routine): Likewise. * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_{BIND,NOHOST}. diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 8183775..efd5450 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -5772,6 +5772,7 @@ extern void finish_scope (void); extern void push_switch (tree); extern void pop_switch (void); extern tree make_lambda_name (void); +extern int bind_decls_match (tree, tree); extern int decls_match (tree, tree); extern tree duplicate_decls (tree, tree, bool); extern tree declare_local_label (tree); diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c index 185c98b..7b242781 100644 --- a/gcc/cp/decl.c +++ b/gcc/cp/decl.c @@ -1198,6 +1198,138 @@ decls_match (tree newdecl, tree olddecl) return types_match; } +/* Similiar to decls_match, but only applies to FUNCTION_DECLS. Functions + in separate namespaces may match. +*/ + +int +bind_decls_match (tree newdecl, tree olddecl) +{ + int types_match; + + if (newdecl == olddecl) + return 1; + + if (TREE_CODE (newdecl) != TREE_CODE (olddecl)) + /* If the two DECLs are not even the same kind of thing, we're not + interested in their types. */ + return 0; + + gcc_assert (DECL_P (newdecl)); + gcc_assert (TREE_CODE (newdecl) == FUNCTION_DECL); + + tree f1 = TREE_TYPE (newdecl); + tree f2 = TREE_TYPE (olddecl); + tree p1 = TYPE_ARG_TYPES (f1); + tree p2 = TYPE_ARG_TYPES (f2); + tree r2; + + /* Specializations of different templates are different functions + even if they have the same type. */ + tree t1 = (DECL_USE_TEMPLATE (newdecl) + ? DECL_TI_TEMPLATE (newdecl) + : NULL_TREE); + tree t2 = (DECL_USE_TEMPLATE (olddecl) + ? DECL_TI_TEMPLATE (olddecl) + : NULL_TREE); + if (t1 != t2) + return 0; + + if (CP_DECL_CONTEXT (newdecl) != CP_DECL_CONTEXT (olddecl) + && TREE_CODE (CP_DECL_CONTEXT (newdecl)) != NAMESPACE_DECL + && TREE_CODE (CP_DECL_CONTEXT (olddecl)) != NAMESPACE_DECL + && ! (DECL_EXTERN_C_P (newdecl) + && DECL_EXTERN_C_P (olddecl))) + return 0; + + /* A new declaration doesn't match a built-in one unless it + is also extern "C". */ + if (DECL_IS_BUILTIN (olddecl) + && DECL_EXTERN_C_P (olddecl) && !DECL_EXTERN_C_P (newdecl)) + return 0; + + if (TREE_CODE (f1) != TREE_CODE (f2)) + return 0; + + /* A declaration with deduced return type should use its pre-deduction + type for declaration matching. */ + r2 = fndecl_declared_return_type (olddecl); + + if (same_type_p (TREE_TYPE (f1), r2)) + { + if (!prototype_p (f2) && DECL_EXTERN_C_P (olddecl) + && (DECL_BUILT_IN (olddecl) +#ifndef NO_IMPLICIT_EXTERN_C + || (DECL_IN_SYSTEM_HEADER (newdecl) && !DECL_CLASS_SCOPE_P (newdecl)) + || (DECL_IN_SYSTEM_HEADER (olddecl) && !DECL_CLASS_SCOPE_P (olddecl)) +#endif + )) + { + types_match = self_promoting_args_p (p1); + if (p1 == void_list_node) + TREE_TYPE (newdecl) = TREE_TYPE (olddecl); + } +#ifndef NO_IMPLICIT_EXTERN_C + else if (!prototype_p (f1) + && (DECL_EXTERN_C_P (olddecl) + && DECL_IN_SYSTEM_HEADER (olddecl) + && !DECL_CLASS_SCOPE_P (olddecl)) + && (DECL_EXTERN_C_P (newdecl) + && DECL_IN_SYSTEM_HEADER (newdecl) + && !DECL_CLASS_SCOPE_P (newdecl))) + { + types_match = self_promoting_args_p (p2); + TREE_TYPE (newdecl) = TREE_TYPE (olddecl); + } +#endif + else + types_match = + compparms (p1, p2) + && type_memfn_rqual (f1) == type_memfn_rqual (f2) + && (TYPE_ATTRIBUTES (TREE_TYPE (newdecl)) == NULL_TREE + || comp_type_attributes (TREE_TYPE (newdecl), + TREE_TYPE (olddecl)) != 0); + } + else + types_match = 0; + + /* The decls dont match if they correspond to two different versions + of the same function. Disallow extern "C" functions to be + versions for now. */ + if (types_match + && !DECL_EXTERN_C_P (newdecl) + && !DECL_EXTERN_C_P (olddecl) + && targetm.target_option.function_versions (newdecl, olddecl)) + { + /* Mark functions as versions if necessary. Modify the mangled decl + name if necessary. */ + if (DECL_FUNCTION_VERSIONED (newdecl) + && DECL_FUNCTION_VERSIONED (olddecl)) + return 0; + if (!DECL_FUNCTION_VERSIONED (newdecl)) + { + DECL_FUNCTION_VERSIONED (newdecl) = 1; + if (DECL_ASSEMBLER_NAME_SET_P (newdecl)) + mangle_decl (newdecl); + } + if (!DECL_FUNCTION_VERSIONED (olddecl)) + { + DECL_FUNCTION_VERSIONED (olddecl) = 1; + if (DECL_ASSEMBLER_NAME_SET_P (olddecl)) + mangle_decl (olddecl); + } + cgraph_node::record_function_versions (olddecl, newdecl); + return 0; + } + + /* Normal functions can be constrained, as can variable partial + specializations. */ + if (types_match && VAR_OR_FUNCTION_DECL_P (newdecl)) + types_match = equivalently_constrained (newdecl, olddecl); + + return types_match; +} + /* If NEWDECL is `static' and an `extern' was seen previously, warn about it. OLDDECL is the previous declaration. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7b95dba..33ee6a5 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -30223,6 +30223,10 @@ cp_parser_omp_clause_name (cp_parser *parser) else if (!strcmp ("async", p)) result = PRAGMA_OACC_CLAUSE_ASYNC; break; + case 'b': + if (!strcmp ("bind", p)) + result = PRAGMA_OACC_CLAUSE_BIND; + break; case 'c': if (!strcmp ("collapse", p)) result = PRAGMA_OMP_CLAUSE_COLLAPSE; @@ -30302,6 +30306,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; else if (!strcmp ("nowait", p)) result = PRAGMA_OMP_CLAUSE_NOWAIT; + else if (!strcmp ("nohost", p)) + result = PRAGMA_OACC_CLAUSE_NOHOST; else if (flag_cilkplus && !strcmp ("nomask", p)) result = PRAGMA_CILK_CLAUSE_NOMASK; else if (!strcmp ("num_gangs", p)) @@ -30778,13 +30784,13 @@ cp_parser_oacc_single_int_clause (cp_parser *parser, omp_clause_code code, */ static tree -cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind, +cp_parser_oacc_shape_clause (cp_parser *parser, location_t loc, + omp_clause_code kind, const char *str, tree list) { const char *id = "num"; cp_lexer *lexer = parser->lexer; tree ops[2] = { NULL_TREE, NULL_TREE }, c; - location_t loc = cp_lexer_peek_token (lexer)->location; if (kind == OMP_CLAUSE_VECTOR) id = "length"; @@ -32571,6 +32577,82 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list) return list; } +/* OpenACC 2.0: + bind ( identifier ) + bind ( string-literal ) */ + +static tree +cp_parser_oacc_clause_bind (cp_parser *parser, tree list) +{ + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + + check_no_duplicate_clause (list, OMP_CLAUSE_BIND, "bind", loc); + + bool save_translate_strings_p = parser->translate_strings_p; + parser->translate_strings_p = false; + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + { + parser->translate_strings_p = save_translate_strings_p; + return list; + } + tree name = error_mark_node; + cp_token *token = cp_lexer_peek_token (parser->lexer); + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + tree id = cp_parser_id_expression (parser, /*template_p=*/false, + /*check_dependency_p=*/true, + /*template_p=*/NULL, + /*declarator_p=*/false, + /*optional_p=*/false); + tree 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 || decl == error_mark_node) + error_at (token->location, "%qE has not been declared", + token->u.value); + else if (is_overloaded_fn (decl) + && (TREE_CODE (decl) != FUNCTION_DECL + || DECL_FUNCTION_TEMPLATE_P (decl))) + error_at (token->location, "%qE names a set of overloads", + token->u.value); + else if (TREE_CODE (decl) != FUNCTION_DECL) + error_at (token->location, + "%qE does not refer to a function", + token->u.value); + else + name = decl; + } + else if (cp_lexer_next_token_is (parser->lexer, CPP_STRING)) + { + name = token->u.value; + cp_lexer_consume_token (parser->lexer); + + /* This shouldn't be an empty string. */ + if (strcmp (TREE_STRING_POINTER (name), "\"\"") == 0) + error_at (token->location, + "bind argument must not be an empty string"); + + parser->translate_strings_p = save_translate_strings_p; + } + else + { + cp_parser_error (parser, + "expected identifier or character string literal"); + cp_parser_skip_to_closing_parenthesis (parser, false, false, true); + return list; + } + cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN); + if (name != error_mark_node) + { + tree c = build_omp_clause (loc, OMP_CLAUSE_BIND); + OMP_CLAUSE_BIND_NAME (c) = name; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + return list; +} + /* Parse all OpenACC clauses. The set clauses allowed by the directive is a bitmask in MASK. Return the list of clauses found. */ @@ -32607,6 +32689,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses, here); c_name = "auto"; break; + case PRAGMA_OACC_CLAUSE_BIND: + clauses = cp_parser_oacc_clause_bind (parser, clauses); + c_name = "bind"; + break; case PRAGMA_OACC_CLAUSE_COLLAPSE: clauses = cp_parser_omp_clause_collapse (parser, clauses, here); c_name = "collapse"; @@ -32654,7 +32740,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, break; case PRAGMA_OACC_CLAUSE_GANG: c_name = "gang"; - clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG, + clauses = cp_parser_oacc_shape_clause (parser, here, OMP_CLAUSE_GANG, c_name, clauses); break; case PRAGMA_OACC_CLAUSE_HOST: @@ -32675,6 +32761,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "link"; break; + case PRAGMA_OACC_CLAUSE_NOHOST: + clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_NOHOST, + clauses, here); + c_name = "nohost"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: code = OMP_CLAUSE_NUM_GANGS; c_name = "num_gangs"; @@ -32736,7 +32827,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; - clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, + clauses = cp_parser_oacc_shape_clause (parser, here, + OMP_CLAUSE_VECTOR, c_name, clauses); break; case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: @@ -32751,7 +32843,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, break; case PRAGMA_OACC_CLAUSE_WORKER: c_name = "worker"; - clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_WORKER, + clauses = cp_parser_oacc_shape_clause (parser, here, + OMP_CLAUSE_WORKER, c_name, clauses); break; default: @@ -36992,7 +37085,10 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok, ( (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)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST)) + /* Parse the OpenACC routine pragma. This has an optional '( name )' @@ -37051,6 +37147,9 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine", cp_lexer_peek_token (parser->lexer)); + /* The clauses are in reverse order; fix that to make later diagnostic + emission easier. */ + data.clauses = nreverse (data.clauses); if (decl && is_overloaded_fn (decl) && (TREE_CODE (decl) != FUNCTION_DECL @@ -37062,16 +37161,6 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, return; } - /* Perhaps we should use the same rule as declarations in different - namespaces? */ - if (!DECL_NAMESPACE_SCOPE_P (decl)) - { - error_at (name_loc, - "%qD does not refer to a namespace scope function", decl); - parser->oacc_routine = NULL; - return; - } - if (TREE_CODE (decl) != FUNCTION_DECL) { error_at (name_loc, "%qD does not refer to a function", decl); @@ -37147,6 +37236,9 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) parser->oacc_routine->clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine", pragma_tok); + /* The clauses are in reverse order; fix that to make later diagnostic + emission easier. */ + parser->oacc_routine->clauses = nreverse (parser->oacc_routine->clauses); cp_parser_pop_lexer (parser); /* Later, cp_finalize_oacc_routine will process the clauses, and then set fndecl_seen. */ @@ -37181,31 +37273,70 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) return; } - if (get_oacc_fn_attrib (fndecl)) + /* Process the bind clause, if present. */ + for (tree c = parser->oacc_routine->clauses; + c; + c = OMP_CLAUSE_CHAIN (c)) { - error_at (parser->oacc_routine->loc, - "%<#pragma acc routine%> already applied to %qD", fndecl); - parser->oacc_routine = NULL; - return; + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_BIND) + continue; + + tree bind_decl = OMP_CLAUSE_BIND_NAME (c); + + /* String arguments don't require any special treatment. */ + if (TREE_CODE (bind_decl) != FUNCTION_DECL) + break; + + if (!bind_decls_match (bind_decl, fndecl)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "bind identifier %qE is not compatible with " + "function %qE", bind_decl, fndecl); + parser->oacc_routine = NULL; + return; + } + + tree name_id = decl_assembler_name (bind_decl); + tree name = build_string (IDENTIFIER_LENGTH (name_id), + IDENTIFIER_POINTER (name_id)); + OMP_CLAUSE_BIND_NAME (c) = name; + + break; } - if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + int compatible + = verify_oacc_routine_clauses (fndecl, &parser->oacc_routine->clauses, + parser->oacc_routine->loc, + "#pragma acc routine"); + if (compatible < 0) { - error_at (parser->oacc_routine->loc, - "%<#pragma acc routine%> must be applied before %s", - TREE_USED (fndecl) ? "use" : "definition"); parser->oacc_routine = NULL; return; } + if (compatible > 0) + { + } + else + { + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + { + error_at (parser->oacc_routine->loc, + "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + parser->oacc_routine = NULL; + return; + } - /* Process the routine's dimension clauses. */ - tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses); - replace_oacc_fn_attrib (fndecl, dims); - - /* Add an "omp declare target" attribute. */ - DECL_ATTRIBUTES (fndecl) - = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (fndecl)); + /* Set the routine's level of parallelism. */ + tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses); + replace_oacc_fn_attrib (fndecl, dims); + + /* Add an "omp declare target" attribute. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + parser->oacc_routine->clauses, + DECL_ATTRIBUTES (fndecl)); + } /* Don't unset parser->oacc_routine here: we may still need it to diagnose wrong usage. But, remember that we've used this "#pragma acc diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 1a7c478..8958347 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7073,6 +7073,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_AUTO: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: break; case OMP_CLAUSE_TILE: