From patchwork Tue Nov 22 19:57:33 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 83483 Delivered-To: patch@linaro.org Received: by 10.140.97.165 with SMTP id m34csp2286729qge; Tue, 22 Nov 2016 11:58:12 -0800 (PST) X-Received: by 10.99.135.200 with SMTP id i191mr46814576pge.18.1479844691969; Tue, 22 Nov 2016 11:58:11 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id b5si29861732pgg.315.2016.11.22.11.58.11 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Tue, 22 Nov 2016 11:58:11 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-442280-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-442280-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-442280-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=b8lZ68+L3v0PJbec8 8TpxXpmH0ZseKfBNtICrx4zJLLGRiQJp0kOxEjZ0HJ+hEAl0oHTSjr2+ZndBaRLH tEet/eub30KXRF1JyYgRFBR0sZoszLNoJLxeApLr+Ps3r6NuASS3YV+mPCt7pAFR 1/0Fo0JjyPjfWIASryFUA0CN8M= 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=7b8YLckWIfxmJHzUaFCENpk mlyA=; b=g7xXZphvlJYdxJG2+36QK94wkjgimWp3a3/Bn7ZLIjJ0HHLOl+lI1Ob dEdB3Y8g6H+AQzpaYpwfanlFzriU2mSQITlz5jE9shw4s7z9a8D/SPvWDPnl+e7s 3YaFPlQQ+Wiy/4kt8wRuR2312+rjPrrx+Ltk83dh3L0yorQP2J/c= Received: (qmail 109060 invoked by alias); 22 Nov 2016 19:57:49 -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 108552 invoked by uid 89); 22 Nov 2016 19:57:49 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.7 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS, UNWANTED_LANGUAGE_BODY, URIBL_RED autolearn=ham version=3.3.2 spammy=ort, acc, unneeded, DECL_ATTRIBUTES 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 19:57:38 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1c9HCD-0002eX-DV from Cesar_Philippidis@mentor.com ; Tue, 22 Nov 2016 11:57:37 -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 11:57:34 -0800 Subject: Re: [PATCH] OpenACC routines -- c front end To: Jakub Jelinek References: <4ece4875-06cd-bfe7-d7db-7c002d98b8fb@codesourcery.com> <20161118122155.GY3541@tucnak.redhat.com> CC: "gcc-patches@gcc.gnu.org" , Thomas Schwinge From: Cesar Philippidis Message-ID: Date: Tue, 22 Nov 2016 11:57:33 -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: <20161118122155.GY3541@tucnak.redhat.com> X-ClientProxiedBy: svr-orw-mbx-03.mgc.mentorg.com (147.34.90.203) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) On 11/18/2016 04:21 AM, Jakub Jelinek wrote: > On Fri, Nov 11, 2016 at 03:43:23PM -0800, Cesar Philippidis wrote: >> @@ -11801,12 +11807,11 @@ c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind, >> } >> >> location_t expr_loc = c_parser_peek_token (parser)->location; >> - c_expr cexpr = c_parser_expr_no_commas (parser, NULL); >> - cexpr = convert_lvalue_to_rvalue (expr_loc, cexpr, false, true); >> - tree expr = cexpr.value; >> + tree expr = c_parser_expr_no_commas (parser, NULL).value; >> if (expr == error_mark_node) >> goto cleanup_error; >> >> + mark_exp_read (expr); >> expr = c_fully_fold (expr, false, NULL); >> >> /* Attempt to statically determine when the number isn't a > > Why? Are the arguments of the clauses lvalues? The spec is unclear if those args must be constants or not. The only time it explicitly mentions constant int-expr is for the tile clause, which was added late. Gang, worker and vector were added early in the 1.0 spec, where things were defined somewhat loosely. >> @@ -11867,12 +11872,12 @@ c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind, >> seq */ >> >> static tree >> -c_parser_oacc_simple_clause (c_parser *parser, enum omp_clause_code code, >> - tree list) >> +c_parser_oacc_simple_clause (c_parser * /* parser */, location_t loc, > > Just leave it as c_parser *, or better yet remove the argument if you don't > need it. I removed that argument. >> + else >> + { >> + //TODO? TREE_USED (decl) = 1; > > This would be /* FIXME: TREE_USED (decl) = 1; */ > but wouldn't it be better to figure out if you want to do that or not? Thomas has more state on that, but it seems unneeded. The c++ FE doesn't do that either, so I removed that comment. Is this patch ok for trunk? Cesar 2016-11-22 Cesar Philippidis Thomas Schwinge gcc/c/ * c-parser.c (c_parser_omp_clause_name): Handle OpenACC bind and nohost. (c_parser_oacc_shape_clause): New location_t loc argument. Use it to report more accurate diagnostics. (c_parser_oacc_simple_clause): Likewise. (c_parser_oacc_clause_bind): New function. (c_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}. (c_parser_oacc_routine): Update diagnostics. (c_finish_oacc_routine): Likewise. * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_{BIND,NOHOST}. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 00fe731..fd87b54 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -10408,6 +10408,10 @@ c_parser_omp_clause_name (c_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; @@ -10489,6 +10493,8 @@ c_parser_omp_clause_name (c_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 (!strcmp ("num_gangs", p)) result = PRAGMA_OACC_CLAUSE_NUM_GANGS; else if (!strcmp ("num_tasks", p)) @@ -11676,12 +11682,12 @@ c_parser_omp_clause_num_workers (c_parser *parser, tree list) */ static tree -c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind, +c_parser_oacc_shape_clause (c_parser *parser, location_t loc, + omp_clause_code kind, const char *str, tree list) { const char *id = "num"; tree ops[2] = { NULL_TREE, NULL_TREE }, c; - location_t loc = c_parser_peek_token (parser)->location; if (kind == OMP_CLAUSE_VECTOR) id = "length"; @@ -11746,12 +11752,11 @@ c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind, } location_t expr_loc = c_parser_peek_token (parser)->location; - c_expr cexpr = c_parser_expr_no_commas (parser, NULL); - cexpr = convert_lvalue_to_rvalue (expr_loc, cexpr, false, true); - tree expr = cexpr.value; + tree expr = c_parser_expr_no_commas (parser, NULL).value; if (expr == error_mark_node) goto cleanup_error; + mark_exp_read (expr); expr = c_fully_fold (expr, false, NULL); /* Attempt to statically determine when the number isn't a @@ -11812,12 +11817,12 @@ c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind, seq */ static tree -c_parser_oacc_simple_clause (c_parser *parser, enum omp_clause_code code, +c_parser_oacc_simple_clause (location_t loc, enum omp_clause_code code, tree list) { check_no_duplicate_clause (list, code, omp_clause_code_name[code]); - tree c = build_omp_clause (c_parser_peek_token (parser)->location, code); + tree c = build_omp_clause (loc, code); OMP_CLAUSE_CHAIN (c) = list; return c; @@ -11859,6 +11864,62 @@ c_parser_oacc_clause_async (c_parser *parser, tree list) } /* OpenACC 2.0: + bind ( identifier ) + bind ( string-literal ) */ + +static tree +c_parser_oacc_clause_bind (c_parser *parser, tree list) +{ + check_no_duplicate_clause (list, OMP_CLAUSE_BIND, "bind"); + + location_t loc = c_parser_peek_token (parser)->location; + + parser->lex_untranslated_string = true; + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + { + parser->lex_untranslated_string = false; + return list; + } + tree name = error_mark_node; + c_token *token = c_parser_peek_token (parser); + if (c_parser_next_token_is (parser, CPP_NAME)) + { + tree decl = lookup_name (token->value); + if (!decl) + error_at (token->location, "%qE has not been declared", + token->value); + else if (TREE_CODE (decl) != FUNCTION_DECL) + error_at (token->location, "%qE does not refer to a function", + token->value); + else + { + tree name_id = DECL_NAME (decl); + name = build_string (IDENTIFIER_LENGTH (name_id), + IDENTIFIER_POINTER (name_id)); + } + c_parser_consume_token (parser); + } + else if (c_parser_next_token_is (parser, CPP_STRING)) + { + name = token->value; + c_parser_consume_token (parser); + } + else + c_parser_error (parser, + "expected identifier or character string literal"); + parser->lex_untranslated_string = false; + c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + 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; +} + +/* OpenACC 2.0: tile ( size-expr-list ) */ static tree @@ -13160,10 +13221,14 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "async"; break; case PRAGMA_OACC_CLAUSE_AUTO: - clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_AUTO, + clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_AUTO, clauses); c_name = "auto"; break; + case PRAGMA_OACC_CLAUSE_BIND: + clauses = c_parser_oacc_clause_bind (parser, clauses); + c_name = "bind"; + break; case PRAGMA_OACC_CLAUSE_COLLAPSE: clauses = c_parser_omp_clause_collapse (parser, clauses); c_name = "collapse"; @@ -13210,7 +13275,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, break; case PRAGMA_OACC_CLAUSE_GANG: c_name = "gang"; - clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG, + clauses = c_parser_oacc_shape_clause (parser, here, OMP_CLAUSE_GANG, c_name, clauses); break; case PRAGMA_OACC_CLAUSE_HOST: @@ -13222,14 +13287,19 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "if"; break; case PRAGMA_OACC_CLAUSE_INDEPENDENT: - clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_INDEPENDENT, - clauses); + clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_INDEPENDENT, + clauses); c_name = "independent"; break; case PRAGMA_OACC_CLAUSE_LINK: clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "link"; break; + case PRAGMA_OACC_CLAUSE_NOHOST: + clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_NOHOST, + clauses); + c_name = "nohost"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: clauses = c_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; @@ -13271,8 +13341,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "self"; break; case PRAGMA_OACC_CLAUSE_SEQ: - clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, - clauses); + clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_SEQ, clauses); c_name = "seq"; break; case PRAGMA_OACC_CLAUSE_TILE: @@ -13285,7 +13354,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; - clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, + clauses = c_parser_oacc_shape_clause (parser, here, OMP_CLAUSE_VECTOR, c_name, clauses); break; case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: @@ -13298,7 +13367,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, break; case PRAGMA_OACC_CLAUSE_WORKER: c_name = "worker"; - clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_WORKER, + clauses = c_parser_oacc_shape_clause (parser, here, OMP_CLAUSE_WORKER, c_name, clauses); break; default: @@ -14092,7 +14161,9 @@ c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, ( (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 an OpenACC routine directive. For named directives, we apply immediately to the named function. For unnamed ones we then parse @@ -14142,6 +14213,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context) data.clauses = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine"); + /* The clauses are in reverse order; fix that to make later diagnostic + emission easier. */ + data.clauses = nreverse (data.clauses); if (TREE_CODE (decl) != FUNCTION_DECL) { @@ -14156,6 +14230,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context) data.clauses = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine"); + /* The clauses are in reverse order; fix that to make later diagnostic + emission easier. */ + data.clauses = nreverse (data.clauses); /* Emit a helpful diagnostic if there's another pragma following this one. Also don't allow a static assertion declaration, as in the @@ -14219,31 +14296,37 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl, return; } - if (get_oacc_fn_attrib (fndecl)) + int compatible + = verify_oacc_routine_clauses (fndecl, &data->clauses, data->loc, + "#pragma acc routine"); + if (compatible < 0) { - error_at (data->loc, - "%<#pragma acc routine%> already applied to %qD", fndecl); data->error_seen = true; return; } - - if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + if (compatible > 0) { - error_at (data->loc, - "%<#pragma acc routine%> must be applied before %s", - TREE_USED (fndecl) ? "use" : "definition"); - data->error_seen = true; - return; } + else + { + if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl))) + { + error_at (data->loc, + "%<#pragma acc routine%> must be applied before %s", + TREE_USED (fndecl) ? "use" : "definition"); + data->error_seen = true; + return; + } - /* Process the routine's dimension clauses. */ - tree dims = build_oacc_routine_dims (data->clauses); - replace_oacc_fn_attrib (fndecl, dims); + /* Set the routine's level of parallelism. */ + tree dims = build_oacc_routine_dims (data->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)); + /* Add an "omp declare target" attribute. */ + DECL_ATTRIBUTES (fndecl) + = tree_cons (get_identifier ("omp declare target"), + data->clauses, DECL_ATTRIBUTES (fndecl)); + } /* Remember that we've used this "#pragma acc routine". */ data->fndecl_seen = true; diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index f0917ed..fadec8c 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13580,6 +13580,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: case OMP_CLAUSE_TILE: pc = &OMP_CLAUSE_CHAIN (c); continue;