| Message ID | e3b2342a0837402e281029142022ac39c5147018.1592343756.git.julian@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | Refcounting and manual deep copy improvements | expand |
Hi Julian! On 2020-06-16T15:39:43-0700, Julian Brown <julian@codesourcery.com> wrote: > 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 ACK. > 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. That's a rather complex scenario. ;-) If I'm understanding this right, what we need to show is that an object is created as a persistent, visible device copy, with state initialized by 'enter data', and then any 'GOMP_MAP_TO_PSET' etc. that come with each OpenACC 'parallel' etc. are no-ops (because the object is present already). My attached (new) 'libgomp.oacc-fortran/dynamic-pointer-1.f90' would seem to be a conceptually simple test case for this, using a Fortran 'pointer'. (I hope I got my Fortran right, please verify.) This test case doesn't work in current master and releases/gcc-10 branches (because we don't create the persistent, visible device copy), and is "enabled" by your patch posted here. I'm intentionally not saying "regression fixed" or something like that, because it also doesn't work before all the "OpenACC 2.6 deep copy: middle-end parts" etc. changes... (Maybe because of wrong handling of 'GOMP_MAP_TO_PSET' back then, too? Just mentioning that for completeness; I don't think we need to investigate that now.) Please include some such rationale in the commit log, or "even" as source code comments, as makes sense. This code surely is complicated/complex to grasp. > gcc/ > * gimplify.c (gimplify_scan_omp_clauses): Do not strip > GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data > directives. Should reference PR92929 here. Please include my attached (new) 'libgomp.oacc-fortran/dynamic-pointer-1.f90' (assuming that one makes sense to you), and then this is OK for master and releases/gcc-10 branches. We then (later) still need to resolve other items discussed in PR92929 "OpenACC/OpenMP 'target' 'exit data'/'update' optimizations". Grüße Thomas > 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 ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter ! Verify that a 'enter data'ed 'pointer' object creates a persistent, visible device copy ! { dg-do run } ! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } module m implicit none contains subroutine verify_a (a_ref, a) implicit none integer, dimension (:, :, :), allocatable :: a_ref integer, dimension (:, :, :), pointer :: a !$acc routine seq if (any (lbound (a) /= lbound (a_ref))) stop 101 if (any (ubound (a) /= ubound (a_ref))) stop 102 if (size (a) /= size (a_ref)) stop 103 end subroutine verify_a end module m program main use m use openacc implicit none integer, parameter :: n = 30 integer, dimension (:, :, :), allocatable, target :: a1, a2 integer, dimension (:, :, :), pointer :: p allocate (a1(1:n, 0:n-1, 10:n/2)) !$acc enter data create(a1) allocate (a2(3:n/3, 10:n, n-10:n+10)) !$acc enter data create(a2) p => a1 call verify_a(a1, p) ! 'p' object isn't present on the device. !$acc parallel ! Implicit 'copy(p)'; creates 'p' object... call verify_a(a1, p) !$acc end parallel ! ..., and deletes it again. p => a2 call verify_a(a2, p) ! 'p' object isn't present on the device. !$acc parallel ! Implicit 'copy(p)'; creates 'p' object... call verify_a(a2, p) !$acc end parallel ! ..., and deletes it again. p => a1 !$acc enter data create(p) ! 'p' object is now present on the device (visible device copy). !TODO PR96080 if (.not. acc_is_present (p)) stop 1 !$acc parallel ! On the device, got created as 'p => a1'. call verify_a(a1, p) !$acc end parallel call verify_a(a1, p) !$acc parallel p => a2 ! On the device, 'p => a2' is now set. call verify_a(a2, p) !$acc end parallel ! On the host, 'p => a1' persists. call verify_a(a1, p) !$acc parallel ! On the device, 'p => a2' persists. call verify_a(a2, p) !$acc end parallel ! On the host, 'p => a1' still persists. call verify_a(a1, p) p => a2 !$acc parallel p => a1 ! On the device, 'p => a1' is now set. call verify_a(a1, p) !$acc end parallel ! On the host, 'p => a2' persists. call verify_a(a2, p) !$acc parallel ! On the device, 'p => a1' persists. call verify_a(a1, p) !$acc end parallel ! On the host, 'p => a2' still persists. call verify_a(a2, p) end program main
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