[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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 642faea1e44f063e448f328cd7784531d7ddea41..be8dfa3fecfe8403ef7385d9c75d8d9bc6e57878 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,8 @@
+2019-12-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+       * gimplify.c (gimplify_omp_target_update): Elaborate 'exit data'
+       'finalize' handling.
+
 2019-12-18  Tobias Burnus  <tobias@codesourcery.com>
 
        PR middle-end/86416

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 9073680cb3179f196979a0ea10b8db0188462eb8..60a80cb8098a3538fa6f7971addc3c9ffb8d0c77 100644 (file)
--- 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 942448061cb4d115793fb44f7145f505f8ce225f..f1bf3452243eeba76e05a3e2a425b7819a2d80d0 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2019-12-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+       * c-c++-common/goacc/finalize-1.c: Extend.
+       * gfortran.dg/goacc/finalize-1.f: Likewise.
+
 2019-12-18  Harald Anlauf  <anlauf@gmx.de>
 
        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 94820290b94335fd5feca6b2e16df73d2522e34b..3d64b2e7cb32f75cc5c0077d35902d1edd8be0c6 100644 (file)
--- 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 5c7a921a2e37f4ae45a46cb9a526609ddf2dc791..ca642156e9fdef6b8e78bdd61d9a8807bb0dcfa0 100644 (file)
--- 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 d9aba5bee18554cd6b2571e52d8f46dd5a2ab586..c4283fdfe1d50ebc1d5402211c7f4f9dafe0f172 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,8 @@
 2019-12-18  Thomas Schwinge  <thomas@codesourcery.com>
 
+       * 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 5442798234142d97dca276c5cc50bf49c04b780c..b21d83c37d84edc8b0dd6955fbfba2b47791381b 100644 (file)
--- 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)