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
Thanks!
> --- libgomp/oacc-parallel.c (revision 248095)
> +++ libgomp/oacc-parallel.c (revision 248096)
> void
> GOACC_enter_exit_data (int device, size_t mapnum,
> void **hostaddrs, size_t *sizes, unsigned short *kinds,
> - int async, int num_waits, ...)
> + int async, int finalize, int num_waits, ...)
> --- gcc/omp-low.c (revision 248095)
> +++ gcc/omp-low.c (revision 248096)
> @@ -14216,6 +14218,13 @@
> if (t_async)
> args.safe_push (t_async);
>
> + if (start_ix == BUILT_IN_GOACC_ENTER_EXIT_DATA)
> + {
> + c = find_omp_clause (clauses, OMP_CLAUSE_FINALIZE);
> + tree t_finalize = c ? integer_one_node : integer_zero_node;
> + args.safe_push (t_finalize);
> + }
> +
> /* Save the argument index, and ... */
> unsigned t_wait_idx = args.length ();
> unsigned num_waits = 0;
This breaks the ABI. (Also this didn't update
gcc/omp-builtins.def:BUILT_IN_GOACC_ENTER_EXIT_DATA. If I remember
correctly, I noted before that we really should add some consistency
checking for definition vs. usage of builtin functions.)
So I changed that to do similar to what Cesar recently had done for the
OpenACC update construct's if_present clause, and also added test cases.
Committed to gomp-4_0-branch in r248149:
commit 98cb728fdc1d848b3beadae5a81b663a30915925
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed May 17 12:03:00 2017 +0000
Revert GOACC_enter_exit_data ABI change
Instead, use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
semantics apply to all mappings of this OpenACC directive.
gcc/
* omp-low.c (expand_omp_target): Revert GOACC_enter_exit_data ABI
change.
* gimplify.c (gimplify_omp_target_update): Handle
OMP_CLAUSE_FINALIZE for OACC_EXIT_DATA.
gcc/testsuite/
* c-c++-common/goacc/finalize-1.c: New file.
* gfortran.dg/goacc/finalize-1.f: Likewise.
libgomp/
* oacc-parallel.c (GOACC_enter_exit_data): Locally compute
"finalize", and remove the formal parameter. Adjust all users.
(GOACC_declare): Don't replace GOMP_MAP_FROM with
GOMP_MAP_FORCE_FROM.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@248149
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 5 +++++
gcc/gimplify.c | 27 ++++++++++++++++++++++++++
gcc/omp-low.c | 7 -------
gcc/testsuite/ChangeLog.gomp | 3 +++
gcc/testsuite/c-c++-common/goacc/finalize-1.c | 28 +++++++++++++++++++++++++++
gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 27 ++++++++++++++++++++++++++
libgomp/ChangeLog.gomp | 5 +++++
libgomp/libgomp_g.h | 2 +-
libgomp/oacc-parallel.c | 22 +++++++++++++++------
9 files changed, 112 insertions(+), 14 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index d89897d..084ac9b 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,10 @@
2017-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ * omp-low.c (expand_omp_target): Revert GOACC_enter_exit_data ABI
+ change.
+ * gimplify.c (gimplify_omp_target_update): Handle
+ OMP_CLAUSE_FINALIZE for OACC_EXIT_DATA.
+
* tree-nested.c (convert_nonlocal_omp_clauses)
(convert_local_omp_clauses): Handle "OMP_CLAUSE_FINALIZE".
* tree-pretty-print.c (dump_omp_clause): Handle
diff --git gcc/gimplify.c gcc/gimplify.c
index 7812471..54c5d43 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -10060,6 +10060,33 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq
*pre_p)
break;
}
}
+ else if (TREE_CODE (expr) == OACC_EXIT_DATA
+ && find_omp_clause (OMP_STANDALONE_CLAUSES (expr),
+ OMP_CLAUSE_FINALIZE))
+ {
+ /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
+ semantics apply to all mappings of this OpenACC directive. */
+ bool finalize_marked = false;
+ for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM);
+ finalize_marked = true;
+ break;
+ case GOMP_MAP_RELEASE:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
+ finalize_marked = true;
+ break;
+ default:
+ /* Check consistency: libgomp relies on the very first data
+ mapping clause being marked, so make sure we did that before
+ any other mapping clauses. */
+ gcc_assert (finalize_marked);
+ break;
+ }
+ }
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
gimplify_seq_add_stmt (pre_p, stmt);
diff --git gcc/omp-low.c gcc/omp-low.c
index 394dd47..048d9fb 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -14218,13 +14218,6 @@ expand_omp_target (struct omp_region *region)
if (t_async)
args.safe_push (t_async);
- if (start_ix == BUILT_IN_GOACC_ENTER_EXIT_DATA)
- {
- c = find_omp_clause (clauses, OMP_CLAUSE_FINALIZE);
- tree t_finalize = c ? integer_one_node : integer_zero_node;
- args.safe_push (t_finalize);
- }
-
/* Save the argument index, and ... */
unsigned t_wait_idx = args.length ();
unsigned num_waits = 0;
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 960ad15..9ff5476 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,8 @@
2017-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ * c-c++-common/goacc/finalize-1.c: New file.
+ * gfortran.dg/goacc/finalize-1.f: Likewise.
+
* c-c++-common/goacc/data-2.c: Update.
* g++.dg/goacc/template.C: Likewise.
* gcc.dg/goacc/nested-function-1.c: Likewise.
diff --git gcc/testsuite/c-c++-common/goacc/finalize-1.c
gcc/testsuite/c-c++-common/goacc/finalize-1.c
new file mode 100644
index 0000000..9482029
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/finalize-1.c
@@ -0,0 +1,28 @@
+/* Test valid usage and processing of the finalize clause. */
+
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
+extern int del_r;
+extern float del_f[3];
+extern double cpo_r[8];
+extern long cpo_f;
+
+void f ()
+{
+#pragma acc exit data delete (del_r)
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(release:del_r\\);$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple"
} } */
+
+#pragma acc exit data finalize delete (del_f)
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(release:del_f\\) finalize;$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1
"gimple" } } */
+
+#pragma acc exit data copyout (cpo_r)
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(from:cpo_r\\);$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
*/
+
+#pragma acc exit data copyout (cpo_f) finalize
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize
map\\(from:cpo_f\\);$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$"
1 "gimple" } } */
+}
+
diff --git gcc/testsuite/gfortran.dg/goacc/finalize-1.f
gcc/testsuite/gfortran.dg/goacc/finalize-1.f
new file mode 100644
index 0000000..5c7a921
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -0,0 +1,27 @@
+! Test valid usage and processing of the finalize clause.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+ SUBROUTINE f
+ IMPLICIT NONE
+ INTEGER :: del_r
+ REAL, DIMENSION (3) :: del_f
+ DOUBLE PRECISION, DIMENSION (8) :: cpo_r
+ LOGICAL :: cpo_f
+
+!$ACC EXIT DATA DELETE (del_r)
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(release:del_r\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple"
} }
+
+!$ACC EXIT DATA FINALIZE DELETE (del_f)
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(release:del_f\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1
"gimple" } }
+
+!$ACC EXIT DATA COPYOUT (cpo_r)
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(from:cpo_r\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
+
+!$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(from:cpo_f\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$"
1 "gimple" } }
+ END SUBROUTINE f
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 61d0358..2ea7215 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,10 @@
2017-05-17 Thomas Schwinge <tho...@codesourcery.com>
+ * oacc-parallel.c (GOACC_enter_exit_data): Locally compute
+ "finalize", and remove the formal parameter. Adjust all users.
+ (GOACC_declare): Don't replace GOMP_MAP_FROM with
+ GOMP_MAP_FORCE_FROM.
+
* oacc-parallel.c (GOACC_enter_exit_data, GOACC_declare): Handle
"GOMP_MAP_RELEASE".
diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
index 00f7a57..864ef3d 100644
--- libgomp/libgomp_g.h
+++ libgomp/libgomp_g.h
@@ -304,7 +304,7 @@ extern void GOACC_data_start (int, size_t, void **, size_t
*,
unsigned short *);
extern void GOACC_data_end (void);
extern void GOACC_enter_exit_data (int, size_t, void **,
- size_t *, unsigned short *, int, int, int,
...);
+ size_t *, unsigned short *, int, int, ...);
extern void GOACC_update (int, size_t, void **, size_t *,
unsigned short *, int, int, ...);
extern void GOACC_wait (int, int, ...);
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index 5b8435b..ff6e96c 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -601,7 +601,7 @@ GOACC_data_end (void)
void
GOACC_enter_exit_data (int device, size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned short *kinds,
- int async, int finalize, int num_waits, ...)
+ int async, int num_waits, ...)
{
struct goacc_thread *thr;
struct gomp_device_descr *acc_dev;
@@ -609,6 +609,17 @@ GOACC_enter_exit_data (int device, size_t mapnum,
bool data_enter = false;
size_t i;
+ /* Determine whether "finalize" semantics apply to all mappings of this
+ OpenACC directive. */
+ bool finalize = false;
+ if (mapnum > 0)
+ {
+ unsigned char kind = kinds[0] & 0xff;
+ if (kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_FORCE_FROM)
+ finalize = true;
+ }
+
/* Determine if this is an "acc enter data". */
for (i = 0; i < mapnum; ++i)
{
@@ -1102,7 +1113,7 @@ GOACC_declare (int device, size_t mapnum,
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
- &kinds[i], 0, 0, 0);
+ &kinds[i], 0, 0);
break;
case GOMP_MAP_FORCE_DEVICEPTR:
@@ -1111,19 +1122,18 @@ GOACC_declare (int device, size_t mapnum,
case GOMP_MAP_ALLOC:
if (!acc_is_present (hostaddrs[i], sizes[i]))
GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
- &kinds[i], 0, 0, 0);
+ &kinds[i], 0, 0);
break;
case GOMP_MAP_TO:
GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
- &kinds[i], 0, 0, 0);
+ &kinds[i], 0, 0);
break;
case GOMP_MAP_FROM:
- kinds[i] = GOMP_MAP_FORCE_FROM;
GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
- &kinds[i], 0, 0, 0);
+ &kinds[i], 0, 0);
break;
case GOMP_MAP_FORCE_PRESENT:
Grüße
Thomas
