| Message ID | 877fe7sthf.fsf@kepler.schwinge.homeip.net |
|---|---|
| State | New |
| Headers | show |
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 <thomas@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
On Wed, Jun 08, 2016 at 03:28:57PM +0200, Thomas Schwinge wrote: > > [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_. Ok. > > --- 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? */ Please try to not add GNU extensions on top of OpenACC, unless strictly necessary. It is better if the compiler is strict and there is interoperability. If you think it should accept something that it doesn't, talk to the OpenACC committee. Jakub
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"