From patchwork Tue Nov 22 20:09:38 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 83486 Delivered-To: patch@linaro.org Received: by 10.140.97.165 with SMTP id m34csp2292556qge; Tue, 22 Nov 2016 12:10:15 -0800 (PST) X-Received: by 10.84.218.8 with SMTP id q8mr3103765pli.138.1479845415380; Tue, 22 Nov 2016 12:10:15 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id f7si1331301plm.323.2016.11.22.12.10.15 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Tue, 22 Nov 2016 12:10:15 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-442284-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-442284-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-442284-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=STIYuNI5s6U+Twu1a UJ47KAChPzKcT9WaEwFOWRBm8vlrWFd3p8SIFj3HcXO1zFtST0LYYIAeq4nBrQqj 7KUB33XD+KuUdVL4CG4Es74SxrTPR43IkvLBN5jRIeednZRi+ou1YT+sz1DgXoYu YcWs+j5AD2IOhW9dSgThUOGAu0= 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=B7KlbCvTQG96FHOS0Trgb/f PLIU=; b=brQGBVIPeUtQnjTm7f3dbebXTn6hQGZ9moJfSi9SjpqoWi3v77ya37V q166sjAoPAVqkqjYscuH8ddmWHypc92DMBAOT6tFt1uof/kM9gb7Ep4S/Nx+vlf3 z3I8iRwXBSG9XpYWRc7WUNLmjio2FSmYcyorAIgNS01OVF693paw= Received: (qmail 36694 invoked by alias); 22 Nov 2016 20:09:54 -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 36088 invoked by uid 89); 22 Nov 2016 20:09:54 -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=Updates, TYPE_ATTRIBUTES, type_attributes, sk:transla 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; Tue, 22 Nov 2016 20:09:43 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1c9HNt-0005s9-V4 from Cesar_Philippidis@mentor.com ; Tue, 22 Nov 2016 12:09:42 -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; Tue, 22 Nov 2016 12:09:39 -0800 Subject: Re: [PATCH] OpenACC routines -- c++ front end To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek References: <5c7a3063-d600-a688-94ee-bfd6915ee044@codesourcery.com> CC: Thomas Schwinge From: Cesar Philippidis Message-ID: <11dce8a7-68f7-697c-5b9a-f701d3e1d4d4@codesourcery.com> Date: Tue, 22 Nov 2016 12:09:38 -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: <5c7a3063-d600-a688-94ee-bfd6915ee044@codesourcery.com> X-ClientProxiedBy: svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) On 11/11/2016 03:43 PM, Cesar Philippidis wrote: > Like it's c FE counterpart, this contains the following changes: > > * Updates c_parser_oacc_shape_clause to accept a location_t > argument in order to make the diagnostics more precise. > > * Adds support for the bind and nohost clauses. > > * Adds more diagnostics for invalid acc routines. > > Is this patch OK for trunk? Here is the updated version of the c++ OpenACC routine patch. It's mostly the same as before, but now cp_parser_oacc_shape_clause no has a dummy cp_parser argument like its c FE counterpart. Is this patch ok for trunk? Cesar 2016-11-22 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_simple_clause): Remove unused cp_parser argument. (cp_parser_oacc_shape_clause): New location_t loc argument. Use it to report more accurate diagnostics. Remove parser argument. (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 5674886..c9dbc4f 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -5785,6 +5785,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 6893eae..09f9ffc 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 9da4b30..15ead5a 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -30497,6 +30497,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; @@ -30576,6 +30580,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)) @@ -30991,8 +30997,7 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list) seq */ static tree -cp_parser_oacc_simple_clause (cp_parser * /* parser */, - enum omp_clause_code code, +cp_parser_oacc_simple_clause (enum omp_clause_code code, tree list, location_t location) { check_no_duplicate_clause (list, code, omp_clause_code_name[code], location); @@ -31052,13 +31057,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"; @@ -32845,6 +32850,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. */ @@ -32877,10 +32958,14 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "async"; break; case PRAGMA_OACC_CLAUSE_AUTO: - clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_AUTO, - clauses, here); + clauses = cp_parser_oacc_simple_clause (OMP_CLAUSE_AUTO, 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"; @@ -32928,7 +33013,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: @@ -32940,8 +33025,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "if"; break; case PRAGMA_OACC_CLAUSE_INDEPENDENT: - clauses = cp_parser_oacc_simple_clause (parser, - OMP_CLAUSE_INDEPENDENT, + clauses = cp_parser_oacc_simple_clause (OMP_CLAUSE_INDEPENDENT, clauses, here); c_name = "independent"; break; @@ -32949,6 +33033,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 (OMP_CLAUSE_NOHOST, clauses, + here); + c_name = "nohost"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: code = OMP_CLAUSE_NUM_GANGS; c_name = "num_gangs"; @@ -32995,8 +33084,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "self"; break; case PRAGMA_OACC_CLAUSE_SEQ: - clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, - clauses, here); + clauses = cp_parser_oacc_simple_clause (OMP_CLAUSE_SEQ, clauses, + here); c_name = "seq"; break; case PRAGMA_OACC_CLAUSE_TILE: @@ -33010,7 +33099,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: @@ -33025,7 +33115,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: @@ -37266,7 +37357,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 )' @@ -37325,6 +37419,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 @@ -37336,16 +37433,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); @@ -37421,6 +37508,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. */ @@ -37455,31 +37545,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 389e7f1..02c9ba2 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7074,6 +7074,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: