[gcc/devel/omp/gcc-12] OpenMP/Fortran: Fix unmapping of GOMP_MAP_POINTER for scalar allocatables/pointers
Tobias Burnus
burnus@gcc.gnu.org
Thu Mar 23 09:27:16 GMT 2023
https://gcc.gnu.org/g:8ea805840200f7dfd2c11b37abf5fbfe479c2fe2
commit 8ea805840200f7dfd2c11b37abf5fbfe479c2fe2
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Thu Mar 23 08:57:45 2023 +0100
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:
---
gcc/fortran/ChangeLog.omp | 5 +++
gcc/fortran/trans-openmp.cc | 46 +++++++++++++++------
gcc/testsuite/ChangeLog.omp | 4 ++
gcc/testsuite/gfortran.dg/gomp/map-10.f90 | 69 +++++++++++++++++++++++++++++++
4 files changed, 111 insertions(+), 13 deletions(-)
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index 38e69db65b3..663102d9329 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 acd8ce645bb..7a94bdcc870 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 1ed72703466..691c7f160f8 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 00000000000..33048d2a690
--- /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" } }
More information about the Gcc-cvs
mailing list