Consider the following OMP directive, assuming tiles is allocatable:
!$omp target enter data &
!$omp map(to: chunk%tiles(1)%field%density0) &
!$omp map(to: chunk%left_rcv_buffer)
libgomp reports an illegal memory access error at runtime. This is because
density0 is referenced through tiles, which requires its descriptor to be mapped
along its content.
This patch ensures that such intervening allocatable in a reference chain is
properly mapped. For the above example, the frontend has to create the following
three additional map clauses:
(1) map (alloc: *(struct tile_type[0:] * restrict) chunk.tiles.data [len: 0])
(2) map (to: chunk.tiles [pointer set, len: 64])
(3) map (attach_detach: (struct tile_type[0:] * restrict) chunk.tiles.data
[bias: -1])
(1) will turn into a no-op at runtime because the inner component is explicitly
to-mapped but alloc is required at compile time for attaching. (2) ensures that
the array descriptor will be available at runtime to compute offsets and strides
in various dimensions. The gimplifier will turn (3) into a regular attach of the
data pointer and compute the bias.
PR fortran/120505
gcc/fortran/ChangeLog:
* trans-openmp.cc (gfc_map_array_descriptor): New function.
(gfc_trans_omp_clauses): Emit map clauses for an intermediate array
descriptor.
gcc/ChangeLog:
* gimplify.cc (omp_mapped_by_containing_struct): Handle Fortran array
descriptors.
(omp_build_struct_sibling_lists): Allow attach_detach bias to be
adjusted on non-target regions.
* tree-core.h (OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT): Define.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/map-subarray-11.f90: New test.
* testsuite/libgomp.fortran/map-subarray-13.f90: New test.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/map-subarray-3.f90: New test.
* gfortran.dg/gomp/map-subarray-5.f90: New test.
---
gcc/fortran/trans-openmp.cc | 345 ++++++++++--------
gcc/gimplify.cc | 37 +-
.../gfortran.dg/gomp/map-subarray-3.f90 | 49 +++
.../gfortran.dg/gomp/map-subarray-5.f90 | 47 +++
gcc/tree-core.h | 1 +
.../libgomp.fortran/map-subarray-11.f90 | 56 +++
.../libgomp.fortran/map-subarray-13.f90 | 50 +++
7 files changed, 434 insertions(+), 151 deletions(-)
create mode 100644 gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/map-subarray-11.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/map-subarray-13.f90
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 254fc934af1..6bfee916365 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3531,6 +3531,162 @@ get_symbol_rooted_namelist (hash_map<gfc_symbol *,
return NULL;
}
+/* Helper function for gfc_trans_omp_clauses. */
+
+static bool
+gfc_map_array_descriptor (
+ tree &node, tree &node2, tree &node3, tree &node4, tree descr, bool openacc,
+ location_t map_loc, stmtblock_t *block, gfc_exec_op op, gfc_omp_namelist *n,
+ hash_map<gfc_symbol *, gfc_omp_namelist *> *&sym_rooted_nl, gfc_se se,
+ gfc_omp_clauses *clauses, bool mid_desc_p)
+{
+ tree type = TREE_TYPE (descr);
+ tree ptr = gfc_conv_descriptor_data_get (descr);
+ ptr = build_fold_indirect_ref (ptr);
+ OMP_CLAUSE_DECL (node) = ptr;
+ int rank = GFC_TYPE_ARRAY_RANK (type);
+ OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, descr, rank);
+ tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+
+ gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (node);
+ if (GOMP_MAP_COPY_TO_P (map_kind) || map_kind == GOMP_MAP_ALLOC)
+ {
+ if (mid_desc_p)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_SIZE (node) = size_int (0);
+ }
+ else
+ map_kind
+ = ((GOMP_MAP_ALWAYS_P (map_kind) || gfc_expr_attr (n->expr).pointer)
+ ? GOMP_MAP_ALWAYS_TO
+ : GOMP_MAP_TO);
+ }
+ else if (n->u.map.op == OMP_MAP_RELEASE || n->u.map.op == OMP_MAP_DELETE)
+ ;
+ else if (op == EXEC_OMP_TARGET_EXIT_DATA || op == EXEC_OACC_EXIT_DATA)
+ map_kind = GOMP_MAP_RELEASE;
+ else
+ map_kind = GOMP_MAP_ALLOC;
+
+ if (!openacc && n->expr->ts.type == BT_CHARACTER && n->expr->ts.deferred)
+ {
+ gcc_assert (se.string_length);
+ tree len = fold_convert (size_type_node, se.string_length);
+ elemsz = gfc_get_char_type (n->expr->ts.kind);
+ elemsz = TYPE_SIZE_UNIT (elemsz);
+ elemsz = fold_build2 (MULT_EXPR, size_type_node, len, elemsz);
+ node4 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
+ OMP_CLAUSE_DECL (node4) = se.string_length;
+ OMP_CLAUSE_SIZE (node4) = TYPE_SIZE_UNIT (gfc_charlen_type_node);
+ }
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ OMP_CLAUSE_SIZE (node) = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ OMP_CLAUSE_SIZE (node), elemsz);
+
+ node2 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
+ if (map_kind == GOMP_MAP_RELEASE || map_kind == GOMP_MAP_DELETE)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
+ OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) = 1;
+ }
+ else
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_DECL (node2) = descr;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+
+ if (!openacc)
+ {
+ if (n->expr->ts.type == BT_DERIVED
+ && n->expr->ts.u.derived->attr.alloc_comp)
+ {
+ /* Save array descriptor for use
+ in gfc_omp_deep_mapping{,_p,_cnt}; force
+ evaluate to ensure that it is
+ not gimplified + is a decl. */
+ tree tmp = OMP_CLAUSE_SIZE (node);
+ tree var = gfc_create_var (TREE_TYPE (tmp), NULL);
+ gfc_add_modify_loc (map_loc, block, var, tmp);
+ OMP_CLAUSE_SIZE (node) = var;
+ gfc_allocate_lang_decl (var);
+ GFC_DECL_SAVED_DESCRIPTOR (var) = descr;
+ }
+
+ gfc_omp_namelist *n2 = clauses->lists[OMP_LIST_MAP];
+
+ /* If we don't have a mapping of a smaller part
+ of the array -- or we can't prove that we do
+ statically -- set this flag. If there is a
+ mapping of a smaller part of the array after
+ all, this will turn into a no-op at
+ runtime. */
+ OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (node) = 1;
+
+ bool sym_based;
+ n2 = get_symbol_rooted_namelist (sym_rooted_nl, n, n2, &sym_based);
+
+ bool drop_mapping = false;
+
+ for (; n2 != NULL; n2 = n2->next)
+ {
+ if ((!sym_based && n == n2) || (sym_based && n == n2->u2.duplicate_of)
+ || !n2->expr)
+ continue;
+
+ if (!gfc_omp_expr_prefix_same (n->expr, n2->expr))
+ continue;
+
+ gfc_ref *ref1 = n->expr->ref;
+ gfc_ref *ref2 = n2->expr->ref;
+
+ /* We know ref1 and ref2 overlap. We're
+ interested in whether ref2 describes a
+ smaller part of the array than ref1, which
+ we already know refers to the full
+ array. */
+
+ while (ref1->next && ref2->next)
+ {
+ ref1 = ref1->next;
+ ref2 = ref2->next;
+ }
+
+ if (ref2->next
+ || (ref2->type == REF_ARRAY
+ && (ref2->u.ar.type == AR_ELEMENT
+ || (ref2->u.ar.type == AR_SECTION))))
+ {
+ drop_mapping = true;
+ break;
+ }
+ }
+ if (drop_mapping)
+ return true;
+ }
+
+ if (mid_desc_p && GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (node)))
+ node = NULL_TREE;
+
+ node3 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH_DETACH);
+ OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (descr);
+ /* Similar to gfc_trans_omp_array_section (details
+ there), we add/keep the cast for OpenMP to prevent
+ that an 'alloc:' gets added for node3 ('desc.data')
+ as that is part of the whole descriptor (node3).
+ TODO: Remove once the ME handles this properly. */
+ if (!openacc)
+ OMP_CLAUSE_DECL (node3) = fold_convert (TREE_TYPE (TREE_OPERAND (ptr, 0)),
+ OMP_CLAUSE_DECL (node3));
+ else
+ STRIP_NOPS (OMP_CLAUSE_DECL (node3));
+ OMP_CLAUSE_SIZE (node3)
+ = size_int (mid_desc_p ? OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT : 0);
+
+ return false;
+}
+
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false,
@@ -3544,6 +3700,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
enum omp_clause_code clause_code;
gfc_omp_namelist *prev = NULL;
gfc_se se;
+ vec<gfc_symbol *> descriptors = vNULL;
if (clauses == NULL)
return NULL_TREE;
@@ -4645,6 +4802,8 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
{
gfc_init_se (&se, NULL);
se.expr = gfc_maybe_dereference_var (n->sym, decl);
+ tree mid_descr = NULL_TREE;
+ gfc_ref *midref = NULL;
for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
{
@@ -4654,6 +4813,12 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
conv_parent_component_references (&se, ref);
gfc_conv_component_ref (&se, ref);
+ if (!mid_descr
+ && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (se.expr)))
+ {
+ mid_descr = se.expr;
+ midref = ref;
+ }
}
else if (ref->type == REF_ARRAY)
{
@@ -4807,156 +4972,11 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
{
- gomp_map_kind map_kind;
- tree type = TREE_TYPE (inner);
- tree ptr = gfc_conv_descriptor_data_get (inner);
- ptr = build_fold_indirect_ref (ptr);
- OMP_CLAUSE_DECL (node) = ptr;
- int rank = GFC_TYPE_ARRAY_RANK (type);
- OMP_CLAUSE_SIZE (node)
- = gfc_full_array_size (block, inner, rank);
- tree elemsz
- = TYPE_SIZE_UNIT (gfc_get_element_type (type));
- map_kind = OMP_CLAUSE_MAP_KIND (node);
- if (GOMP_MAP_COPY_TO_P (map_kind)
- || map_kind == GOMP_MAP_ALLOC)
- map_kind = ((GOMP_MAP_ALWAYS_P (map_kind)
- || gfc_expr_attr (n->expr).pointer)
- ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO);
- else if (n->u.map.op == OMP_MAP_RELEASE
- || n->u.map.op == OMP_MAP_DELETE)
- ;
- else if (op == EXEC_OMP_TARGET_EXIT_DATA
- || op == EXEC_OACC_EXIT_DATA)
- map_kind = GOMP_MAP_RELEASE;
- else
- map_kind = GOMP_MAP_ALLOC;
- if (!openacc
- && n->expr->ts.type == BT_CHARACTER
- && n->expr->ts.deferred)
- {
- gcc_assert (se.string_length);
- tree len = fold_convert (size_type_node,
- se.string_length);
- elemsz = gfc_get_char_type (n->expr->ts.kind);
- elemsz = TYPE_SIZE_UNIT (elemsz);
- elemsz = fold_build2 (MULT_EXPR, size_type_node,
- len, elemsz);
- node4 = build_omp_clause (map_loc,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
- OMP_CLAUSE_DECL (node4) = se.string_length;
- OMP_CLAUSE_SIZE (node4)
- = TYPE_SIZE_UNIT (gfc_charlen_type_node);
- }
- elemsz = fold_convert (gfc_array_index_type, elemsz);
- OMP_CLAUSE_SIZE (node)
- = fold_build2 (MULT_EXPR, gfc_array_index_type,
- OMP_CLAUSE_SIZE (node), elemsz);
- node2 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
- if (map_kind == GOMP_MAP_RELEASE
- || map_kind == GOMP_MAP_DELETE)
- {
- OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
- OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) = 1;
- }
- else
- OMP_CLAUSE_SET_MAP_KIND (node2,
- GOMP_MAP_TO_PSET);
- OMP_CLAUSE_DECL (node2) = inner;
- OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
- if (!openacc)
- {
- if (n->expr->ts.type == BT_DERIVED
- && n->expr->ts.u.derived->attr.alloc_comp)
- {
- /* Save array descriptor for use
- in gfc_omp_deep_mapping{,_p,_cnt}; force
- evaluate to ensure that it is
- not gimplified + is a decl. */
- tree tmp = OMP_CLAUSE_SIZE (node);
- tree var = gfc_create_var (TREE_TYPE (tmp),
- NULL);
- gfc_add_modify_loc (map_loc, block,
- var, tmp);
- OMP_CLAUSE_SIZE (node) = var;
- gfc_allocate_lang_decl (var);
- GFC_DECL_SAVED_DESCRIPTOR (var) = inner;
- }
-
- gfc_omp_namelist *n2
- = clauses->lists[OMP_LIST_MAP];
-
- /* If we don't have a mapping of a smaller part
- of the array -- or we can't prove that we do
- statically -- set this flag. If there is a
- mapping of a smaller part of the array after
- all, this will turn into a no-op at
- runtime. */
- OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (node) = 1;
-
- bool sym_based;
- n2 = get_symbol_rooted_namelist (sym_rooted_nl,
- n, n2,
- &sym_based);
-
- bool drop_mapping = false;
-
- for (; n2 != NULL; n2 = n2->next)
- {
- if ((!sym_based && n == n2)
- || (sym_based && n == n2->u2.duplicate_of)
- || !n2->expr)
- continue;
-
- if (!gfc_omp_expr_prefix_same (n->expr,
- n2->expr))
- continue;
-
- gfc_ref *ref1 = n->expr->ref;
- gfc_ref *ref2 = n2->expr->ref;
-
- /* We know ref1 and ref2 overlap. We're
- interested in whether ref2 describes a
- smaller part of the array than ref1, which
- we already know refers to the full
- array. */
-
- while (ref1->next && ref2->next)
- {
- ref1 = ref1->next;
- ref2 = ref2->next;
- }
-
- if (ref2->next
- || (ref2->type == REF_ARRAY
- && (ref2->u.ar.type == AR_ELEMENT
- || (ref2->u.ar.type
- == AR_SECTION))))
- {
- drop_mapping = true;
- break;
- }
- }
- if (drop_mapping)
- continue;
- }
- node3 = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3,
- GOMP_MAP_ATTACH_DETACH);
- OMP_CLAUSE_DECL (node3)
- = gfc_conv_descriptor_data_get (inner);
- /* Similar to gfc_trans_omp_array_section (details
- there), we add/keep the cast for OpenMP to prevent
- that an 'alloc:' gets added for node3 ('desc.data')
- as that is part of the whole descriptor (node3).
- TODO: Remove once the ME handles this properly. */
- if (!openacc)
- OMP_CLAUSE_DECL (node3)
- = fold_convert (TREE_TYPE (TREE_OPERAND(ptr,
0)),
- OMP_CLAUSE_DECL (node3));
- else
- STRIP_NOPS (OMP_CLAUSE_DECL (node3));
- OMP_CLAUSE_SIZE (node3) = size_int (0);
+ bool drop_mapping = gfc_map_array_descriptor (
+ node, node2, node3, node4, inner, openacc, map_loc,
+ block, op, n, sym_rooted_nl, se, clauses, false);
+ if (drop_mapping)
+ continue;
}
else
OMP_CLAUSE_DECL (node) = inner;
@@ -4972,6 +4992,31 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
}
else
gcc_unreachable ();
+
+ /* Map intermediate array descriptor. */
+ if (!openacc && mid_descr != NULL_TREE && mid_descr != inner
+ && !descriptors.contains (midref->u.c.sym))
+ {
+ descriptors.safe_push (midref->u.c.sym);
+
+ tree node1 = copy_node (node);
+ tree node2 = NULL_TREE;
+ tree node3 = NULL_TREE;
+ tree node4 = NULL_TREE;
+ gfc_map_array_descriptor (node1, node2, node3, node4,
+ mid_descr, openacc, map_loc,
+ block, op, n, sym_rooted_nl, se,
+ clauses, true);
+
+ if (node1 != NULL_TREE)
+ omp_clauses = gfc_trans_add_clause (node1, omp_clauses);
+ if (node2 != NULL_TREE)
+ omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
+ if (node3 != NULL_TREE)
+ omp_clauses = gfc_trans_add_clause (node3, omp_clauses);
+ if (node4 != NULL_TREE)
+ omp_clauses = gfc_trans_add_clause (node4, omp_clauses);
+ }
}
else
sorry_at (gfc_get_location (&n->where), "unhandled expression");
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 751b4697271..8559b7ff2ad 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -11511,9 +11511,13 @@ omp_mapped_by_containing_struct
(hash_map<tree_operand_hash_no_se,
}
if (wholestruct)
{
+ tree desc = OMP_CLAUSE_CHAIN (*(*wholestruct)->grp_start);
+ if (desc != NULL_TREE && omp_map_clause_descriptor_p (desc))
+ goto next;
*mapped_by_group = *wholestruct;
return true;
}
+ next:
decl = wsdecl;
}
@@ -13390,6 +13394,35 @@ omp_build_struct_sibling_lists (enum tree_code code,
tail = added_tail;
}
+ /* Find each attach node whose bias needs to be adjusted and move it to the
+ * group containing its pointee, right after the struct node. */
+ FOR_EACH_VEC_ELT (*groups, i, grp)
+ {
+ tree c = *grp->grp_start;
+ if (c != NULL && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+ && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) == GOMP_MAP_TO_PSET
+ && OMP_CLAUSE_MAP_KIND (grp->grp_end) == GOMP_MAP_ATTACH_DETACH
+ && OMP_CLAUSE_SIZE (grp->grp_end)
+ == size_int (OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT))
+ {
+ tree *cp;
+ for (cp = &OMP_CLAUSE_CHAIN (c); cp != NULL;
+ cp = &OMP_CLAUSE_CHAIN (*cp))
+ if (*cp == grp->grp_end)
+ {
+ c = *cp;
+ break;
+ }
+
+ tree base = OMP_CLAUSE_DECL (c);
+ gcc_assert (TREE_CODE (base) == NOP_EXPR);
+ base = build_fold_indirect_ref (base);
+ tree *struct_node = struct_map_to_clause->get (base);
+ omp_siblist_move_node_after (c, cp, &OMP_CLAUSE_CHAIN (*struct_node));
+ }
+ }
+
/* Now we have finished building the struct sibling lists, reprocess
newly-added "attach" nodes: we need the address of the first
mapped element of each struct sibling list for the bias of the attach
@@ -13416,7 +13449,9 @@ omp_build_struct_sibling_lists (enum tree_code code,
base they attach to). We should only have created the
ATTACH_DETACH node after GOMP_MAP_STRUCT for a target region, so
this should never be true. */
- gcc_assert ((region_type & ORT_TARGET) != 0);
+ // This is no longer true. See zlas in gomp_map_vars_internal
+ // (libgomp/target.c).
+ // gcc_assert ((region_type & ORT_TARGET) != 0);
/* This is the first sorted node in the struct sibling list. Use it
to recalculate the correct bias to use.
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90
b/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90
new file mode 100644
index 00000000000..b2bec2501fa
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90
@@ -0,0 +1,49 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+! PR fortran/120505
+
+! Check that mapping nested allocatable DT components triggers required
+! additional mappings for the outer array descriptor.
+
+module m
+
+ type field_type
+ real(kind=8), allocatable :: density0(:,:), density1(:,:)
+ end type field_type
+
+ type tile_type
+ type(field_type) :: field
+ end type tile_type
+
+ type chunk_type
+ real(kind=8), allocatable :: left_rcv_buffer(:)
+ type(tile_type), allocatable :: tiles(:)
+ end type chunk_type
+
+ type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density0 = reshape([1,2,3,4],[2,2])
+chunk%tiles(1)%field%density1 = reshape([5,6,7,8],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp map(to: chunk%tiles(1)%field%density0) &
+!$omp map(to: chunk%tiles(1)%field%density1) &
+!$omp map(to: chunk%left_rcv_buffer)
+
+! { dg-final { scan-tree-dump-times { #pragma omp target enter data
map\(alloc:\*\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[len:
0\] \[runtime_implicit\]\) map\(to:chunk\.tiles \[pointer set, len: 64\]\)
map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data
\[bias: [0-9]+\]\) } 1 "original" } }
+
+!$omp target exit data &
+!$omp map(from: chunk%tiles(1)%field%density0) &
+!$omp map(from: chunk%tiles(1)%field%density1) &
+!$omp map(from: chunk%left_rcv_buffer)
+
+! { dg-final { scan-tree-dump-times { #pragma omp target exit data
map\(release:chunk\.tiles \[pointer set, len: 64\]\)
map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data
\[bias: [0-9]+\]\) } 1 "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90
b/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90
new file mode 100644
index 00000000000..a44c4d72960
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90
@@ -0,0 +1,47 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+! PR fortran/120505
+
+! Check that the bias into the inner derived type is correctly computed on
+! target enter data. For target exit data, the bias is ignored so just check
+! that detach is present.
+
+module m
+
+ type field_type
+ real(kind=8), allocatable :: density0(:,:), density1(:,:)
+ end type field_type
+
+ type tile_type
+ type(field_type) :: field
+ end type tile_type
+
+ type chunk_type
+ real(kind=8), allocatable :: left_rcv_buffer(:)
+ type(tile_type), allocatable :: tiles(:)
+ end type chunk_type
+
+ type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp map(to: chunk%tiles(1)%field%density1) &
+!$omp map(to: chunk%left_rcv_buffer)
+
+! { dg-final { scan-tree-dump-times { map\(struct_unord:MEM <struct
tile_type\[0:\]> \[\(struct tile_type\[0:\] \*\)_[0-9]+\] \[len: 1\]\)
map\(to:MEM <struct tile_type\[0:\]> \[\(struct tile_type\[0:\]
\*\)_[0-9]+\]\[_[0-9]+\]\.field\.density1 \[pointer set, len: 88\]\)
map\(attach:chunk\.tiles\.data \[bias: _[0-9]+\]\) } 1 "gimple" } }
+
+!$omp target exit data &
+!$omp map(from: chunk%tiles(1)%field%density1) &
+!$omp map(from: chunk%left_rcv_buffer)
+
+! { dg-final { scan-tree-dump-times { map\(release:chunk\.tiles \[pointer set,
len: 64\]\) map\(detach:chunk\.tiles\.data \[bias: [0-9]+\]\)} 1 "gimple" } }
+
+end
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 7bdb474a253..555b54b7ced 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -667,6 +667,7 @@ enum omp_clause_fallback_kind {
OMP_CLAUSE_FALLBACK_NULL
};
+#define OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT -1
/* memory-order-clause on OpenMP atomic/flush constructs or
argument of atomic_default_mem_order clause. */
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90
b/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90
new file mode 100644
index 00000000000..d81a31491bb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90
@@ -0,0 +1,56 @@
+! { dg-do run }
+
+! PR fortran/120505
+
+! Check that mapping nested allocatable DT components triggers required
+! additional mappings for the outer array descriptor.
+
+module m
+
+ type field_type
+ real(kind=8), allocatable :: density0(:,:), density1(:,:)
+ end type field_type
+
+ type tile_type
+ type(field_type) :: field
+ end type tile_type
+
+ type chunk_type
+ real(kind=8), allocatable :: left_rcv_buffer(:)
+ type(tile_type), allocatable :: tiles(:)
+ end type chunk_type
+
+ type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density0 = reshape([1,2,3,4],[2,2])
+chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp map(to: chunk%tiles(1)%field%density0) &
+!$omp map(to: chunk%tiles(1)%field%density1) &
+!$omp map(to: chunk%left_rcv_buffer)
+
+!$omp target
+ if (any (chunk%tiles(1)%field%density0 /= reshape([1,2,3,4],[2,2]))) stop 1
+ if (any (chunk%tiles(1)%field%density1 /= reshape([1,2,3,4],[2,2]))) stop 1
+ chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 + 7
+ chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 + 5
+ chunk%left_rcv_buffer(1) = 42.0_8
+!$omp end target
+
+!$omp target exit data &
+!$omp map(from: chunk%tiles(1)%field%density0) &
+!$omp map(from: chunk%tiles(1)%field%density1) &
+!$omp map(from: chunk%left_rcv_buffer)
+
+if (any (chunk%tiles(1)%field%density0 /= 7 + reshape([1,2,3,4],[2,2]))) stop 1
+if (any (chunk%tiles(1)%field%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 1
+if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 1
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90
b/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90
new file mode 100644
index 00000000000..13fad59415b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90
@@ -0,0 +1,50 @@
+! { dg-do run }
+
+! PR fortran/120505
+
+! Check that a nested allocatable DT component is mapped properly even when the
+! first component is *not* mapped.
+
+module m
+
+ type field_type
+ real(kind=8), allocatable :: density0(:,:), density1(:,:)
+ end type field_type
+
+ type tile_type
+ type(field_type) :: field
+ end type tile_type
+
+ type chunk_type
+ real(kind=8), allocatable :: left_rcv_buffer(:)
+ type(tile_type), allocatable :: tiles(:)
+ end type chunk_type
+
+ type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp map(to: chunk%tiles(1)%field%density1) &
+!$omp map(to: chunk%left_rcv_buffer)
+
+!$omp target
+ if (any (chunk%tiles(1)%field%density1 /= reshape([1,2,3,4],[2,2]))) stop 1
+ chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 + 5
+ chunk%left_rcv_buffer(1) = 42.0_8
+!$omp end target
+
+!$omp target exit data &
+!$omp map(from: chunk%tiles(1)%field%density1) &
+!$omp map(from: chunk%left_rcv_buffer)
+
+if (any (chunk%tiles(1)%field%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 1
+if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 1
+
+end
--
2.51.0
