diff mbox

[gomp4] Implement OpenACC 2.5 reference counting, and finalize clause

Message ID 87fug31yyc.fsf@hertz.schwinge.homeip.net
State New
Headers show

Commit Message

Thomas Schwinge May 17, 2017, 11:53 a.m. UTC 
Hi!

On Tue, 16 May 2017 20:55:46 +0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> finalize clause of the exit data directive

This would run into ICEs in the C++ front end (template handling) as well
as C and Fortran front ends (nested function handling), and didn't
pretty-print the "finalize" clause.  Also test cases.  Committed to
gomp-4_0-branch in r248148:

commit 2d734ec8526f73e69c7bfa9b60ec5e9c5a9e4f13
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Wed May 17 11:52:15 2017 +0000

    Complete compiler-side handling of the OpenACC finalize clause
    
            gcc/cp/
            * pt.c (tsubst_omp_clauses): Handle "OMP_CLAUSE_FINALIZE".
            gcc/
            * tree-nested.c (convert_nonlocal_omp_clauses)
            (convert_local_omp_clauses): Handle "OMP_CLAUSE_FINALIZE".
            * tree-pretty-print.c (dump_omp_clause): Handle
            "OMP_CLAUSE_FINALIZE".
            gcc/testsuite/
            * c-c++-common/goacc/data-2.c: Update.
            * g++.dg/goacc/template.C: Likewise.
            * gcc.dg/goacc/nested-function-1.c: Likewise.
            * gfortran.dg/goacc/enter-exit-data.f95: Likewise.
            * gfortran.dg/goacc/nested-function-1.f90: Likewise.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@248148 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp                                    | 5 +++++
 gcc/cp/ChangeLog.gomp                                 | 2 ++
 gcc/cp/pt.c                                           | 1 +
 gcc/testsuite/ChangeLog.gomp                          | 8 ++++++++
 gcc/testsuite/c-c++-common/goacc/data-2.c             | 3 +++
 gcc/testsuite/g++.dg/goacc/template.C                 | 9 +++++++++
 gcc/testsuite/gcc.dg/goacc/nested-function-1.c        | 4 ++++
 gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95   | 3 +++
 gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 | 4 ++++
 gcc/tree-nested.c                                     | 2 ++
 gcc/tree-pretty-print.c                               | 3 +++
 11 files changed, 44 insertions(+)



Grüße
 Thomas
diff mbox

Patch

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index e858e78..d89897d 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,10 @@ 
 2017-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* tree-nested.c (convert_nonlocal_omp_clauses)
+	(convert_local_omp_clauses): Handle "OMP_CLAUSE_FINALIZE".
+	* tree-pretty-print.c (dump_omp_clause): Handle
+	"OMP_CLAUSE_FINALIZE".
+
 	* gimplify.c (gimplify_oacc_declare_1) <GOMP_MAP_ALLOC>: Use
 	"GOMP_MAP_RELEASE".
 
diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp
index 9a68194..b0c3dbf 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,5 +1,7 @@ 
 2017-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* pt.c (tsubst_omp_clauses): Handle "OMP_CLAUSE_FINALIZE".
+
 	* parser.c (cp_parser_oacc_data_clause)
 	<PRAGMA_OACC_CLAUSE_DELETE>: Use "GOMP_MAP_RELEASE".
 
diff --git gcc/cp/pt.c gcc/cp/pt.c
index 84f64d8..abe8d36 100644
--- gcc/cp/pt.c
+++ gcc/cp/pt.c
@@ -14751,6 +14751,7 @@  tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
 	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	case OMP_CLAUSE_DEVICE_TYPE:
 	  break;
 	case OMP_CLAUSE_BIND:
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 34f0a06..960ad15 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,11 @@ 
+2017-05-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/data-2.c: Update.
+	* g++.dg/goacc/template.C: Likewise.
+	* gcc.dg/goacc/nested-function-1.c: Likewise.
+	* gfortran.dg/goacc/enter-exit-data.f95: Likewise.
+	* gfortran.dg/goacc/nested-function-1.f90: Likewise.
+
 2017-05-15  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* c-c++-common/cpp/openacc-define-3.c: Update.
diff --git gcc/testsuite/c-c++-common/goacc/data-2.c gcc/testsuite/c-c++-common/goacc/data-2.c
index 1043bf8a..8c5d42a 100644
--- gcc/testsuite/c-c++-common/goacc/data-2.c
+++ gcc/testsuite/c-c++-common/goacc/data-2.c
@@ -10,6 +10,9 @@  foo (void)
 #pragma acc exit data delete (a) if (0)
 #pragma acc exit data copyout (b) if (a)
 #pragma acc exit data delete (b)
+#pragma acc exit data delete (a) if (!0) finalize
+#pragma acc exit data copyout (b) finalize if (!a)
+#pragma acc exit data finalize delete (b)
 #pragma acc enter /* { dg-error "expected 'data' after" } */
 #pragma acc exit /* { dg-error "expected 'data' after" } */
 #pragma acc enter data /* { dg-error "has no data movement clause" } */
diff --git gcc/testsuite/g++.dg/goacc/template.C gcc/testsuite/g++.dg/goacc/template.C
index f4d255c..d1acece 100644
--- gcc/testsuite/g++.dg/goacc/template.C
+++ gcc/testsuite/g++.dg/goacc/template.C
@@ -86,6 +86,8 @@  oacc_parallel_copy (T a)
 #pragma acc update self (b)
 #pragma acc update device (b)
 #pragma acc exit data delete (b)
+#pragma acc exit data finalize copyout (b)
+#pragma acc exit data delete (b) finalize
 
   return b;
 }
@@ -133,6 +135,13 @@  oacc_kernels_copy (T a)
     b = a;
   }
 
+#pragma acc update host (b)
+#pragma acc update self (b)
+#pragma acc update device (b)
+#pragma acc exit data delete (b)
+#pragma acc exit data finalize copyout (b)
+#pragma acc exit data delete (b) finalize
+
   return b;
 }
 
diff --git gcc/testsuite/gcc.dg/goacc/nested-function-1.c gcc/testsuite/gcc.dg/goacc/nested-function-1.c
index 5fc2e46..6b76112 100644
--- gcc/testsuite/gcc.dg/goacc/nested-function-1.c
+++ gcc/testsuite/gcc.dg/goacc/nested-function-1.c
@@ -56,6 +56,8 @@  int main ()
 	for (local_j = 0; local_j < N; ++local_j)
 	  ;
       }
+
+#pragma acc exit data copyout(local_a) delete(local_i) finalize
   }
 
   void nonlocal ()
@@ -95,6 +97,8 @@  int main ()
 	for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
 	  ;
       }
+
+#pragma acc exit data copyout(nonlocal_a) delete(nonlocal_i) finalize
   }
 
   local ();
diff --git gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
index 8f1715e..805459c 100644
--- gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
+++ gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
@@ -84,5 +84,8 @@  contains
   !$acc exit data delete (tip) ! { dg-error "POINTER" }
   !$acc exit data delete (tia) ! { dg-error "ALLOCATABLE" }
   !$acc exit data copyout (i) delete (i) ! { dg-error "multiple clauses" }
+  !$acc exit data finalize
+  !$acc exit data finalize copyout (i)
+  !$acc exit data finalize delete (i)
   end subroutine foo
 end module test
diff --git gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90
index bbb53c3..005193f 100644
--- gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90
+++ gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90
@@ -56,6 +56,8 @@  contains
        enddo
     enddo
     !$acc end kernels loop
+
+    !$acc exit data copyout(local_a) delete(local_i) finalize
   end subroutine local
 
   subroutine nonlocal ()
@@ -93,5 +95,7 @@  contains
        enddo
     enddo
     !$acc end kernels loop
+
+    !$acc exit data copyout(nonlocal_a) delete(nonlocal_i) finalize
   end subroutine nonlocal
 end program main
diff --git gcc/tree-nested.c gcc/tree-nested.c
index 3ddfd65..d6635ab 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1203,6 +1203,7 @@  convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  break;
 
 	case OMP_CLAUSE_DEVICE_TYPE:
@@ -1902,6 +1903,7 @@  convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_IF_PRESENT:
+	case OMP_CLAUSE_FINALIZE:
 	  break;
 
 	case OMP_CLAUSE_DEVICE_TYPE:
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index a208e8f..a8a0073 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -1093,6 +1093,9 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
     case OMP_CLAUSE_IF_PRESENT:
       pp_string (pp, "if_present");
       break;
+    case OMP_CLAUSE_FINALIZE:
+      pp_string (pp, "finalize");
+      break;
 
     default:
       pp_string (pp, "unknown");