Hi! 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 > 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 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) : Loosen checking. gcc/cp/ PR c/71381 * parser.c (cp_parser_omp_var_list_no_open) : 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