Re: [PATCH 7/9] [OpenACC] Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for enter/exit data directives

[Thomas Schwinge <thomas@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 6705 bytes --]

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

[-- Attachment #2: dynamic-pointer-1.f90 --]
[-- Type: text/plain, Size: 2325 bytes --]

! 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