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 <jul...@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
