return NULL;
}
+/* Helper function for gfc_trans_omp_clauses. Adjust existing and create new
+ map nodes for derived-type component array descriptors. Return true if the
+ mapping has to be dropped. */
+
+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)
+ {
+ /* For an intermediate descriptor, the pointee (i.e. the actual array
+ content) is mapped in a separate set of nodes. This ALLOC is only
+ emitted to comply with the group layout expected by the gimplifier.
+ */
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_SIZE (node) = size_zero_node;
+ OMP_CLAUSE_MAP_GIMPLE_ONLY (node) = 1;
+ }
+ 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 if (mid_desc_p)
+ {
+ /* For an intermediate descriptor, the pointee (i.e. the actual array
+ content) is mapped in a separate set of nodes. This ALLOC is only
+ emitted to comply with the group layout expected by the gimplifier. */
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_SIZE (node) = size_zero_node;
+ OMP_CLAUSE_MAP_GIMPLE_ONLY (node) = 1;
+ }
+ 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;
+ }
+
+ /* 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 drop_mapping = false;
+
+ if (!mid_desc_p)
+ {
+ gfc_omp_namelist *n2 = clauses->lists[OMP_LIST_MAP];
+
+ bool sym_based;
+ n2 = get_symbol_rooted_namelist (sym_rooted_nl, n, n2, &sym_based);
+
+ 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_zero_node;
+ if (mid_desc_p)
+ OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (node3) = 1;
+
+ return false;
+}
+
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false,
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;
{
gfc_init_se (&se, NULL);
se.expr = gfc_maybe_dereference_var (n->sym, decl);
+ vec<tree> mid_descr = vNULL;
+ vec<gfc_ref *> midref = vNULL;
for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
{
conv_parent_component_references (&se, ref);
gfc_conv_component_ref (&se, ref);
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (se.expr)))
+ {
+ mid_descr.safe_push (se.expr);
+ midref.safe_push (ref);
+ }
}
else if (ref->type == REF_ARRAY)
{
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;
}
else
gcc_unreachable ();
+
+ /* Map intermediate array descriptors. */
+ if (!openacc && !mid_descr.is_empty ())
+ for (size_t i = 0; i < mid_descr.length (); i++)
+ if (mid_descr[i] != inner
+ && !descriptors.contains (midref[i]->u.c.sym))
+ {
+ descriptors.safe_push (midref[i]->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[i], 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");
}
if (wholestruct)
{
+ /* An intermediate descriptor should not match here because the
+ pointee is actually not mapped by this group -- it is just a
+ zero-length alloc. */
+ 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;
}
omp_mapping_group *wholestruct;
if (omp_mapped_by_containing_struct (*grpmap, OMP_CLAUSE_DECL (c),
&wholestruct))
- continue;
+ {
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (c) = 0;
+ continue;
+ }
if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
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, so that it can
+ be picked up by the adjustment code further down in this function. */
+ bool attach_bias_needs_adjustment;
+ attach_bias_needs_adjustment = false;
+ FOR_EACH_VEC_ELT_REVERSE (*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 (c) == GOMP_MAP_STRUCT_UNORD)
+ && 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_MAP_SIZE_NEEDS_ADJUSTMENT (grp->grp_end))
+ {
+ OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (grp->grp_end) = 0;
+ attach_bias_needs_adjustment = true;
+ 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
an "enter data" operation (because for those, variables need to be
mapped separately and attach nodes must be grouped together with the
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);
+ ATTACH_DETACH node either after GOMP_MAP_STRUCT for a target region
+ or for an intermediate descriptor that needs adjustment -- so this
+ should never be true. */
+ gcc_assert ((region_type & ORT_TARGET) != 0
+ || attach_bias_needs_adjustment);
/* This is the first sorted node in the struct sibling list. Use it
to recalculate the correct bias to use.
break;
case OMP_CLAUSE_MAP:
+ if (OMP_CLAUSE_MAP_GIMPLE_ONLY (c))
+ {
+ remove = true;
+ goto end_adjust_omp_map_clause;
+ }
decl = OMP_CLAUSE_DECL (c);
if (!grp_end)
{
/* Fallthrough. */
case OMP_TARGET_EXIT_DATA:
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DETACH);
+ OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (c) = 0;
break;
case OACC_UPDATE:
/* An "attach/detach" operation on an update directive
break;
default:
change_to_attach:
+ gcc_assert (!OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (c));
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH);
if ((ctx->region_type & ORT_TARGET) != 0)
move_attach = true;
--- /dev/null
+! { 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\] \[gimple only\]\) map\(to:chunk\.tiles \[pointer set, len: (?:36|64)\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: 0 \(needs adjustment\)\]\) } 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: (?:36|64)\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: 0 \(needs adjustment\)\]\) } 1 "original" } }
+
+end
--- /dev/null
+! { 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.
+! Pointer set lengths are checked for both 32 and 64 bits.
+
+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: (?:48|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: (?:36|64)\]\) map\(detach:chunk\.tiles\.data \[bias: [0-9]+\]\)} 1 "gimple" } }
+
+
+! { dg-final { scan-tree-dump-not { map\(alloc } "gimple" } }
+! { dg-final { scan-tree-dump-not { gimple only } "gimple" } }
+! { dg-final { scan-tree-dump-not { needs adjustment } "gimple" } }
+
+end
}
dump_generic_node (pp, OMP_CLAUSE_SIZE (clause),
spc, flags, false);
+ if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (clause))
+ pp_string (pp, " (needs adjustment)");
pp_right_bracket (pp);
}
- if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (clause))
- pp_string (pp, " [runtime_implicit]");
+ if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP)
+ {
+ if (OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (clause))
+ pp_string (pp, " [runtime_implicit]");
+ if (OMP_CLAUSE_MAP_GIMPLE_ONLY (clause))
+ pp_string (pp, " [gimple only]");
+ }
pp_right_paren (pp);
break;
#define OMP_CLAUSE_MAP_READONLY(NODE) \
TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* Nonzero if the size (or bias) is not known by the front end and needs to be
+ adjusted in the middle end. */
+#define OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT(NODE) \
+ TREE_CONSTANT (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+
+/* Nonzero on a map clause that is only used internally by the gimplifier and
+ can thus be removed at the end of the GIMPLE pass. */
+#define OMP_CLAUSE_MAP_GIMPLE_ONLY(NODE) \
+ TREE_USED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+
/* Same as above, for use in OpenACC cache directives. */
#define OMP_CLAUSE__CACHE__READONLY(NODE) \
TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_))
--- /dev/null
+! Test the 'present' modifier with derived-type allocatable components.
+
+module m
+ implicit none
+ 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
+implicit none
+allocate(chunk%tiles(1))
+chunk%tiles(1)%field%density0 = reshape([1,2,3,4],[2,2])
+
+!$omp target enter data &
+!$omp map(to: chunk%tiles(1)%field%density0) &
+!$omp map(to: chunk%tiles(1)%field%density1)
+
+!$omp target map(present, alloc: chunk%tiles(1)%field%density0)
+ if (.not. allocated(chunk%tiles(1)%field%density0)) stop 1
+ if (any (chunk%tiles(1)%field%density0 /= reshape([1,2,3,4],[2,2]))) stop 1
+ chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 * 2
+!$omp end target
+
+chunk%tiles(1)%field%density1 = reshape([11,22,33,44],[2,2])
+
+!$omp target map(alloc: chunk%tiles(1)%field%density0)
+ if (.not. allocated(chunk%tiles(1)%field%density0)) stop 1
+ if (any (chunk%tiles(1)%field%density0 /= 2*reshape([1,2,3,4],[2,2]))) stop 1
+ chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 * 7
+!$omp end target
+
+!$omp target exit data &
+!$omp map(from: chunk%tiles(1)%field%density0)
+
+if (any (chunk%tiles(1)%field%density0 /= 7*2*reshape([1,2,3,4],[2,2]))) stop 1
+if (any (chunk%tiles(1)%field%density1 /= reshape([11,22,33,44],[2,2]))) stop 2
+
+end
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { dg-do run }
+
+! PR fortran/120505
+
+! Test multiple levels of nesting of allocatable components, with some
+! components mapped and some not.
+
+module m
+
+ type field_type
+ real(kind=8), allocatable :: density0(:,:), density1(:,:)
+ end type field_type
+
+ type tile_type
+ type(field_type),allocatable :: field(:)
+ end type tile_type
+
+ type tile_type2
+ type(tile_type),allocatable :: tiles_inner(:)
+ end type tile_type2
+
+ type chunk_type
+ real(kind=8), allocatable :: left_rcv_buffer(:)
+ type(tile_type2), allocatable :: tiles(:)
+ end type chunk_type
+
+ type(chunk_type) :: chunk
+
+end
+
+use m
+
+allocate(chunk%tiles(1))
+allocate(chunk%tiles(1)%tiles_inner(1))
+allocate(chunk%tiles(1)%tiles_inner(1)%field(1))
+chunk%tiles(1)%tiles_inner(1)%field(1)%density1 = reshape([1,2,3,4],[2,2])
+allocate(chunk%left_rcv_buffer(1))
+
+!$omp target enter data &
+!$omp map(to: chunk%tiles(1)%tiles_inner(1)%field(1)%density1) &
+!$omp map(to: chunk%left_rcv_buffer)
+
+!$omp target
+ if (any (chunk%tiles(1)%tiles_inner(1)%field(1)%density1 /= reshape([1,2,3,4],[2,2]))) stop 1
+ chunk%tiles(1)%tiles_inner(1)%field(1)%density1 = chunk%tiles(1)%tiles_inner(1)%field(1)%density1 + 5
+ chunk%left_rcv_buffer(1) = 42.0_8
+!$omp end target
+
+!$omp target exit data &
+!$omp map(from: chunk%tiles(1)%tiles_inner(1)%field(1)%density1) &
+!$omp map(from: chunk%left_rcv_buffer)
+
+if (any (chunk%tiles(1)%tiles_inner(1)%field(1)%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 1
+if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 1
+
+end
--- /dev/null
+! { dg-do run }
+
+! PR fortran/120505
+
+! Test 'target enter data' followed by 'target map' with a nested allocatable
+! component.
+
+module m
+
+ type field_type
+ real(kind=8), allocatable :: density0(:,:), density1(:,:)
+ end type field_type
+
+ type tile_type
+ type(field_type),allocatable :: 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))
+allocate(chunk%tiles(1)%field(1))
+chunk%tiles(1)%field(1)%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(1)) &
+!$omp map(to: chunk%tiles(1)%field(1)%density1) &
+!$omp map(to: chunk%left_rcv_buffer)
+
+!$omp target map(tofrom: chunk%tiles(1)%field(1)%density1)
+ if (any (chunk%tiles(1)%field(1)%density1 /= reshape([1,2,3,4],[2,2]))) stop 1
+ chunk%tiles(1)%field(1)%density1 = chunk%tiles(1)%field(1)%density1 + 5
+ chunk%left_rcv_buffer(1) = 42.0_8
+!$omp end target
+
+!$omp target exit data &
+!$omp map(from: chunk%tiles(1)%field(1)%density1) &
+!$omp map(from: chunk%left_rcv_buffer)
+
+if (any (chunk%tiles(1)%field(1)%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 2
+if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 3
+
+end
--- /dev/null
+! { dg-do run }
+! { dg-shouldfail "PR124178 TODO" }
+
+! PR fortran/120505
+
+! Demonstrate a bogus error which occurs when trying to map a nested allocatable
+! component as well as its containing structure on the same target directive.
+! Also check that no ICE happens in this case.
+
+module m
+
+ type field_type
+ real(kind=8), allocatable :: density0(:,:), density1(:,:)
+ end type field_type
+
+ type tile_type
+ type(field_type),allocatable :: 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))
+allocate(chunk%tiles(1)%field(1))
+chunk%tiles(1)%field(1)%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(1)) &
+!$omp map(to: chunk%tiles(1)%field(1)%density1) &
+!$omp map(to: chunk%left_rcv_buffer)
+
+!$omp target map(tofrom: chunk%tiles(1)%field(1)%density1, chunk)
+! { dg-output {libgomp: Trying to map into device \[0x[0-9a-f]+\.\.0x[0-9a-f]+\) object when \[0x[0-9a-f]+\.\.0x[0-9a-f]+\) is already mapped} }
+ if (any (chunk%tiles(1)%field(1)%density1 /= reshape([1,2,3,4],[2,2]))) stop 1
+ chunk%tiles(1)%field(1)%density1 = chunk%tiles(1)%field(1)%density1 + 5
+ chunk%left_rcv_buffer(1) = 42.0_8
+!$omp end target
+
+!$omp target exit data &
+!$omp map(from: chunk%tiles(1)%field(1)%density1) &
+!$omp map(from: chunk%left_rcv_buffer)
+
+if (any (chunk%tiles(1)%field(1)%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 2
+if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 3
+
+end
projects / thirdparty / gcc.git / commitdiff
