From patchwork Fri Nov 11 23:43:23 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 81906 Delivered-To: patch@linaro.org Received: by 10.140.97.165 with SMTP id m34csp9515qge; Fri, 11 Nov 2016 15:43:54 -0800 (PST) X-Received: by 10.99.54.74 with SMTP id d71mr8791840pga.34.1478907834208; Fri, 11 Nov 2016 15:43:54 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id u2si10545559pau.333.2016.11.11.15.43.53 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Fri, 11 Nov 2016 15:43:54 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-441209-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-441209-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-441209-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:from :subject:to:cc:message-id:date:mime-version:content-type; q=dns; s=default; b=Wm9Db23jkXhRnkuDcoUHJAyxZXJ0flBBBe03kXPorNf1sOTOX3 FZmF5AYY9XcsRQ4h6S3xJzBUb0umqKEulyhimjOvlSXcmm9Pw5n1tOIYIrRuZt2r xVE4LJi8aX7csn7A1MvPDvooCQ55b5a1uk1fYepweBdP0lvNDppuDso3E= 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:from :subject:to:cc:message-id:date:mime-version:content-type; s= default; bh=+ITknvAl1yoDiVitVZybt7i7S3M=; b=d6AWXpTWLgenCAyY87YI C1ARlByW3KJ74NWp8tIUR+YQ9Gl43D3CMZV63YqXgWjx1ssNMScISvxj5FfrH/xR ED9qvvI5hk2mdtpqMn4fj/yFwuqTlzrBflT73jU/l5jvhkcgHBr9eWmIq/XZZ1Ku r//9WFQmLN5NX6qPQPtbnnI= Received: (qmail 52598 invoked by alias); 11 Nov 2016 23:43:39 -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 52588 invoked by uid 89); 11 Nov 2016 23:43:38 -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=sk:!c_pars, c_token, CPP_STRING, emission 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; Fri, 11 Nov 2016 23:43:28 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1c5LTj-0002oO-0X from Cesar_Philippidis@mentor.com ; Fri, 11 Nov 2016 15:43:27 -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; Fri, 11 Nov 2016 15:43:24 -0800 From: Cesar Philippidis Subject: [PATCH] OpenACC routines -- c front end To: "gcc-patches@gcc.gnu.org" , Jakub Jelinek CC: Thomas Schwinge Message-ID: <4ece4875-06cd-bfe7-d7db-7c002d98b8fb@codesourcery.com> Date: Fri, 11 Nov 2016 15:43:23 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.3.0 MIME-Version: 1.0 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) This contains the following changes: * Updates c_parser_oacc_{shape,simple}_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? Cesar 2016-11-11 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 6bc42da..072af5d 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -10463,6 +10463,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; @@ -10544,6 +10548,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)) @@ -11731,12 +11737,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"; @@ -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 @@ -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, + 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; @@ -11914,6 +11919,63 @@ 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 + { + //TODO? TREE_USED (decl) = 1; + 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 @@ -13215,10 +13277,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 (parser, 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"; @@ -13265,7 +13331,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: @@ -13277,14 +13343,20 @@ 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 (parser, 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 (parser, 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"; @@ -13326,7 +13398,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 = c_parser_oacc_simple_clause (parser, here, OMP_CLAUSE_SEQ, clauses); c_name = "seq"; break; @@ -13340,7 +13412,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: @@ -13353,7 +13425,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: @@ -14147,7 +14219,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 @@ -14197,6 +14271,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) { @@ -14211,6 +14288,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 @@ -14274,31 +14354,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;