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 <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
signature.asc
Description: PGP signature
