Hi! Ping.
On Thu, 2 Jun 2016 13:47:08 +0200, I wrote: > On Wed, 05 Nov 2014 17:29:19 +0100, I wrote: > > In r217145, I applied Jim's patch to gomp-4_0-branch: > > > > commit 4361f9b6b2c74c2961c3a5290a4945abe2d7a444 > > Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> > > Date: Wed Nov 5 16:26:47 2014 +0000 > > > > OpenACC cache directive for C. > > (That, and the corresponding C++ changes later made it into trunk.) > > > --- gcc/c/c-parser.c > > +++ gcc/c/c-parser.c > > @@ -10053,6 +10053,14 @@ c_parser_omp_variable_list (c_parser *parser, > > { > > switch (kind) > > { > > + case OMP_NO_CLAUSE_CACHE: > > + if (c_parser_peek_token (parser)->type != CPP_OPEN_SQUARE) > > + { > > + c_parser_error (parser, "expected %<[%>"); > > + t = error_mark_node; > > + break; > > + } > > + /* FALL THROUGH. */ > > case OMP_CLAUSE_MAP: > > case OMP_CLAUSE_FROM: > > case OMP_CLAUSE_TO: > > Strictly speaking (OpenACC 2.0a specification), that is correct: the > OpenACC cache directive explicitly only allows "array elements or > subarrays". However, I wonder if it would make sense to allow complete > arrays as a GNU extension? That is, syntactic sugar to allow "cache (a)" > to mean "cache (a[0:LENGTH])"? > > > @@ -10091,6 +10099,29 @@ c_parser_omp_variable_list (c_parser *parser, > > t = error_mark_node; > > break; > > } > > + > > + if (kind == OMP_NO_CLAUSE_CACHE) > > + { > > + mark_exp_read (low_bound); > > + mark_exp_read (length); > > + > > + if (TREE_CODE (low_bound) != INTEGER_CST > > + && !TREE_READONLY (low_bound)) > > + { > > + error_at (clause_loc, > > + "%qD is not a constant", low_bound); > > + t = error_mark_node; > > + } > > WHile OpenACC 2.0a specifies that "the lower bound is a constant", it > also permits the lower bound to be a "loop invariant, or the for loop > index variable plus or minus a constant or loop invariant". So, we're > rejecting valid syntax here. > > > + > > + if (TREE_CODE (length) != INTEGER_CST > > + && !TREE_READONLY (length)) > > + { > > + error_at (clause_loc, > > + "%qD is not a constant", length); > > + t = error_mark_node; > > + } > > + } > > The idea is correct (OpenACC 2.0a: "the length is a constant"), but we > can't reliably check that here; for example: > > #pragma acc cache (a[0:n + 1]) > > ... will run into an ICE, "tree check: expected tree that contains 'decl > minimal' structure, have 'plus_expr' in [...]". > > Currently we're discarding the OpenACC cache directive in the middle end; > I expect checking of the lower bound and length will come automatically > as soon as we start to do something with OACC_CACHE/OMP_CLAUSE__CACHE_. > Until then, I propose we simple remove these checks from the front ends. > OK for trunk and gcc-6-branch? > > commit a620ebe6fa509ec6441ba87276e55078eb2d00fc > Author: Thomas Schwinge <tho...@codesourcery.com> > Date: Thu Jun 2 12:19:49 2016 +0200 > > [PR c/71381] C/C++ OpenACC cache directive rejects valid syntax > > gcc/c/ > PR c/71381 > * c-parser.c (c_parser_omp_variable_list) <OMP_CLAUSE__CACHE_>: > Loosen checking. > gcc/cp/ > PR c/71381 > * parser.c (cp_parser_omp_var_list_no_open) <OMP_CLAUSE__CACHE_>: > Loosen checking. > gcc/fortran/ > PR c/71381 > * openmp.c (gfc_match_oacc_cache): Add comment. > gcc/testsuite/ > PR c/71381 > * c-c++-common/goacc/cache-1.c: Update. Move invalid usage tests > to... > * c-c++-common/goacc/cache-2.c: ... this new file. > * gfortran.dg/goacc/cache-1.f95: Move invalid usage tests to... > * gfortran.dg/goacc/cache-2.f95: ... this new file. > * gfortran.dg/goacc/coarray.f95: Update OpenACC cache directive > usage. > * gfortran.dg/goacc/cray.f95: Likewise. > * gfortran.dg/goacc/loop-1.f95: Likewise. > libgomp/ > PR c/71381 > * testsuite/libgomp.oacc-c-c++-common/cache-1.c: #include > "../../../gcc/testsuite/c-c++-common/goacc/cache-1.c". > * testsuite/libgomp.oacc-fortran/cache-1.f95: New file. > > gcc/ > * omp-low.c (scan_sharing_clauses): Don't expect > OMP_CLAUSE__CACHE_. > --- > gcc/c/c-parser.c | 22 +------- > gcc/cp/parser.c | 22 +------- > gcc/fortran/openmp.c | 5 ++ > gcc/omp-low.c | 6 -- > gcc/testsuite/c-c++-common/goacc/cache-1.c | 66 > ++++++++-------------- > .../c-c++-common/goacc/{cache-1.c => cache-2.c} | 39 ++----------- > gcc/testsuite/gfortran.dg/goacc/cache-1.f95 | 7 +-- > gcc/testsuite/gfortran.dg/goacc/cache-2.f95 | 12 ++++ > gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 2 +- > gcc/testsuite/gfortran.dg/goacc/cray.f95 | 3 +- > gcc/testsuite/gfortran.dg/goacc/loop-1.f95 | 7 ++- > .../testsuite/libgomp.oacc-c-c++-common/cache-1.c | 49 +--------------- > libgomp/testsuite/libgomp.oacc-fortran/cache-1.f95 | 5 ++ > 13 files changed, 66 insertions(+), 179 deletions(-) > > diff --git gcc/c/c-parser.c gcc/c/c-parser.c > index bca8653..1db1886 100644 > --- gcc/c/c-parser.c > +++ gcc/c/c-parser.c > @@ -10601,6 +10601,9 @@ c_parser_omp_variable_list (c_parser *parser, > switch (kind) > { > case OMP_CLAUSE__CACHE_: > + /* The OpenACC cache directive explicitly only allows "array > + elements or subarrays". Would it make sense to allow complete > + arrays as a GNU extension? */ > if (c_parser_peek_token (parser)->type != CPP_OPEN_SQUARE) > { > c_parser_error (parser, "expected %<[%>"); > @@ -10663,25 +10666,6 @@ c_parser_omp_variable_list (c_parser *parser, > break; > } > > - if (kind == OMP_CLAUSE__CACHE_) > - { > - if (TREE_CODE (low_bound) != INTEGER_CST > - && !TREE_READONLY (low_bound)) > - { > - error_at (clause_loc, > - "%qD is not a constant", low_bound); > - t = error_mark_node; > - } > - > - if (TREE_CODE (length) != INTEGER_CST > - && !TREE_READONLY (length)) > - { > - error_at (clause_loc, > - "%qD is not a constant", length); > - t = error_mark_node; > - } > - } > - > t = tree_cons (low_bound, length, t); > } > break; > diff --git gcc/cp/parser.c gcc/cp/parser.c > index 29a1b80..826277d 100644 > --- gcc/cp/parser.c > +++ gcc/cp/parser.c > @@ -29984,6 +29984,9 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, > enum omp_clause_code kind, > switch (kind) > { > case OMP_CLAUSE__CACHE_: > + /* The OpenACC cache directive explicitly only allows "array > + elements or subarrays". Would it make sense to allow complete > + arrays as a GNU extension? */ > if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_SQUARE) > { > error_at (token->location, "expected %<[%>"); > @@ -30035,25 +30038,6 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, > enum omp_clause_code kind, > RT_CLOSE_SQUARE)) > goto skip_comma; > > - if (kind == OMP_CLAUSE__CACHE_) > - { > - if (TREE_CODE (low_bound) != INTEGER_CST > - && !TREE_READONLY (low_bound)) > - { > - error_at (token->location, > - "%qD is not a constant", low_bound); > - decl = error_mark_node; > - } > - > - if (TREE_CODE (length) != INTEGER_CST > - && !TREE_READONLY (length)) > - { > - error_at (token->location, > - "%qD is not a constant", length); > - decl = error_mark_node; > - } > - } > - > decl = tree_cons (low_bound, length, decl); > } > break; > diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c > index 2689d30..d97c193 100644 > --- gcc/fortran/openmp.c > +++ gcc/fortran/openmp.c > @@ -1688,6 +1688,11 @@ match > gfc_match_oacc_cache (void) > { > gfc_omp_clauses *c = gfc_get_omp_clauses (); > + /* The OpenACC cache directive explicitly only allows "array elements or > + subarrays", which we're currently not checking here. Either check this > + after the call of gfc_match_omp_variable_list, or add something like a > + only_sections variant next to its allow_sections parameter. Or, would > it > + make sense to allow complete arrays as a GNU extension? */ > match m = gfc_match_omp_variable_list (" (", > &c->lists[OMP_LIST_CACHE], true, > NULL, NULL, true); > diff --git gcc/omp-low.c gcc/omp-low.c > index 77bdb18..91d5fcf 100644 > --- gcc/omp-low.c > +++ gcc/omp-low.c > @@ -2201,9 +2201,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, > break; > > case OMP_CLAUSE__CACHE_: > - sorry ("Clause not supported yet"); > - break; > - > default: > gcc_unreachable (); > } > @@ -2368,9 +2365,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, > break; > > case OMP_CLAUSE__CACHE_: > - sorry ("Clause not supported yet"); > - break; > - > default: > gcc_unreachable (); > } > diff --git gcc/testsuite/c-c++-common/goacc/cache-1.c > gcc/testsuite/c-c++-common/goacc/cache-1.c > index 9503341..1d4759e 100644 > --- gcc/testsuite/c-c++-common/goacc/cache-1.c > +++ gcc/testsuite/c-c++-common/goacc/cache-1.c > @@ -1,3 +1,7 @@ > +/* OpenACC cache directive: valid usage. */ > +/* For execution testing, this file is "#include"d from > + libgomp/testsuite/libgomp.oacc-c-c++-common/cache-1.c. */ > + > int > main (int argc, char **argv) > { > @@ -21,57 +25,31 @@ main (int argc, char **argv) > int n = 1; > const int len = n; > > -#pragma acc cache /* { dg-error "expected '\\\(' before end of line" } */ > - > -#pragma acc cache a[0:N] /* { dg-error "expected '\\\(' before 'a'" } */ > - /* { dg-bogus "expected end of line before 'a'" "" { xfail c++ } 26 } */ > - > -#pragma acc cache (a) /* { dg-error "expected '\\\['" } */ > - > -#pragma acc cache ( /* { dg-error "expected (identifier|unqualified-id) > before end of line" } */ > - > -#pragma acc cache () /* { dg-error "expected (identifier|unqualified-id) > before '\\\)' token" } */ > - > -#pragma acc cache (,) /* { dg-error "expected (identifier|unqualified-id) > before '(,|\\\))' token" } */ > - > -#pragma acc cache (a[0:N] /* { dg-error "expected '\\\)' before end of line" > } */ > - > -#pragma acc cache (a[0:N],) /* { dg-error "expected > (identifier|unqualified-id) before '(,|\\\))' token" "" { xfail c } } */ > - > -#pragma acc cache (a[0:N]) copyin (a[0:N]) /* { dg-error "expected end of > line before 'copyin'" } */ > - > -#pragma acc cache () /* { dg-error "expected (identifier|unqualified-id) > before '\\\)' token" } */ > - > -#pragma acc cache (a[0:N] b[0:N]) /* { dg-error "expected '\\\)' before 'b'" > } */ > - > -#pragma acc cache (a[0:N] b[0:N}) /* { dg-error "expected '\\\)' before 'b'" > } */ > - /* { dg-bogus "expected end of line before '\\\}' token" "" { xfail c++ > } 47 } */ > - > -#pragma acc cache (a[0:N] /* { dg-error "expected '\\\)' before end of line" > } */ > - > -#pragma acc cache (a[ii]) /* { dg-error "'ii' is not a constant" } */ > - > -#pragma acc cache (a[idx:n]) /* { dg-error "'n' is not a constant" } */ > - > -#pragma acc cache (a[0:N]) ( /* { dg-error "expected end of line before > '\\(' token" } */ > - > -#pragma acc cache (a[0:N]) ii /* { dg-error "expected end of line before > 'ii'" } */ > - > -#pragma acc cache (a[0:N] ii) /* { dg-error "expected '\\)' before 'ii'" } */ > - > + /* Have at it, GCC! */ > #pragma acc cache (a[0:N]) > - > #pragma acc cache (a[0:N], a[0:N]) > - > #pragma acc cache (a[0:N], b[0:N]) > - > #pragma acc cache (a[0]) > - > #pragma acc cache (a[0], a[1], b[0:N]) > - > +#pragma acc cache (a[i - 5]) > +#pragma acc cache (a[i + 5:len]) > +#pragma acc cache (a[i + 5:len - 1]) > +#pragma acc cache (b[i]) > +#pragma acc cache (b[i:len]) > +#pragma acc cache (a[ii]) > +#pragma acc cache (a[ii:len]) > +#pragma acc cache (b[ii - 1]) > +#pragma acc cache (b[ii - 1:len]) > +#pragma acc cache (b[i - ii + 1]) > +#pragma acc cache (b[i + ii - 1:len]) > +#pragma acc cache (b[i * ii - 1:len + 1]) > +#pragma acc cache (a[idx + 2]) > +#pragma acc cache (a[idx:len + 2]) > #pragma acc cache (a[idx]) > - > #pragma acc cache (a[idx:len]) > +#pragma acc cache (a[idx + 2:len]) > +#pragma acc cache (a[idx + 2 + i:len]) > +#pragma acc cache (a[idx + 2 + i + ii:len]) > > b[ii] = a[ii]; > } > diff --git gcc/testsuite/c-c++-common/goacc/cache-1.c > gcc/testsuite/c-c++-common/goacc/cache-2.c > similarity index 83% > copy from gcc/testsuite/c-c++-common/goacc/cache-1.c > copy to gcc/testsuite/c-c++-common/goacc/cache-2.c > index 9503341..f717515 100644 > --- gcc/testsuite/c-c++-common/goacc/cache-1.c > +++ gcc/testsuite/c-c++-common/goacc/cache-2.c > @@ -1,3 +1,5 @@ > +/* OpenACC cache directive: invalid usage. */ > + > int > main (int argc, char **argv) > { > @@ -22,57 +24,24 @@ main (int argc, char **argv) > const int len = n; > > #pragma acc cache /* { dg-error "expected '\\\(' before end of line" } */ > - > #pragma acc cache a[0:N] /* { dg-error "expected '\\\(' before 'a'" } */ > - /* { dg-bogus "expected end of line before 'a'" "" { xfail c++ } 26 } */ > - > + /* { dg-bogus "expected end of line before 'a'" "" { xfail c++ } 27 } */ > #pragma acc cache (a) /* { dg-error "expected '\\\['" } */ > - > #pragma acc cache ( /* { dg-error "expected (identifier|unqualified-id) > before end of line" } */ > - > #pragma acc cache () /* { dg-error "expected (identifier|unqualified-id) > before '\\\)' token" } */ > - > #pragma acc cache (,) /* { dg-error "expected (identifier|unqualified-id) > before '(,|\\\))' token" } */ > - > #pragma acc cache (a[0:N] /* { dg-error "expected '\\\)' before end of line" > } */ > - > #pragma acc cache (a[0:N],) /* { dg-error "expected > (identifier|unqualified-id) before '(,|\\\))' token" "" { xfail c } } */ > - > #pragma acc cache (a[0:N]) copyin (a[0:N]) /* { dg-error "expected end of > line before 'copyin'" } */ > - > #pragma acc cache () /* { dg-error "expected (identifier|unqualified-id) > before '\\\)' token" } */ > - > #pragma acc cache (a[0:N] b[0:N]) /* { dg-error "expected '\\\)' before 'b'" > } */ > - > #pragma acc cache (a[0:N] b[0:N}) /* { dg-error "expected '\\\)' before 'b'" > } */ > - /* { dg-bogus "expected end of line before '\\\}' token" "" { xfail c++ > } 47 } */ > - > + /* { dg-bogus "expected end of line before '\\\}' token" "" { xfail c++ > } 38 } */ > #pragma acc cache (a[0:N] /* { dg-error "expected '\\\)' before end of line" > } */ > - > -#pragma acc cache (a[ii]) /* { dg-error "'ii' is not a constant" } */ > - > -#pragma acc cache (a[idx:n]) /* { dg-error "'n' is not a constant" } */ > - > #pragma acc cache (a[0:N]) ( /* { dg-error "expected end of line before > '\\(' token" } */ > - > #pragma acc cache (a[0:N]) ii /* { dg-error "expected end of line before > 'ii'" } */ > - > #pragma acc cache (a[0:N] ii) /* { dg-error "expected '\\)' before 'ii'" } */ > > -#pragma acc cache (a[0:N]) > - > -#pragma acc cache (a[0:N], a[0:N]) > - > -#pragma acc cache (a[0:N], b[0:N]) > - > -#pragma acc cache (a[0]) > - > -#pragma acc cache (a[0], a[1], b[0:N]) > - > -#pragma acc cache (a[idx]) > - > -#pragma acc cache (a[idx:len]) > - > b[ii] = a[ii]; > } > } > diff --git gcc/testsuite/gfortran.dg/goacc/cache-1.f95 > gcc/testsuite/gfortran.dg/goacc/cache-1.f95 > index 2aa9e05..39fbf2c 100644 > --- gcc/testsuite/gfortran.dg/goacc/cache-1.f95 > +++ gcc/testsuite/gfortran.dg/goacc/cache-1.f95 > @@ -1,4 +1,6 @@ > -! { dg-do compile } > +! OpenACC cache directive: valid usage. > +! For execution testing, this file is "#include"d from > +! libgomp/testsuite/libgomp.oacc-fortran/cache-1.f95. > ! { dg-additional-options "-std=f2008" } > > program test > @@ -6,11 +8,8 @@ program test > integer :: i, d(10), e(5,13) > > do concurrent (i=1:5) > - !$acc cache (d) > !$acc cache (d(1:3)) > !$acc cache (d(i:i+2)) > - > - !$acc cache (e) > !$acc cache (e(1:3,2:4)) > !$acc cache (e(i:i+2,i+1:i+3)) > enddo > diff --git gcc/testsuite/gfortran.dg/goacc/cache-2.f95 > gcc/testsuite/gfortran.dg/goacc/cache-2.f95 > new file mode 100644 > index 0000000..be81878 > --- /dev/null > +++ gcc/testsuite/gfortran.dg/goacc/cache-2.f95 > @@ -0,0 +1,12 @@ > +! OpenACC cache directive: invalid usage. > +! { dg-additional-options "-std=f2008" } > + > +program test > + implicit none > + integer :: i, d(10), e(5,13) > + > + do concurrent (i=1:5) > + !$acc cache (d) ! { dg-error "" "TODO" { xfail *-*-* } } > + !$acc cache (e) ! { dg-error "" "TODO" { xfail *-*-* } } > + enddo > +end > diff --git gcc/testsuite/gfortran.dg/goacc/coarray.f95 > gcc/testsuite/gfortran.dg/goacc/coarray.f95 > index 932e1f7..f30917b8 100644 > --- gcc/testsuite/gfortran.dg/goacc/coarray.f95 > +++ gcc/testsuite/gfortran.dg/goacc/coarray.f95 > @@ -24,7 +24,7 @@ contains > !$acc end parallel loop > !$acc parallel loop > do i = 1,5 > - !$acc cache (a) > + !$acc cache (a) ! { dg-error "" "TODO" { xfail *-*-* } } > enddo > !$acc end parallel loop > !$acc update device (a) > diff --git gcc/testsuite/gfortran.dg/goacc/cray.f95 > gcc/testsuite/gfortran.dg/goacc/cray.f95 > index a35ab0d..705c18c 100644 > --- gcc/testsuite/gfortran.dg/goacc/cray.f95 > +++ gcc/testsuite/gfortran.dg/goacc/cray.f95 > @@ -44,7 +44,8 @@ contains > !$acc end parallel loop > !$acc parallel loop > do i = 1,5 > - !$acc cache (ptr) ! TODO: This must fail, as in openacc-1_0-branch > + !TODO: This must fail, as in openacc-1_0-branch. > + !$acc cache (ptr) ! { dg-error "" "TODO" { xfail *-*-* } } > enddo > !$acc end parallel loop > !$acc update device (ptr) > diff --git gcc/testsuite/gfortran.dg/goacc/loop-1.f95 > gcc/testsuite/gfortran.dg/goacc/loop-1.f95 > index b5f9e03..a605f03 100644 > --- gcc/testsuite/gfortran.dg/goacc/loop-1.f95 > +++ gcc/testsuite/gfortran.dg/goacc/loop-1.f95 > @@ -158,15 +158,16 @@ subroutine test1 > enddo > > > - !$acc cache (a) ! { dg-error "inside of loop" } > + !$acc cache (a(1:10)) ! { dg-error "ACC CACHE directive must be inside of > loop" } > > do i = 1,10 > - !$acc cache(a) > + !$acc cache(a(i:i+1)) > enddo > > do i = 1,10 > + !$acc cache(a(i:i+1)) > a(i) = i > - !$acc cache(a) > + !$acc cache(a(i+2:i+2+1)) > enddo > > end subroutine test1 > diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/cache-1.c > libgomp/testsuite/libgomp.oacc-c-c++-common/cache-1.c > index 3f1f0bb..16aaed5 100644 > --- libgomp/testsuite/libgomp.oacc-c-c++-common/cache-1.c > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/cache-1.c > @@ -1,48 +1,3 @@ > -int > -main (int argc, char **argv) > -{ > -#define N 2 > - int a[N], b[N]; > - int i; > +/* OpenACC cache directive. */ > > - for (i = 0; i < N; i++) > - { > - a[i] = 3; > - b[i] = 0; > - } > - > -#pragma acc parallel copyin (a[0:N]) copyout (b[0:N]) > -{ > - int ii; > - > - for (ii = 0; ii < N; ii++) > - { > - const int idx = ii; > - int n = 1; > - const int len = n; > - > -#pragma acc cache (a[0:N]) > - > -#pragma acc cache (a[0:N], b[0:N]) > - > -#pragma acc cache (a[0]) > - > -#pragma acc cache (a[0], a[1], b[0:N]) > - > -#pragma acc cache (a[idx]) > - > -#pragma acc cache (a[idx:len]) > - > - b[ii] = a[ii]; > - } > -} > - > - > - for (i = 0; i < N; i++) > - { > - if (a[i] != b[i]) > - __builtin_abort (); > - } > - > - return 0; > -} > +#include "../../../gcc/testsuite/c-c++-common/goacc/cache-1.c" > diff --git libgomp/testsuite/libgomp.oacc-fortran/cache-1.f95 > libgomp/testsuite/libgomp.oacc-fortran/cache-1.f95 > new file mode 100644 > index 0000000..a68c970 > --- /dev/null > +++ libgomp/testsuite/libgomp.oacc-fortran/cache-1.f95 > @@ -0,0 +1,5 @@ > +! OpenACC cache directive. > +! { dg-additional-options "-std=f2008" } > +! { dg-additional-options "-cpp" } > + > +#include "../../../gcc/testsuite/gfortran.dg/goacc/cache-1.f95" Grüße Thomas