[Bug middle-end/107517] New: [OpenMP][5.0] 'target update' with strides — for C/C++ and Fortran

[Bug middle-end/107517] New: [OpenMP][5.0] 'target update' with strides — for C/C++ and Fortran

[burnus at gcc dot gnu.org]

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

            Bug ID: 107517
           Summary: [OpenMP][5.0] 'target update' with strides — for C/C++
                    and Fortran
           Product: gcc
           Version: 13.0
            Status: UNCONFIRMED
          Keywords: openmp, rejects-valid
          Severity: normal
          Priority: P3
         Component: middle-end
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
  Target Milestone: ---

Cf. https://gcc.gnu.org/pipermail/gcc-patches/2022-November/604938.html

"on the library side we certainly have omp_target_memcpy_rect which under the
hood just does multiple copies of the contiguous subparts, but I don't remember
something similar done in GOMP_target_update.  And I think it is not ok to copy
bytes that aren't requested to be copied."


OpenMP 5.0 (and later) seems to permit list items with strides
for the 'to' and 'from' clause of 'target update', e.g.

For C/C++:
   int array[1];
    ....
   #pragma omp target update from(arr[::2])

For Fortran:

   arr(::2)  ! stride, but must be positive
   arr([1,7,2,4]) ! first index must be first in element order
but not
   dt(:)%comp ! Invalid - not last part-ref

OpenMP 5.2 wording for 'to'/'from':
"5.9 Data-Motion Clauses": "The list items may include array sections with
stride expressions."


OpenMP 5.0 - same in 5.2 (in "3.2.5 Array Sections")
C/C++:
"The stride must evaluate to a positive integer."

Fortran:
• If a stride expression is specified, it must be positive.
• The upper bound for the last dimension of an assumed-size dummy array must be
specified.
• If a list item is an array section with vector subscripts, the first array
element must be the lowest in the array element order of the array section.
• If a list item is an array section, the last part-ref of the list item must
have a section subscript list


[Fortran] Regarding substrings: "variable list item" does not permit them – and
this implies extended list items and locator list items may not have them
neither.
(See OpenMP 5.2's "3.2.1 OpenMP Argument Lists")


Unrelated to this issue but related to mapping:

[Fortran] In "5.1.1 Variables Referenced in a Construct", OpenMP 5.2 states:
* "A type parameter inquiry or complex part designator that is referenced in a
construct is treated as if its designator is referenced."

[Fortran] For the map clause, OpenMP 5.2 states:
* "A list item must not be a complex part designator."

[Bug middle-end/107517] [OpenMP][5.0] 'target update' with strides — for C/C++ and Fortran

[jakub at gcc dot gnu.org]

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

Jakub Jelinek <jakub at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |jakub at gcc dot gnu.org

--- Comment #1 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
OpenMP 4.5 had:
"If a list item is an array section it must specify contiguous storage."
restriction on target update, which got dropped in 5.0.
For the implementation, I think we have multiple possibilities.
One is to do what we'll need to do e.g. for map clause with iterators (that is
5.1 feature which we still don't implement), so basically arrange for having a
VLA for GOMP_target* arrays and fill them with runtime loops.
Or add some new GOMP_MAP_* to represent record that will describe some strided
mapping or to/from update, perhaps pointing to arguments like
omp_target_memcpy_rect has to describe, and let the library handle it.
Or a mixture thereof.
If it is easily describable, advantage of handling it on the library side would
be that we could use some smarter host to device or vice versa copying if a
plugin supports them (I think Cuda has some support but I think we don't use it
even for omp_target_memcpy_rect).