OpenMP/Fortran: Fix unmapping of GOMP_MAP_POINTER for scalar allocatables/pointers

target exit data: Do unmap GOMP_MAP_POINTER for scalar allocatables/pointers
to prevent stale mappings.

While for allocatable/pointer arrays, there is a PSET followed by POINTER,
for allocatable/pointer scalars there is only a POINTER. Before the below
mentioned OG12 patch: For exit data, PSET was converted to RELEASE/DELETE
in gimplify.cc while all POINTER were removed; correct for arrays but leaving
POINTER behind for scalars. Since that commit, all in trans-openmp.cc but
the scalar case was still mishandled before this follow-up commit.

This is a follow up to OG12's 55a18d4744258e3909568e425f9f473c49f9d13f
While the problem is independent, it will be merged into v4 of the
mainline patch
'Fortran/OpenMP: Fix mapping of array descriptors and deferred-length strings'

gcc/fortran/

	* trans-openmp.cc (gfc_trans_omp_clauses): Fix unmapping of
	GOMP_MAP_POINTER for scalar allocatables/pointers.

gcc/testsuite/

	* gfortran.dg/gomp/map-10.f90: New test.

diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 38e69db65b3dfb194d3147eb1225b640e7838f64..663102d93295e41be0d65c9218562359c67a40a0 100644 (file)
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,8 @@
+2023-03-23  Tobias Burnus  <tobias@codesourcery.com>
+
+       * trans-openmp.cc (gfc_trans_omp_clauses): Fix unmapping of
+       GOMP_MAP_POINTER for scalar allocatables/pointers.
+
 2023-03-01  Tobias Burnus  <tobias@codesourcery.com>
 
        Backported from master:

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index acd8ce645bbbd3d6d0e899ac13c8e92c2b4e75bd..7a94bdcc870c3ec21cc97d10755f2fa2beb3e35d 100644 (file)
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -4678,15 +4678,27 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
                                                             cond, tmp,
                                                             NULL_TREE));
                        }
+                     /* For descriptor types, the unmapping happens below.  */
                      if (op != EXEC_OMP_TARGET_EXIT_DATA
-                         && n->u.map_op != OMP_MAP_RELEASE
-                         && n->u.map_op != OMP_MAP_DELETE)
+                         || !GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
                        {
                          node4 = build_omp_clause (input_location,
                                                    OMP_CLAUSE_MAP);
+                         if (gmk == GOMP_MAP_POINTER
+                             && op == EXEC_OMP_TARGET_EXIT_DATA
+                             && n->u.map_op == OMP_MAP_DELETE)
+                           gmk = GOMP_MAP_DELETE;
+                         else if (gmk == GOMP_MAP_POINTER
+                                  && op == EXEC_OMP_TARGET_EXIT_DATA)
+                           gmk = GOMP_MAP_RELEASE;
+                         tree size;
+                         if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE)
+                           size = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+                         else
+                           size = size_int (0);
                          OMP_CLAUSE_SET_MAP_KIND (node4, gmk);
                          OMP_CLAUSE_DECL (node4) = decl;
-                         OMP_CLAUSE_SIZE (node4) = size_int (0);
+                         OMP_CLAUSE_SIZE (node4) = size;
                        }
                      decl = build_fold_indirect_ref (decl);
                      if ((TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE
@@ -4694,16 +4706,24 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
                          && (GFC_DECL_GET_SCALAR_POINTER (orig_decl)
                              || GFC_DECL_GET_SCALAR_ALLOCATABLE (orig_decl)))
                        {
-                         if (op != EXEC_OMP_TARGET_EXIT_DATA
-                             && n->u.map_op != OMP_MAP_RELEASE
-                             && n->u.map_op != OMP_MAP_DELETE)
-                           {
-                             node3 = build_omp_clause (input_location,
-                                                       OMP_CLAUSE_MAP);
-                             OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
-                             OMP_CLAUSE_DECL (node3) = decl;
-                             OMP_CLAUSE_SIZE (node3) = size_int (0);
-                           }
+
+                         if (op == EXEC_OMP_TARGET_EXIT_DATA
+                             && n->u.map_op == OMP_MAP_DELETE)
+                           gmk = GOMP_MAP_DELETE;
+                         else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+                           gmk = GOMP_MAP_RELEASE;
+                         else
+                           gmk = GOMP_MAP_POINTER;
+                         tree size;
+                         if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE)
+                           size = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+                         else
+                           size = size_int (0);
+                         node3 = build_omp_clause (input_location,
+                                                   OMP_CLAUSE_MAP);
+                         OMP_CLAUSE_SET_MAP_KIND (node3, gmk);
+                         OMP_CLAUSE_DECL (node3) = decl;
+                         OMP_CLAUSE_SIZE (node3) = size;
                          decl = build_fold_indirect_ref (decl);
                        }
                    }

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 1ed7270346671b1f399607a7200226734ae59ca3..691c7f160f8892038cdaa65bbd388111a3d1bd43 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,7 @@
+2023-03-23  Tobias Burnus  <tobias@codesourcery.com>
+
+       * gfortran.dg/gomp/map-10.f90: New test.
+
 2023-03-22  Andrew Jenner  <andrew@codesourcery.com>
 
        * gcc.target/gcn/complex.c: New test.

diff --git a/gcc/testsuite/gfortran.dg/gomp/map-10.f90 b/gcc/testsuite/gfortran.dg/gomp/map-10.f90
new file mode 100644 (file)
index 0000000..33048d2
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-10.f90
@@ -0,0 +1,69 @@
+! { dg-additional-options "-fdump-tree-omplower" }
+
+! If enter data adds a (GOMP_MAP_)POINTER attachment, exit data needs to remove
+! it again. If not there can be all kind of issues, in particular when
+! stack memory was mapped, reused later and mapped again.
+
+subroutine test_aa (aa2, aa3)
+  integer(kind=4), allocatable :: aa1, aa2, aa3
+  optional :: aa3
+  !$omp target enter data map(aa1)
+  !$omp target exit data map(aa1)
+  !$omp target enter data map(aa2)
+  !$omp target exit data map(aa2)
+  !$omp target enter data map(aa3)
+  !$omp target exit data map(aa3)
+end
+
+subroutine test_pp (pp2, pp3)
+  integer(kind=4), allocatable :: pp1, pp2, pp3
+  optional :: pp3
+  !$omp target enter data map(pp1)
+  !$omp target exit data map(pp1)
+  !$omp target enter data map(pp2)
+  !$omp target exit data map(pp2)
+  !$omp target enter data map(pp3)
+  !$omp target exit data map(pp3)
+end
+
+subroutine test_pprelease (rp2, rp3)
+  integer(kind=4), allocatable :: rp1, rp2, rp3
+  optional :: rp3
+  !$omp target enter data map(rp1)
+  !$omp target exit data map(release:rp1)
+  !$omp target enter data map(rp2)
+  !$omp target exit data map(release:rp2)
+  !$omp target enter data map(rp3)
+  !$omp target exit data map(release:rp3)
+end
+
+subroutine test_ppdelete (dp2, dp3)
+  integer(kind=4), allocatable :: dp1, dp2, dp3
+  optional :: dp3
+  !$omp target enter data map(dp1)
+  !$omp target exit data map(delete:dp1)
+  !$omp target enter data map(dp2)
+  !$omp target exit data map(delete:dp2)
+  !$omp target enter data map(dp3)
+  !$omp target exit data map(delete:dp3)
+end
+
+
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*aa1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:aa1 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*aa1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa1 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*aa2.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:aa2 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:\\*aa2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:aa2 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*aa3.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:aa3 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:\\*aa3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:aa3 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*pp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:pp1 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*pp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp1 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*pp2.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:pp2 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:\\*pp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:pp2 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*pp3.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:pp3 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:\\*pp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:pp3 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:\\*rp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:rp1 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:\\*rp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:rp2 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:\\*rp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:rp3 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:\\*dp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(delete:dp1 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(delete:\\*dp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(delete:dp2 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(delete:\\*dp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(delete:dp3 \\\[len: .\\\]\\)" "omplower" } }