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.

4 files changed, 111 insertions, 13 deletions

diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 38e69db..663102d 100644
--- 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 acd8ce6..7a94bdc 100644
--- 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 1ed7270..691c7f1 100644
--- 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
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" } }