Hi! Julian, Tobias, regarding the following OpenACC 'exit data' 'finalize' handling:
On 2018-05-25T13:01:58-0700, Cesar Philippidis <ce...@codesourcery.com> wrote:
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -10859,6 +10849,53 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq
> *pre_p)
> + else if (TREE_CODE (expr) == OACC_EXIT_DATA
> + && omp_find_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;
> + }
> + }
> --- a/libgomp/oacc-parallel.c
> +++ b/libgomp/oacc-parallel.c
> @@ -286,6 +360,17 @@ GOACC_enter_exit_data (int device, size_t mapnum,
> + /* 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;
> + }
> +
> @@ -360,22 +458,28 @@ GOACC_enter_exit_data (int device, size_t mapnum,
> switch (kind)
> {
> - case GOMP_MAP_POINTER:
> - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
> - == GOMP_MAP_FORCE_FROM,
> - async, 1);
> - break;
> + case GOMP_MAP_RELEASE:
> case GOMP_MAP_DELETE:
> - acc_delete (hostaddrs[i], sizes[i]);
> + if (acc_is_present (hostaddrs[i], sizes[i]))
> + {
> + if (finalize)
> + acc_delete_finalize (hostaddrs[i], sizes[i]);
> + else
> + acc_delete (hostaddrs[i], sizes[i]);
> + }
> break;
> + case GOMP_MAP_FROM:
> case GOMP_MAP_FORCE_FROM:
> - acc_copyout (hostaddrs[i], sizes[i]);
> + if (finalize)
> + acc_copyout_finalize (hostaddrs[i], sizes[i]);
> + else
> + acc_copyout (hostaddrs[i], sizes[i]);
> break;
> default:
> gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
> @@ -385,10 +489,12 @@ GOACC_enter_exit_data (int device, size_t mapnum,
> }
> else
> {
> - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
> - == GOMP_MAP_FORCE_FROM, async, 3);
> +[...]
> + gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async,
> + finalize, pointer);
... does the attached patch "[OpenACC] Elaborate/simplify 'exit data'
'finalize' handling" (with "No functional changes") match your
understanding of what's going on? If approving this patch, please
respond with "Reviewed-by: NAME <EMAIL>" so that your effort will be
recorded in the commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>.
(It will be a separate discussion to change the 'GOMP_MAP_POINTER',
'GOMP_MAP_TO_PSET' stuff later on -- thinking about the changes from
Julian's big "OpenACC reference count overhaul" as well as
<https://gcc.gnu.org/PR92929> "OpenACC/OpenMP 'target' 'exit
data'/'update' optimizations". That patch here is just meant to document
what's going at present, and simplify things as a preparation for other
changes.)
Grüße
Thomas
From 283577c63b374c3e368e3c0b68b90e19085f193c Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Fri, 13 Dec 2019 13:56:51 +0100
Subject: [PATCH] [OpenACC] Elaborate/simplify 'exit data' 'finalize' handling
No functional changes.
---
gcc/gimplify.c | 23 +++++++++++--------
gcc/testsuite/c-c++-common/goacc/finalize-1.c | 11 ++++++++-
gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 10 ++++++++
libgomp/oacc-mem.c | 14 +++--------
4 files changed, 36 insertions(+), 22 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 9073680cb31..60a80cb8098 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -12738,27 +12738,30 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
&& omp_find_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;
+ /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote "finalize"
+ semantics. */
+ bool have_clause = 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;
+ have_clause = true;
break;
case GOMP_MAP_RELEASE:
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
- finalize_marked = true;
+ have_clause = 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);
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_TO_PSET:
+ /* TODO PR92929: we may see these here, but they'll always follow
+ one of the clauses above, and will be handled by libgomp as
+ one group, so no handling required here. */
+ gcc_assert (have_clause);
break;
+ default:
+ gcc_unreachable ();
}
}
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
index 94820290b94..3d64b2e7cb3 100644
--- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
@@ -4,8 +4,10 @@
extern int del_r;
extern float del_f[3];
+extern char *del_f_p;
extern double cpo_r[8];
extern long cpo_f;
+extern char *cpo_f_p;
void f ()
{
@@ -17,6 +19,10 @@ void 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 finalize delete (del_f_p[2:5])
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) 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" } } */
@@ -24,5 +30,8 @@ void f ()
#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" } } */
-}
+#pragma acc exit data copyout (cpo_f_p[4:10]) finalize
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index 5c7a921a2e3..ca642156e9f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -6,8 +6,10 @@
IMPLICIT NONE
INTEGER :: del_r
REAL, DIMENSION (3) :: del_f
+ INTEGER (1), DIMENSION (:), ALLOCATABLE :: del_f_p
DOUBLE PRECISION, DIMENSION (8) :: cpo_r
LOGICAL :: cpo_f
+ INTEGER (1), DIMENSION (:), ALLOCATABLE :: cpo_f_p
!$ACC EXIT DATA DELETE (del_r)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
@@ -17,6 +19,10 @@
! { 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 FINALIZE DELETE (del_f_p(2:5))
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) 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" } }
@@ -24,4 +30,8 @@
!$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" } }
+
+!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
END SUBROUTINE f
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 311f9585f77..291ef9192b9 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1058,17 +1058,6 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
thr = goacc_thread ();
acc_dev = thr->dev;
- /* 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)
{
@@ -1221,6 +1210,9 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
{
unsigned char kind = kinds[i] & 0xff;
+ bool finalize = (kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_FORCE_FROM);
+
int pointer = find_pointer (i, mapnum, kinds);
if (!pointer)
--
2.17.1
signature.asc
Description: PGP signature
