From: Tobias Burnus Date: Thu, 23 Mar 2023 07:57:45 +0000 (+0100) Subject: OpenMP/Fortran: Fix unmapping of GOMP_MAP_POINTER for scalar allocatables/pointers X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=8ea805840200f7dfd2c11b37abf5fbfe479c2fe2;p=thirdparty%2Fgcc.git 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 38e69db65b3d..663102d93295 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,8 @@ +2023-03-23 Tobias Burnus + + * trans-openmp.cc (gfc_trans_omp_clauses): Fix unmapping of + GOMP_MAP_POINTER for scalar allocatables/pointers. + 2023-03-01 Tobias Burnus Backported from master: diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index acd8ce645bbb..7a94bdcc870c 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 1ed727034667..691c7f160f88 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,7 @@ +2023-03-23 Tobias Burnus + + * gfortran.dg/gomp/map-10.f90: New test. + 2023-03-22 Andrew Jenner * 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 000000000000..33048d2a6904 --- /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" } }