When attaching pointers in Fortran, OpenACC 2.6 specifies that a descriptor must be copied to the target at the same time (see next patch). That means that stripping GOMP_MAP_TO_PSET (and lesserly, GOMP_MAP_POINTER), which was behaviour introduced by the manual deep-copy middle-end support patch, was probably wrong.
That arguably answers some of the questions at the end of: https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547424.html It appears that the user can (but certainly should not!) map a synthesized array descriptor using an "enter data" operation that can go out of scope before that data is unmapped. It would be nice to give a warning for an attempt to do such a thing, though I have no idea if that's possible in practice. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data directives. gcc/testsuite/ * gfortran.dg/goacc/finalize-1.f: Update expected dump output. --- gcc/gimplify.c | 11 ++--------- gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 ++-- 2 files changed, 4 insertions(+), 11 deletions(-) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 9851edfc4db..aa6853f0dcc 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8767,6 +8767,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: case OACC_HOST_DATA: if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -8775,15 +8777,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, mapped, but not the pointer to it. */ remove = true; break; - case OACC_ENTER_DATA: - case OACC_EXIT_DATA: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) - remove = true; - break; default: break; } diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f index 1e2e3e94b8a..ca642156e9f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -21,7 +21,7 @@ !$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: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { 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" } } @@ -33,5 +33,5 @@ !$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: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { 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 -- 2.23.0
