This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.

Index Nav:	[Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav:	[Date Prev] [Date Next]	[Thread Prev] [Thread Next]
Other format:	[Raw text]

[gomp4] OpenACC: Complete changes to disallow the independent clause after device_type (was: Remove clause from device_type options)

From: Thomas Schwinge <thomas at codesourcery dot com>
To: James Norris <jnorris at codesourcery dot com>, <gcc-patches at gcc dot gnu dot org>
Date: Tue, 7 Jul 2015 14:58:20 +0200
Subject: [gomp4] OpenACC: Complete changes to disallow the independent clause after device_type (was: Remove clause from device_type options)
Authentication-results: sourceware.org; auth=none
References: <5593E27D dot 1080205 at codesourcery dot com>

Hi!

On Wed, 1 Jul 2015 07:52:13 -0500, James Norris <jnorris@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  <thomas@codesourcery.com>
+
+	* parser.c (OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK): Remove
+	PRAGMA_OACC_CLAUSE_INDEPENDENT.
+
 2015-06-15  Cesar Philippidis  <cesar@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  <thomas@codesourcery.com>
+
+	* openmp.c (OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK): Remove
+	OMP_CLAUSE_INDEPENDENT.
+
 2015-06-18  James Norris  <jnorris@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  <thomas@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  <tom@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

Attachment: signature.asc
Description: PGP signature

Follow-Ups:
- [gomp4] OpenACC device_type clause (was: OpenACC: Complete changes to disallow the independent clause after device_type)
  - From: Thomas Schwinge

References:
- [gomp4] Remove clause from device_type options
  - From: James Norris

Index Nav:	[Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav:	[Date Prev] [Date Next]	[Thread Prev] [Thread Next]