[Bug fortran/104949] New: [OpenMP] omp target: firstprivate of allocatable array – only descriptor firstprivatized

[Bug fortran/104949] New: [OpenMP] omp target: firstprivate of allocatable array – only descriptor firstprivatized

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104949

            Bug ID: 104949
           Summary: [OpenMP] omp target: firstprivate of allocatable array
                    – only descriptor firstprivatized
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Keywords: openmp, wrong-code
          Severity: normal
          Priority: P3
         Component: fortran
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
  Target Milestone: ---

Found via
https://github.com/SOLLVE/sollve_vv/blob/master/tests/5.0/target/test_target_defaultmap_firstprivate.F90

The line:
  !$omp target firstprivate(A)

is translated as

        .omp_data_arr.19.a = &a;
        #pragma omp target num_teams(1) thread_limit(0) firstprivate(a)
               [child fn: MAIN__._omp_fn.1 (.omp_data_arr.19,
.omp_data_sizes.20,
                                            .omp_data_kinds.21)]

Thus, it does not handle the descriptor. See also PR 90742.

 * * *

implicit none (type,external)
integer, allocatable :: A(:)
A = [1,2,3,4,5,6]

!$omp parallel firstprivate(A)
!$omp master
  if (any (A /= [1,2,3,4,5])) error stop
  A(:) = [99,88,77,66,55]
!$omp end master
!$omp end parallel

!$omp target firstprivate(A)
  if (any (A /= [1,2,3,4,5])) error stop
  A(:) = [99,88,77,66,55]
!$omp end target
if (any (A /= [1,2,3,4,5])) error stop
end

[Bug fortran/104949] [OpenMP] omp target: firstprivate of allocatable array – only descriptor firstprivatized

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104949

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
The following addition to testcase is needed.

! -------------------------------

!$omp parallel default(A)
!$omp master
  if (any (A /= [1,2,3,4,5])) error stop
  A(:) = [99,88,77,66,55]
!$omp end master
!$omp end parallel
if (any (A /= [1,2,3,4,5])) error stop

!$omp target defaultmap(firstprivate)
  if (any (A /= [1,2,3,4,5])) error stop
  A(:) = [99,88,77,66,55]
!$omp end target
if (any (A /= [1,2,3,4,5])) error stop

! -------------------------------

Reason: A different code path for
(a) defaultmap(firstprivate)  and  default(firstprivate)
    For them gfc_omp_finish_clause is invoked
(b) explicit firstprivate(a)
    Handled via gfc_trans_omp_clauses -> gfc_trans_omp_variable_list.
    but gfc_omp_finish_clause is not called (should it?)

The 'parallel' variant works – and for both variants the
gfc_omp_clause_copy_ctor is invoked for both via lower_rec_input_clauses()'s
the block which follows the 'do_firstprivate:' label.

Both 'target' variants only create 'firstprivate(a)' - but additionally
'firstprivate(a.data)' is needed, including doing a pointer attach.

  *  *  * 

Currently, we support either code using attach like:

  map(force_to:var) map(force_to:*var.data [len...]) map(attach_detach:var.data
[bias: 0])

or pointer-set/pointer assign like for

  map(to:var.data [len:...])
  map(to:var [pointer set, len: 64])
  map(alloc: var.data [pointer assign, bias: 0])

In my understanding, either code requires that 'var.data' can be found as
host->device lookup, which does not work for firstprivate. And just internally
adding 'a' to the host->device mapping for firstprivate purpose internally does
not handle the following.

a = 5
!$omp target enter data map(to: a)
a = 7

! adding internally 'a' (host)->'a'(firstprivate) mapping to find the 'a.var'
! for the usage in attach/pointer assign does not work as 'a' already exists.

!$omp target firstprivate(a) 
  ! a should be now 7
!$omp end firstprivate
!$omp target
   ! a is 5
!$omp end target


For the usage in
  type t
    integer, allocatable :: a, b
  end type t
  type(t) :: var
  !$omp target firstprivate(var)
    ...
multiple attachments are needed, i.e. just using firstprivate(a)
firstprivate(b) pointer_attach (a.data) and working with index i, i+1, i+2 does
not work.

[Bug fortran/104949] [OpenMP] omp target: firstprivate of allocatable array – only descriptor firstprivatized

[cvs-commit at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104949

--- Comment #2 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tobias Burnus <burnus@gcc.gnu.org>:

https://gcc.gnu.org/g:49d1a2f91325fa8cc011149e27e5093a988b3a49

commit r13-706-g49d1a2f91325fa8cc011149e27e5093a988b3a49
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Mon May 23 10:54:32 2022 +0200

    OpenMP: Handle descriptors in target's firstprivate [PR104949]

    For allocatable/pointer arrays, a firstprivate to a device
    not only needs to privatize the descriptor but also the actual
    data. This is implemented as:
      firstprivate(x) firstprivate(x.data) attach(x [bias: &x.data-&x)
    where the address of x in device memory is saved in hostaddrs[i]
    by libgomp and the middle end actually passes hostaddrs[i]' to
    attach.

    As side effect, has_device_addr(array_desc) had to be changed:
    before, it was converted to firstprivate in the front end; now
    it is handled in omp-low.cc as has_device_addr requires a shallow
    firstprivate (not touching the data pointer) while the normal
    firstprivate requires (now) a deep firstprivate.

    gcc/fortran/ChangeLog:

            PR fortran/104949
            * f95-lang.cc (LANG_HOOKS_OMP_ARRAY_SIZE): Redefine.
            * trans-openmp.cc (gfc_omp_array_size): New.
            (gfc_trans_omp_variable_list): Never turn has_device_addr
            to firstprivate.
            * trans.h (gfc_omp_array_size): New.

    gcc/ChangeLog:

            PR fortran/104949
            * langhooks-def.h (lhd_omp_array_size): New.
            (LANG_HOOKS_OMP_ARRAY_SIZE): Define.
            (LANG_HOOKS_DECLS): Add it.
            * langhooks.cc (lhd_omp_array_size): New.
            * langhooks.h (struct lang_hooks_for_decls): Add hook.
            * omp-low.cc (scan_sharing_clauses, lower_omp_target):
            Handle GOMP_MAP_FIRSTPRIVATE for array descriptors.

    libgomp/ChangeLog:

            PR fortran/104949
            * target.c (gomp_map_vars_internal, copy_firstprivate_data):
            Support attach for GOMP_MAP_FIRSTPRIVATE.
            * testsuite/libgomp.fortran/target-firstprivate-1.f90: New test.
            * testsuite/libgomp.fortran/target-firstprivate-2.f90: New test.
            * testsuite/libgomp.fortran/target-firstprivate-3.f90: New test.

[Bug fortran/104949] [OpenMP] omp target: firstprivate of allocatable array – only descriptor firstprivatized

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104949

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |FIXED
             Status|UNCONFIRMED                 |RESOLVED

--- Comment #3 from Tobias Burnus <burnus at gcc dot gnu.org> ---
FIXED on mainline (GCC 13).

[Bug fortran/104949] [OpenMP] omp target: firstprivate of allocatable array – only descriptor firstprivatized

[cvs-commit at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104949

--- Comment #4 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Thomas Schwinge <tschwinge@gcc.gnu.org>:

https://gcc.gnu.org/g:e8fec6998b656dac02d4bc6c69b35a0fb5611e87

commit r13-6852-ge8fec6998b656dac02d4bc6c69b35a0fb5611e87
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Thu Mar 23 12:32:35 2023 +0100

    Add caveat/safeguard to OpenMP: Handle descriptors in target's firstprivate
[PR104949]

    Follow-up to commit 49d1a2f91325fa8cc011149e27e5093a988b3a49
    "OpenMP: Handle descriptors in target's firstprivate [PR104949]".

            PR fortran/104949
            libgomp/
            * target.c (gomp_map_vars_internal) <GOMP_MAP_FIRSTPRIVATE>: Add
            caveat/safeguard.