OpenACC executable directives

Message ID	cf5dfd4e-e53a-5632-a3bf-1b41f2e6d611@mentor.com
State	New
Headers	show Delivered-To: patch@linaro.org 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; 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= Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk Sender: gcc-patches-owner@gcc.gnu.org Subject: Re: [PATCH] OpenACC executable directives To: "gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org> References: <12adb858-eb15-0a6d-2049-d000bdd96de7@codesourcery.com> CC: Jakub Jelinek <jakub@redhat.com> From: Cesar Philippidis <cesar_philippidis@mentor.com> Message-ID: <cf5dfd4e-e53a-5632-a3bf-1b41f2e6d611@mentor.com> 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> Content-Type: multipart/mixed; boundary="------------3DF9D214C7C04B38AEC00296"

2016-12-02 Cesar Philippidis <cesar@codesourcery.com> James Norris <jnorris@codesourcery.com> 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 %<data%> in %<#pragma acc enter data%>" - : "expected %<data%> 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 %<data%> after %<#pragma acc enter%>" + : "expected %<data%> 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 %<data%> in %<#pragma acc enter data%>" - : "expected %<data%> 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 %<data%> after %<#pragma acc enter%>" + : "expected %<data%> 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<typename T> @@ -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 } */

OpenACC executable directives

Commit Message

Comments

Patch