From patchwork Fri Dec 2 14:38:25 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 86309 Delivered-To: patch@linaro.org Received: by 10.182.112.6 with SMTP id im6csp350351obb; Fri, 2 Dec 2016 06:38:59 -0800 (PST) X-Received: by 10.84.215.138 with SMTP id l10mr96990432pli.166.1480689539331; Fri, 02 Dec 2016 06:38:59 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id v21si5467724pfa.41.2016.12.02.06.38.58 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Fri, 02 Dec 2016 06:38:59 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-443336-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-443336-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-443336-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=sWJYiygNX2esCJ6BI VkbrJO0tVTfgqjp0Tv9PX6d28IhBiTPc6y8P/idrIWyD95v98UgAPIlo0cY5ofVP S5owQ4NglepKeWpkEhfmZ5yvEhQP1A5oa2wtkeGfcXne8AC3ZCmVwMh1nZs9StC2 r5Y1mzuf5V+JjfZEUbq0oZh1f8= 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=iWm4XGAMibaQW4NmxUiKU+v J984=; b=HGSyHtSxWo6NOWZfC61uM+CQNO6FVWQ0QF+JL1RwU0uDHM8eBcgwLbS JD0etAu2Ne9PE2NE0EDdty6W/EqMi6xAzQ60O90DvCby218VO2+XcvXP/+F5j7u7 Tl9kewWMKhJd6kAMBXWFGTJH1HlopZZX8tqHTzXAg+xqvJ2vPLnU= Received: (qmail 62815 invoked by alias); 2 Dec 2016 14:38:41 -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 62782 invoked by uid 89); 2 Dec 2016 14:38:40 -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=Cesar, norris, Norris, 2712 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, 02 Dec 2016 14:38:29 +0000 Received: from svr-orw-mbx-04.mgc.mentorg.com ([147.34.90.204]) by relay1.mentorg.com with esmtp id 1cCoyq-0007Lj-Bw from Cesar_Philippidis@mentor.com ; Fri, 02 Dec 2016 06:38:28 -0800 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Fri, 2 Dec 2016 06:38:25 -0800 Subject: Re: [PATCH] OpenACC executable directives To: "gcc-patches@gcc.gnu.org" References: <12adb858-eb15-0a6d-2049-d000bdd96de7@codesourcery.com> CC: Jakub Jelinek From: Cesar Philippidis Message-ID: Date: Fri, 2 Dec 2016 06:38:25 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.4.0 MIME-Version: 1.0 In-Reply-To: <12adb858-eb15-0a6d-2049-d000bdd96de7@codesourcery.com> X-ClientProxiedBy: svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) To SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) On 12/02/2016 06:37 AM, Cesar Philippidis wrote: > This patch teaches the C and C++ FEs to expect ACC ENTER/EXIT DATA, ACC > UPDATE and ACC WAIT executable directives to be used inside compound > statements. This is to prevent situations such as > > if (needs_wait) > #pragma acc wait > > // do something else here > > from generating unexpected code when the program is built without > -fopenacc. The C and C++ FEs already guard against such usages for > OpenMP executable directives. > > This patch also included a tweak for the diagnostics generated by > oacc_enter_exit_data parser. The error message now more accurately > reports if an error is detected in an enter or exit data directive. > > Is this patch OK for trunk? And here's the patch. Cesar 2016-12-02 Cesar Philippidis James Norris gcc/c/ * c-parser.c (c_parser_pragma): Error when PRAGMA_OACC_{ENTER_DATA, EXIT_DATA,WAIT} are not used in compound statements. (c_parser_oacc_enter_exit_data): Update diagnostics. gcc/cp/ * parser.c (cp_parser_oacc_enter_exit_data): Update diagnostics. (cp_parser_pragma): Error when PRAGMA_OACC_{ENTER_DATA, EXIT_DATA,WAIT} are not used in compound statements. gcc/testsuite/ * c-c++-common/goacc/data-2.c: Adjust test. * c-c++-common/goacc/executeables-1.c: New test. * g++.dg/goacc/data-1.C: Adjust test. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 00fe731..f7bf9c4 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -10145,10 +10145,24 @@ c_parser_pragma (c_parser *parser, enum pragma_context context, bool *if_p) return false; case PRAGMA_OACC_ENTER_DATA: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc enter data%> may only be " + "used in compound statements"); + goto bad_stmt; + } c_parser_oacc_enter_exit_data (parser, true); return false; case PRAGMA_OACC_EXIT_DATA: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc exit data%> may only be " + "used in compound statements"); + goto bad_stmt; + } c_parser_oacc_enter_exit_data (parser, false); return false; @@ -10305,6 +10319,16 @@ c_parser_pragma (c_parser *parser, enum pragma_context context, bool *if_p) c_parser_cilk_grainsize (parser, if_p); return false; + case PRAGMA_OACC_WAIT: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc enter data%> may only be " + "used in compound statements"); + goto bad_stmt; + } + /* FALL THROUGH. */ + default: if (id < PRAGMA_FIRST_EXTERNAL) { @@ -13871,28 +13895,26 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter) { location_t loc = c_parser_peek_token (parser)->location; tree clauses, stmt; + const char *p = ""; c_parser_consume_pragma (parser); - if (!c_parser_next_token_is (parser, CPP_NAME)) + if (c_parser_next_token_is (parser, CPP_NAME)) { - c_parser_error (parser, enter - ? "expected % in %<#pragma acc enter data%>" - : "expected % in %<#pragma acc exit data%>"); - c_parser_skip_to_pragma_eol (parser); - return; + p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + c_parser_consume_token (parser); } - const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); if (strcmp (p, "data") != 0) { - c_parser_error (parser, "invalid pragma"); + error_at (loc, enter + ? "expected % after %<#pragma acc enter%>" + : "expected % after %<#pragma acc exit%>"); + parser->error = true; c_parser_skip_to_pragma_eol (parser); return; } - c_parser_consume_token (parser); - if (enter) clauses = c_parser_oacc_all_clauses (parser, OACC_ENTER_DATA_CLAUSE_MASK, "#pragma acc enter data"); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 843cbe2..08f5f9e 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -36175,23 +36175,18 @@ static tree cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok, bool enter) { + location_t loc = pragma_tok->location; tree stmt, clauses; + const char *p = ""; - if (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA_EOL) - || cp_lexer_next_token_is_not (parser->lexer, CPP_NAME)) - { - cp_parser_error (parser, enter - ? "expected % in %<#pragma acc enter data%>" - : "expected % in %<#pragma acc exit data%>"); - cp_parser_skip_to_pragma_eol (parser, pragma_tok); - return NULL_TREE; - } + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + p = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); - const char *p = - IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); if (strcmp (p, "data") != 0) { - cp_parser_error (parser, "invalid pragma"); + error_at (loc, enter + ? "expected % after %<#pragma acc enter%>" + : "expected % after %<#pragma acc exit%>"); cp_parser_skip_to_pragma_eol (parser, pragma_tok); return NULL_TREE; } @@ -36207,8 +36202,8 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok, if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) { - error_at (pragma_tok->location, - "%<#pragma acc enter data%> has no data movement clause"); + error_at (loc, "%<#pragma acc %s data%> has no data movement clause", + enter ? "enter" : "exit"); return NULL_TREE; } @@ -38083,6 +38078,30 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p) cp_parser_oacc_declare (parser, pragma_tok); return false; + case PRAGMA_OACC_ENTER_DATA: + if (context == pragma_stmt) + { + cp_parser_error (parser, "%<#pragma acc enter data%> may only be " + "used in compound statements"); + break; + } + else if (context != pragma_compound) + goto bad_stmt; + cp_parser_omp_construct (parser, pragma_tok, if_p); + return true; + + case PRAGMA_OACC_EXIT_DATA: + if (context == pragma_stmt) + { + cp_parser_error (parser, "%<#pragma acc exit data%> may only be " + "used in compound statements"); + break; + } + else if (context != pragma_compound) + goto bad_stmt; + cp_parser_omp_construct (parser, pragma_tok, if_p); + return true; + case PRAGMA_OACC_ROUTINE: if (context != pragma_external) { @@ -38093,17 +38112,37 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p) cp_parser_oacc_routine (parser, pragma_tok, context); return false; + case PRAGMA_OACC_UPDATE: + if (context == pragma_stmt) + { + cp_parser_error (parser, "%<#pragma acc update%> may only be " + "used in compound statements"); + break; + } + else if (context != pragma_compound) + goto bad_stmt; + cp_parser_omp_construct (parser, pragma_tok, if_p); + return true; + + case PRAGMA_OACC_WAIT: + if (context == pragma_stmt) + { + cp_parser_error (parser, "%<#pragma acc wait%> may only be " + "used in compound statements"); + break; + } + else if (context != pragma_compound) + goto bad_stmt; + cp_parser_omp_construct (parser, pragma_tok, if_p); + return true; + case PRAGMA_OACC_ATOMIC: case PRAGMA_OACC_CACHE: case PRAGMA_OACC_DATA: - case PRAGMA_OACC_ENTER_DATA: - case PRAGMA_OACC_EXIT_DATA: case PRAGMA_OACC_HOST_DATA: case PRAGMA_OACC_KERNELS: case PRAGMA_OACC_PARALLEL: case PRAGMA_OACC_LOOP: - case PRAGMA_OACC_UPDATE: - case PRAGMA_OACC_WAIT: case PRAGMA_OMP_ATOMIC: case PRAGMA_OMP_CRITICAL: case PRAGMA_OMP_DISTRIBUTE: diff --git a/gcc/testsuite/c-c++-common/goacc/data-2.c b/gcc/testsuite/c-c++-common/goacc/data-2.c index a67d8a4..1043bf8a 100644 --- a/gcc/testsuite/c-c++-common/goacc/data-2.c +++ b/gcc/testsuite/c-c++-common/goacc/data-2.c @@ -10,12 +10,14 @@ foo (void) #pragma acc exit data delete (a) if (0) #pragma acc exit data copyout (b) if (a) #pragma acc exit data delete (b) -#pragma acc enter /* { dg-error "expected 'data' in" } */ -#pragma acc exit /* { dg-error "expected 'data' in" } */ +#pragma acc enter /* { dg-error "expected 'data' after" } */ +#pragma acc exit /* { dg-error "expected 'data' after" } */ #pragma acc enter data /* { dg-error "has no data movement clause" } */ -#pragma acc exit data /* { dg-error "has no data movement clause" } */ -#pragma acc enter Data /* { dg-error "invalid pragma before" } */ -#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */ +#pragma acc exit data /* { dg-error "no data movement clause" } */ +#pragma acc enter Data /* { dg-error "expected 'data' after" } */ +#pragma acc exit copyout (b) /* { dg-error "expected 'data' after" } */ +#pragma acc enter for /* { dg-error "expected 'data' after" } */ +#pragma acc enter data2 /* { dg-error "expected 'data' after" } */ } /* { dg-error "has no data movement clause" "" { target *-*-* } 8 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/executeables-1.c b/gcc/testsuite/c-c++-common/goacc/executeables-1.c new file mode 100644 index 0000000..da89437 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/executeables-1.c @@ -0,0 +1,74 @@ +/* { dg-do compile } */ + +void +foo (void) +{ + const int N = 32; + float x[N], y[N]; + int flag = 0; + + if (flag) +#pragma acc update host (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 1; + + while (flag) +#pragma acc update host (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 2; + +#pragma acc enter data create (x[0:N]) + { + if (flag) +#pragma acc update host (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 1; + } + + if (flag) + while (flag) +#pragma acc update host (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 2; + + if (flag) +#pragma acc wait /* { dg-error "may only be used in compound statements" } */ + flag = 1; + + while (flag) +#pragma acc wait /* { dg-error "may only be used in compound statements" } */ + flag = 2; + +#pragma acc enter data create (x[0:N]) + { + if (flag) +#pragma acc wait /* { dg-error "may only be used in compound statements" } */ + flag = 1; + } + + if (flag) +#pragma acc enter data create (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 1; + + while (flag) +#pragma acc enter data create (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 2; + +#pragma acc enter data create (x[0:N]) + { + if (flag) +#pragma acc enter data create (y[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 1; + } + + if (flag) +#pragma acc exit data delete (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 1; + + while (flag) +#pragma acc exit data delete (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 2; + +#pragma acc enter data create (x[0:N]) + { + if (flag) +#pragma acc exit data delete (x[0:N]) /* { dg-error "may only be used in compound statements" } */ + flag = 1; + } +} diff --git a/gcc/testsuite/g++.dg/goacc/data-1.C b/gcc/testsuite/g++.dg/goacc/data-1.C index 54676dc..2b210dc 100644 --- a/gcc/testsuite/g++.dg/goacc/data-1.C +++ b/gcc/testsuite/g++.dg/goacc/data-1.C @@ -8,12 +8,12 @@ foo (int &a, int (&b)[100], int &n) #pragma acc exit data delete (a) if (0) #pragma acc exit data copyout (b) if (a) #pragma acc exit data delete (b) -#pragma acc enter /* { dg-error "expected 'data' in" } */ -#pragma acc exit /* { dg-error "expected 'data' in" } */ +#pragma acc enter /* { dg-error "expected 'data' after" } */ +#pragma acc exit /* { dg-error "expected 'data' after" } */ #pragma acc enter data /* { dg-error "has no data movement clause" } */ #pragma acc exit data /* { dg-error "has no data movement clause" } */ -#pragma acc enter Data /* { dg-error "invalid pragma before" } */ -#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */ +#pragma acc enter Data /* { dg-error "expected 'data' after" } */ +#pragma acc exit copyout (b) /* { dg-error "expected 'data' after" } */ } template @@ -27,12 +27,12 @@ foo (T &a, T (&b)[100], T &n) #pragma acc exit data delete (a) if (0) #pragma acc exit data copyout (b) if (a) #pragma acc exit data delete (b) -#pragma acc enter /* { dg-error "expected 'data' in" } */ -#pragma acc exit /* { dg-error "expected 'data' in" } */ +#pragma acc enter /* { dg-error "expected 'data' after" } */ +#pragma acc exit /* { dg-error "expected 'data' after" } */ #pragma acc enter data /* { dg-error "has no data movement clause" } */ #pragma acc exit data /* { dg-error "has no data movement clause" } */ -#pragma acc enter Data /* { dg-error "invalid pragma before" } */ -#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */ +#pragma acc enter Data /* { dg-error "expected 'data' after" } */ +#pragma acc exit copyout (b) /* { dg-error "expected 'data' after" } */ } /* { dg-error "has no data movement clause" "" { target *-*-* } 6 } */