Hi! On Wed, 1 Jul 2015 07:52:13 -0500, James Norris <jnor...@codesourcery.com> wrote: > The independent clause is not available for use > with device_type clauses associated with loop > directives. This patch removes the usage.
Thanks for noticing this!
> Committed to gomp-4_0-branch
> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -13069,7 +13069,6 @@ c_parser_oacc_host_data (location_t loc, c_parser
> *parser)
> | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \
> | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \
> | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \
> - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \
> | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \
> | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE) )
In r225514, I completed this change as follows:
commit c4e50167b783540d3cf13685380808b73009dc74
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue Jul 7 12:56:37 2015 +0000
OpenACC: Complete changes to disallow the independent clause after
device_type
In r225247, this was done for the C front end, but it likewise applies to
C++
and Fortran. Also, the test cases need to be adjusted to avoid regressions
(as
has happened for C):
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c (test for excess errors)
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop auto private\\(i4\\)" 2
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop collapse\\(1\\) worker private\\(i2\\)" 2
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop private\\(i5\\)" 2
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop seq private\\(i1\\.2\\) private\\(i1\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop seq private\\(i2\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop seq private\\(i4\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop seq private\\(i5\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop seq private\\(i6\\)" 3
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop tile\\(1\\) gang private\\(i1\\.0\\) private\\(i1\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop tile\\(1\\) gang private\\(i1\\.1\\) private\\(i1\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc loop vector private\\(i3\\)" 2
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "acc_parallel wait\\(1\\) vector_length\\(1\\) num_workers\\(1\\)
num_gangs\\(1\\) async\\(1\\) wait\\(3\\) vector_length\\(128\\)
num_workers\\(300\\) num_gangs\\(300\\) async\\(3" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "oacc_kernels async\\(-1\\) wait\\(0\\) async\\(0\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "oacc_kernels async\\(-1\\) wait\\(2\\) async\\(2\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "oacc_kernels async\\(-1\\)" 4
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "oacc_parallel wait\\(1\\) vector_length\\(1\\) num_workers\\(1\\)
num_gangs\\(1\\) async\\(1\\) wait\\(2\\) vector_length\\(64\\)
num_workers\\(200\\) num_gangs\\(200\\) async\\(2\\)" 1
[-PASS:-]{+FAIL:+} c-c++-common/goacc/dtype-1.c scan-tree-dump-times
omplower "oacc_parallel wait\\(1\\) vector_length\\(32\\) num_workers\\(100\\)
num_gangs\\(100\\) async\\(1\\)" 1
gcc/cp/
* parser.c (OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK): Remove
PRAGMA_OACC_CLAUSE_INDEPENDENT.
gcc/fortran/
* openmp.c (OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK): Remove
OMP_CLAUSE_INDEPENDENT.
gcc/testsuite/
* c-c++-common/goacc/dtype-1.c: Don't specify the OpenACC
independent clause after device_type.
* gfortran.dg/goacc/dtype-1.f95: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@225514
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/cp/ChangeLog.gomp | 5 +++++
gcc/cp/parser.c | 1 -
gcc/fortran/ChangeLog.gomp | 5 +++++
gcc/fortran/openmp.c | 3 +--
gcc/testsuite/ChangeLog.gomp | 6 ++++++
gcc/testsuite/c-c++-common/goacc/dtype-1.c | 10 +++++-----
gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 | 6 +++---
7 files changed, 25 insertions(+), 11 deletions(-)
diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp
index 85c3d87..eb39c40 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-07-07 Thomas Schwinge <tho...@codesourcery.com>
+
+ * parser.c (OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK): Remove
+ PRAGMA_OACC_CLAUSE_INDEPENDENT.
+
2015-06-15 Cesar Philippidis <ce...@codesourcery.com>
* parser.c (cp_parser_oacc_all_clauses): Call
diff --git gcc/cp/parser.c gcc/cp/parser.c
index c233595..2c60136 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -32370,7 +32370,6 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser,
cp_token *pragma_tok,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE) )
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index aa02ffe..2d8d1e7 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-07-07 Thomas Schwinge <tho...@codesourcery.com>
+
+ * openmp.c (OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK): Remove
+ OMP_CLAUSE_INDEPENDENT.
+
2015-06-18 James Norris <jnor...@codesourcery.com>
* trans-decl.c (find_module_oacc_declare_clauses): Fix setting of
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index b98a933..cd9565c 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -1385,8 +1385,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t
mask,
#define OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK \
(OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
- | OMP_CLAUSE_VECTOR | OMP_CLAUSE_AUTO | OMP_CLAUSE_INDEPENDENT \
- | OMP_CLAUSE_SEQ | OMP_CLAUSE_TILE)
+ | OMP_CLAUSE_VECTOR | OMP_CLAUSE_AUTO | OMP_CLAUSE_SEQ | OMP_CLAUSE_TILE)
#define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \
(OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT)
#define OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK \
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index ee35ded..066b862 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2015-07-07 Thomas Schwinge <tho...@codesourcery.com>
+
+ * c-c++-common/goacc/dtype-1.c: Don't specify the OpenACC
+ independent clause after device_type.
+ * gfortran.dg/goacc/dtype-1.f95: Likewise.
+
2015-07-01 Tom de Vries <t...@codesourcery.com>
PR tree-optimization/66716
diff --git gcc/testsuite/c-c++-common/goacc/dtype-1.c
gcc/testsuite/c-c++-common/goacc/dtype-1.c
index 1925c1e..57906c7 100644
--- gcc/testsuite/c-c++-common/goacc/dtype-1.c
+++ gcc/testsuite/c-c++-common/goacc/dtype-1.c
@@ -53,7 +53,7 @@ test ()
for (i3 = 1; i3 < 10; i3++)
#pragma acc loop dtype (nVidia) auto
for (i4 = 1; i4 < 10; i4++)
-#pragma acc loop dtype (nVidia) independent
+#pragma acc loop dtype (nVidia)
for (i5 = 1; i5 < 10; i5++)
#pragma acc loop device_type (nVidia) seq
for (i6 = 1; i6 < 10; i6++)
@@ -69,7 +69,7 @@ test ()
for (i3 = 1; i3 < 10; i3++)
#pragma acc loop dtype (nVidia) auto device_type (*) seq
for (i4 = 1; i4 < 10; i4++)
-#pragma acc loop device_type (nVidia) independent device_type (*) seq
+#pragma acc loop device_type (nVidia) device_type (*) seq
for (i5 = 1; i5 < 10; i5++)
#pragma acc loop device_type (nVidia) seq
for (i6 = 1; i6 < 10; i6++)
@@ -85,7 +85,7 @@ test ()
for (i3 = 1; i3 < 10; i3++)
#pragma acc loop device_type (nVidiaGPU) auto device_type (*) seq
for (i4 = 1; i4 < 10; i4++)
-#pragma acc loop dtype (nVidiaGPU) independent dtype (*) seq
+#pragma acc loop dtype (nVidiaGPU) dtype (*) seq
for (i5 = 1; i5 < 10; i5++)
#pragma acc loop device_type (nVidiaGPU) seq device_type (*) seq
for (i6 = 1; i6 < 10; i6++)
@@ -147,7 +147,7 @@ test ()
/* { dg-final { scan-tree-dump-times "acc loop auto private\\(i4\\)" 2
"omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop private\\(i5\\)" 2 "omplower" }
} */
+/* { dg-final { scan-tree-dump-times "acc loop private\\(i5\\)" 1 "omplower" }
} */
/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i6\\)" 3
"omplower" } } */
@@ -155,6 +155,6 @@ test ()
/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i4\\)" 1
"omplower" } } */
-/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i5\\)" 1
"omplower" } } */
+/* { dg-final { scan-tree-dump-times "acc loop seq private\\(i5\\)" 2
"omplower" } } */
/* { dg-final { cleanup-tree-dump "omplower" } } */
diff --git gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
index 15760c3..6b89fe8 100644
--- gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
+++ gcc/testsuite/gfortran.dg/goacc/dtype-1.f95
@@ -54,7 +54,7 @@ program dtype
do i3 = 1, 10
!$acc loop device_type (nVidia) auto
do i4 = 1, 10
- !$acc loop dtype (nVidia) independent
+ !$acc loop dtype (nVidia)
do i5 = 1, 10
!$acc loop dtype (nVidia) seq
do i6 = 1, 10
@@ -76,7 +76,7 @@ program dtype
do i3 = 1, 10
!$acc loop device_type (nVidia) auto dtype (*) seq
do i4 = 1, 10
- !$acc loop dtype (nVidia) independent &
+ !$acc loop dtype (nVidia) &
!$acc& dtype (*) seq
do i5 = 1, 10
!$acc loop device_type (nVidia) seq
@@ -99,7 +99,7 @@ program dtype
do i3 = 1, 10
!$acc loop dtype (nVidiaGPU) auto device_type (*) seq
do i4 = 1, 10
- !$acc loop dtype (nVidiaGPU) independent &
+ !$acc loop dtype (nVidiaGPU) &
!$acc& dtype (*) seq
do i5 = 1, 10
!$acc loop dtype (nVidiaGPU) seq device_type (*) seq
Grüße,
Thomas
signature.asc
Description: PGP signature
