Re: [PATCH v7 3/5] OpenMP: Pointers and member mappings

[Tobias Burnus <tobias@codesourcery.com>]

Hi Julian,

LGTM, except for:

* The 'target exit data' handling - comments below - looks a bit
fishy/inconsistent.

I intent to have a closer look with more code context, but maybe you
should have a look at it as well.

BTW: Fortran deep-mapping is not yet on mainline. Are you aware of
changes or in particular testcases on OG13 related to your patch series
that should be included when upstreaming that auto-mapping of
allocatable components patch?

* * *

On 19.08.23 00:47, Julian Brown wrote:
> This patch changes the mapping node arrangement used for array components
> of derived types in order to accommodate for changes made in the previous
> patch, particularly the use of "GOMP_MAP_ATTACH_DETACH" for pointer-typed
> derived-type members instead of "GOMP_MAP_ALWAYS_POINTER".
>
> We change the mapping nodes used for a derived-type mapping like this:
>
>    type T
>    integer, pointer, dimension(:) :: arrptr
>    end type T
>
>    type(T) :: tvar
>    [...]
>    !$omp target map(tofrom: tvar%arrptr)
>
> So that the nodes used look like this:
>
>    1) map(to: tvar%arrptr)   -->
>    GOMP_MAP_TO [implicit]  *tvar%arrptr%data  (the array data)
>    GOMP_MAP_TO_PSET        tvar%arrptr        (the descriptor)
>    GOMP_MAP_ATTACH_DETACH  tvar%arrptr%data
>
>    2) map(tofrom: tvar%arrptr(3:8)   -->
>    GOMP_MAP_TOFROM         *tvar%arrptr%data(3)  (size 8-3+1, etc.)

[Clarification for the patch readers - not to the patch writer]

At least to me the wording or implications of (1) were not fully clear to
me. However, it is intended that the implicit mapping ensures that the
pointer is actually mapped.

Fortunately, that's indeed the case :-)
as testing and also the new testcase libgomp/testsuite/libgomp.fortran/map-subarray-4.f90  shows.

The related part comes from OpenMP 5.1 [349:9-12] or slightly refined
later, here the TR12 wording:

If a list item in a map clause is an associated pointer that is not attach-ineligible and the pointer is
not the base pointer of another list item in a map clause on the same construct, then it is treated as if
its pointer target is implicitly mapped in the same clause. For the purposes of the map clause, the
mapped pointer target is treated as if its base pointer is the associated pointer."  [213:18-21]

* * *

> [...]
> In the middle end, we process mappings using the struct sibling-list
> handling machinery by moving the "GOMP_MAP_TO_PSET" node from the middle
> of the group of three mapping nodes to the proper sorted position after
> the GOMP_MAP_STRUCT mapping:
>
>    GOMP_MAP_STRUCT   tvar     (len: 1)
>    GOMP_MAP_TO_PSET  tvar%arr (size: 64, etc.)  <--. moved here
>    [...]                                           |
>    GOMP_MAP_TOFROM         *tvar%arrptr%data(3) ___|
>    GOMP_MAP_ATTACH_DETACH  tvar%arrptr%data
>
> In another case, if we have an array of derived-type values "dtarr",
> and mappings like:
>
>    i = 1
>    j = 1
>    map(to: dtarr(i)%arrptr) map(tofrom: dtarr(j)%arrptr(3:8))
>
> We still map the same way, but this time we cannot prove that the base
> expressions "dtarr(i) and "dtarr(j)" are the same in the front-end.
> So we keep both mappings, but we move the "[implicit]" mapping of the
> full-array reference to the end of the clause list in gimplify.cc (by
> adjusting the topological sorting algorithm):
>
>    GOMP_MAP_STRUCT         dtvar  (len: 2)
>    GOMP_MAP_TO_PSET        dtvar(i)%arrptr
>    GOMP_MAP_TO_PSET        dtvar(j)%arrptr
>    [...]
>    GOMP_MAP_TOFROM         *dtvar(j)%arrptr%data(3)  (size: 8-3+1)
>    GOMP_MAP_ATTACH_DETACH  dtvar(j)%arrptr%data
>    GOMP_MAP_TO [implicit]  *dtvar(i)%arrptr%data(1)  (size: whole array)
>    GOMP_MAP_ATTACH_DETACH  dtvar(i)%arrptr%data
>
> Always moving "[implicit]" full-array mappings after array-section
> mappings (without that bit set) means that we'll avoid copying the whole
> array unnecessarily -- even in cases where we can't prove that the arrays
> are the same.
>
> The patch also fixes some bugs with "enter data" and "exit data"
> directives with this new mapping arrangement.  Also now if you have
> mappings like this:
>
>    #pragma omp target enter data map(to: dv, dv%arr(1:20))
>
> The whole of the derived-type variable "dv" is mapped, so the
> GOMP_MAP_TO_PSET for the array-section mapping can be dropped:
>
>    GOMP_MAP_TO            dv
>
>    GOMP_MAP_TO            *dv%arr%data
>    GOMP_MAP_TO_PSET       dv%arr <-- deleted (array section mapping)
>    GOMP_MAP_ATTACH_DETACH dv%arr%data
>
> To accommodate for recent changes to mapping nodes made by
> Tobias, this version of the patch avoids using GOMP_MAP_TO_PSET
> for "exit data" directives, in favour of using the "correct"
> GOMP_MAP_RELEASE/GOMP_MAP_DELETE kinds during early expansion.  A new
> flag is introduced so the middle-end knows when the latter two kinds
> are being used specifically for an array descriptor.
>
> This version of the patch is based on the version posted for the og13
> branch:
>
> https://gcc.gnu.org/pipermail/gcc-patches/2023-June/622222.html
[...]

> +      if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE
> +       || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
> +       || op == EXEC_OMP_TARGET_EXIT_DATA)
>       {
> [...]
> +       gomp_map_kind map_kind
> +         = (op == EXEC_OMP_TARGET_EXIT_DATA) ? GOMP_MAP_RELEASE
> +                                             : OMP_CLAUSE_MAP_KIND (node);
> +       OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
> +       OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) = 1;

For '!$omp target exit data map(delete: array)' this looks wrong as it
replaces 'delete' by 'release' for the descriptor - while for
  '!$omp target (data) map(delete: array)'
it remains 'delete'.

Thus, I wonder whether that shouldn't be instead
   OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE
   ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE;

* * *

We later have the following; just reading the patch, I wonder whether GOMP_TO_PSET is correct
for a generic 'target exit data' here or not.  It seems at a glance as
if an "|| op == 'target exit data'" is missing here:

> -                       if (openacc)
> -                         OMP_CLAUSE_SET_MAP_KIND (desc_node,
> +                       if (openacc
> +                           || (map_kind != GOMP_MAP_RELEASE
> +                               && map_kind != GOMP_MAP_DELETE))
> +                         OMP_CLAUSE_SET_MAP_KIND (node2,
>                                                    GOMP_MAP_TO_PSET);
>                         else
> -                         OMP_CLAUSE_SET_MAP_KIND (desc_node, map_kind);
> -                       OMP_CLAUSE_DECL (desc_node) = inner;
> -                       OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
> -                       if (openacc)
> -                         node2 = desc_node;
> -                       else
> +                         OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);

(Here, without setting OMP_CLAUSE_RELEASE_DESCRIPTOR.)

And for completeness, we later have:

> +omp_map_clause_descriptor_p (tree c)
> +{
> +  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
> +    return false;
> +
> +  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
> +    return true;
> +
> +  if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_RELEASE
> +       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DELETE)
> +      && OMP_CLAUSE_RELEASE_DESCRIPTOR (c))
> +    return true;


* * *

> +                               /* Save array descriptor for use
> +                                  in gfc_omp_deep_mapping{,_p,_cnt}; force
> +                                  evaluate to ensure that it is
> +                                  not gimplified + is a decl.  */
This is part of my Fortran allocatable-components deep-mapping patch that is currently
only on OG9 (?) to OG13, badly needs to be upstreamed but required that Jakub had a
look at it - well, I still would like that he has a look at the omp-low.cc parts and
it needs to be re-diffed.

Hence, while I wouldn't mind to keep it to avoid more diffing work, I think it will
be cleaner not to keep it here.

* * *

Thanks,

Tobias

> 2023-08-18  Julian Brown  <julian@codesourcery.com>
>
> gcc/fortran/
>          * dependency.cc (gfc_omp_expr_prefix_same): New function.
>          * dependency.h (gfc_omp_expr_prefix_same): Add prototype.
>          * gfortran.h (gfc_omp_namelist): Add "duplicate_of" field to "u2"
>          union.
>          * trans-openmp.cc (dependency.h): Include.
>          (gfc_trans_omp_array_section): Adjust mapping node arrangement for
>          array descriptors.  Use GOMP_MAP_TO_PSET or
>          GOMP_MAP_RELEASE/GOMP_MAP_DELETE with the OMP_CLAUSE_RELEASE_DESCRIPTOR
>          flag set.
>          (gfc_symbol_rooted_namelist): New function.
>          (gfc_trans_omp_clauses): Check subcomponent and subarray/element
>          accesses elsewhere in the clause list for pointers to derived types or
>          array descriptors, and adjust or drop mapping nodes appropriately.
>          Adjust for changes to mapping node arrangement.
>          (gfc_trans_oacc_executable_directive): Pass code op through.
>
> gcc/
>          * gimplify.cc (omp_map_clause_descriptor_p): New function.
>          (build_omp_struct_comp_nodes, omp_get_attachment, omp_group_base): Use
>          above function.
>          (omp_tsort_mapping_groups): Process nodes that have
>          OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P set after those that don't.  Add
>          enter_exit_data parameter.
>          (omp_resolve_clause_dependencies): Remove GOMP_MAP_TO_PSET mappings if
>          we're mapping the whole containing derived-type variable.
>          (omp_accumulate_sibling_list): Adjust GOMP_MAP_TO_PSET handling.
>          Remove GOMP_MAP_ALWAYS_POINTER handling.
>          (gimplify_scan_omp_clauses): Pass enter_exit argument to
>          omp_tsort_mapping_groups.  Don't adjust/remove GOMP_MAP_TO_PSET
>          mappings for derived-type components here.
>          * tree.h (OMP_CLAUSE_RELEASE_DESCRIPTOR): New macro.
>
> gcc/testsuite/
>          * gfortran.dg/gomp/map-9.f90: Adjust scan output.
>          * gfortran.dg/gomp/map-subarray-2.f90: New test.
>          * gfortran.dg/gomp/map-subarray.f90: New test.
>
> libgomp/
>          * testsuite/libgomp.fortran/map-subarray.f90: New test.
>          * testsuite/libgomp.fortran/map-subarray-2.f90: New test.
>          * testsuite/libgomp.fortran/map-subarray-3.f90: New test.
>          * testsuite/libgomp.fortran/map-subarray-4.f90: New test.
>          * testsuite/libgomp.fortran/map-subarray-6.f90: New test.
>          * testsuite/libgomp.fortran/map-subarray-7.f90: New test.
>          * testsuite/libgomp.fortran/map-subarray-8.f90: New test.
>          * testsuite/libgomp.fortran/map-subcomponents.f90: New test.
>          * testsuite/libgomp.fortran/struct-elem-map-1.f90: Adjust for
>          descriptor-mapping changes.  Remove XFAIL.
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955