From 32128577aed11aa21f3227edc2276da677e97636 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 18 Dec 2019 18:00:51 +0100 Subject: [PATCH] [OpenACC] Elaborate/simplify 'exit data' 'finalize' handling No functional changes. gcc/ * gimplify.c (gimplify_omp_target_update): Elaborate 'exit data' 'finalize' handling. gcc/testsuite/ * c-c++-common/goacc/finalize-1.c: Extend. * gfortran.dg/goacc/finalize-1.f: Likewise. libgomp/ * oacc-mem.c (GOACC_enter_exit_data): Simplify 'exit data' 'finalize' handling. From-SVN: r279531 --- gcc/ChangeLog | 5 ++++ gcc/gimplify.c | 23 +++++++++++-------- gcc/testsuite/ChangeLog | 5 ++++ gcc/testsuite/c-c++-common/goacc/finalize-1.c | 11 ++++++++- gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 10 ++++++++ libgomp/ChangeLog | 3 +++ libgomp/oacc-mem.c | 14 +++-------- 7 files changed, 49 insertions(+), 22 deletions(-) diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 642faea1e44..be8dfa3fecf 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2019-12-18 Thomas Schwinge + + * gimplify.c (gimplify_omp_target_update): Elaborate 'exit data' + 'finalize' handling. + 2019-12-18 Tobias Burnus PR middle-end/86416 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/ChangeLog b/gcc/testsuite/ChangeLog index 942448061cb..f1bf3452243 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,8 @@ +2019-12-18 Thomas Schwinge + + * c-c++-common/goacc/finalize-1.c: Extend. + * gfortran.dg/goacc/finalize-1.f: Likewise. + 2019-12-18 Harald Anlauf PR fortran/70853 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/ChangeLog b/libgomp/ChangeLog index d9aba5bee18..c4283fdfe1d 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,8 @@ 2019-12-18 Thomas Schwinge + * oacc-mem.c (GOACC_enter_exit_data): Simplify 'exit data' + 'finalize' handling. + PR libgomp/92848 * oacc-mem.c (acc_map_data, present_create_copy) (goacc_insert_pointer): Use 'GOMP_MAP_VARS_ENTER_DATA'. diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 54427982341..b21d83c37d8 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1061,17 +1061,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) { @@ -1224,6 +1213,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.30.2