Hi!
On Tue, 16 May 2017 20:55:46 +0800, Chung-Lin Tang <chunglin_t...@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(+)
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 <tho...@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 <tho...@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 <tho...@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 <tho...@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");
Grüße
Thomas
