return fin_dep == GFC_DEP_OVERLAP;
}
+
+/* Check if two refs are equal, for the purposes of checking if one might be
+ the base of the other for OpenMP (target directives). Derived from
+ gfc_dep_resolver. This function is stricter, e.g. indices arr(i) and
+ arr(j) compare as non-equal. */
+
+bool
+gfc_omp_expr_prefix_same (gfc_expr *lexpr, gfc_expr *rexpr)
+{
+ gfc_ref *lref, *rref;
+
+ if (lexpr->symtree && rexpr->symtree)
+ {
+ /* See are_identical_variables above. */
+ if (lexpr->symtree->n.sym->attr.dummy
+ && rexpr->symtree->n.sym->attr.dummy)
+ {
+ /* Dummy arguments: Only check for equal names. */
+ if (lexpr->symtree->n.sym->name != rexpr->symtree->n.sym->name)
+ return false;
+ }
+ else
+ {
+ if (lexpr->symtree->n.sym != rexpr->symtree->n.sym)
+ return false;
+ }
+ }
+ else if (lexpr->base_expr && rexpr->base_expr)
+ {
+ if (gfc_dep_compare_expr (lexpr->base_expr, rexpr->base_expr) != 0)
+ return false;
+ }
+ else
+ return false;
+
+ lref = lexpr->ref;
+ rref = rexpr->ref;
+
+ while (lref && rref)
+ {
+ gfc_dependency fin_dep = GFC_DEP_EQUAL;
+
+ if (lref && lref->type == REF_COMPONENT && lref->u.c.component
+ && strcmp (lref->u.c.component->name, "_data") == 0)
+ lref = lref->next;
+
+ if (rref && rref->type == REF_COMPONENT && rref->u.c.component
+ && strcmp (rref->u.c.component->name, "_data") == 0)
+ rref = rref->next;
+
+ gcc_assert (lref->type == rref->type);
+
+ switch (lref->type)
+ {
+ case REF_COMPONENT:
+ if (lref->u.c.component != rref->u.c.component)
+ return false;
+ break;
+
+ case REF_ARRAY:
+ if (ref_same_as_full_array (lref, rref))
+ break;
+ if (ref_same_as_full_array (rref, lref))
+ break;
+
+ if (lref->u.ar.dimen != rref->u.ar.dimen)
+ {
+ if (lref->u.ar.type == AR_FULL
+ && gfc_full_array_ref_p (rref, NULL))
+ break;
+ if (rref->u.ar.type == AR_FULL
+ && gfc_full_array_ref_p (lref, NULL))
+ break;
+ return false;
+ }
+
+ for (int n = 0; n < lref->u.ar.dimen; n++)
+ {
+ if (lref->u.ar.dimen_type[n] == DIMEN_VECTOR
+ && rref->u.ar.dimen_type[n] == DIMEN_VECTOR
+ && gfc_dep_compare_expr (lref->u.ar.start[n],
+ rref->u.ar.start[n]) == 0)
+ continue;
+ if (lref->u.ar.dimen_type[n] == DIMEN_RANGE
+ && rref->u.ar.dimen_type[n] == DIMEN_RANGE)
+ fin_dep = check_section_vs_section (&lref->u.ar, &rref->u.ar,
+ n);
+ else if (lref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+ && rref->u.ar.dimen_type[n] == DIMEN_RANGE)
+ fin_dep = gfc_check_element_vs_section (lref, rref, n);
+ else if (rref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+ && lref->u.ar.dimen_type[n] == DIMEN_RANGE)
+ fin_dep = gfc_check_element_vs_section (rref, lref, n);
+ else if (lref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+ && rref->u.ar.dimen_type[n] == DIMEN_ELEMENT)
+ {
+ gfc_array_ref l_ar = lref->u.ar;
+ gfc_array_ref r_ar = rref->u.ar;
+ gfc_expr *l_start = l_ar.start[n];
+ gfc_expr *r_start = r_ar.start[n];
+ int i = gfc_dep_compare_expr (r_start, l_start);
+ if (i == 0)
+ fin_dep = GFC_DEP_EQUAL;
+ else
+ return false;
+ }
+ else
+ return false;
+ if (n + 1 < lref->u.ar.dimen
+ && fin_dep != GFC_DEP_EQUAL)
+ return false;
+ }
+
+ if (fin_dep != GFC_DEP_EQUAL
+ && fin_dep != GFC_DEP_OVERLAP)
+ return false;
+
+ break;
+
+ default:
+ gcc_unreachable ();
+ }
+ lref = lref->next;
+ rref = rref->next;
+ }
+
+ return true;
+}
bool gfc_dep_resolver (gfc_ref *, gfc_ref *, gfc_reverse *,
bool identical = false);
bool gfc_are_equivalenced_arrays (gfc_expr *, gfc_expr *);
+bool gfc_omp_expr_prefix_same (gfc_expr *, gfc_expr *);
gfc_expr * gfc_discard_nops (gfc_expr *);
gfc_namespace *ns;
gfc_expr *allocator;
struct gfc_symbol *traits_sym;
+ struct gfc_omp_namelist *duplicate_of;
} u2;
struct gfc_omp_namelist *next;
locus where;
#include "omp-general.h"
#include "omp-low.h"
#include "memmodel.h" /* For MEMMODEL_ enums. */
+#include "dependency.h"
#undef GCC_DIAG_STYLE
#define GCC_DIAG_STYLE __gcc_tdiag__
}
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
{
- tree desc_node;
tree type = TREE_TYPE (decl);
ptr2 = gfc_conv_descriptor_data_get (decl);
- desc_node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
- OMP_CLAUSE_DECL (desc_node) = decl;
- OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
- if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
- {
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_DELETE);
- node2 = desc_node;
- }
- else if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
- || op == EXEC_OMP_TARGET_EXIT_DATA)
+ node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (node2) = decl;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE
+ || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
+ || op == EXEC_OMP_TARGET_EXIT_DATA
+ || op == EXEC_OACC_EXIT_DATA)
{
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_RELEASE);
- node2 = desc_node;
- }
- else if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
- {
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO);
- node2 = node;
- node = desc_node; /* Needs to come first. */
+ gomp_map_kind map_kind
+ = OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE ? GOMP_MAP_DELETE
+ : GOMP_MAP_RELEASE;
+ OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
+ OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) = 1;
}
else
- {
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
- node2 = desc_node;
- }
- if (op == EXEC_OMP_TARGET_EXIT_DATA)
- return;
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
return list;
}
+/* To alleviate quadratic behaviour in checking each entry of a
+ gfc_omp_namelist against every other entry, we build a hashtable indexed by
+ gfc_symbol pointer, which we can use in the usual case that a map
+ expression has a symbol as its root term. Return a namelist based on the
+ root symbol used by N, building a new table in SYM_ROOTED_NL using the
+ gfc_omp_namelist N2 (all clauses) if we haven't done so already. */
+
+static gfc_omp_namelist *
+get_symbol_rooted_namelist (hash_map<gfc_symbol *,
+ gfc_omp_namelist *> *&sym_rooted_nl,
+ gfc_omp_namelist *n,
+ gfc_omp_namelist *n2, bool *sym_based)
+{
+ /* Early-out if we have a NULL clause list (e.g. for OpenACC). */
+ if (!n2)
+ return NULL;
+
+ gfc_symbol *use_sym = NULL;
+
+ /* We're only interested in cases where we have an expression, e.g. a
+ component access. */
+ if (n->expr && n->expr->expr_type == EXPR_VARIABLE && n->expr->symtree)
+ use_sym = n->expr->symtree->n.sym;
+
+ *sym_based = false;
+
+ if (!use_sym)
+ return n2;
+
+ if (!sym_rooted_nl)
+ {
+ sym_rooted_nl = new hash_map<gfc_symbol *, gfc_omp_namelist *> ();
+
+ for (; n2 != NULL; n2 = n2->next)
+ {
+ if (!n2->expr
+ || n2->expr->expr_type != EXPR_VARIABLE
+ || !n2->expr->symtree)
+ continue;
+
+ gfc_omp_namelist *nl_copy = gfc_get_omp_namelist ();
+ memcpy (nl_copy, n2, sizeof *nl_copy);
+ nl_copy->u2.duplicate_of = n2;
+ nl_copy->next = NULL;
+
+ gfc_symbol *idx_sym = n2->expr->symtree->n.sym;
+
+ bool existed;
+ gfc_omp_namelist *&entry
+ = sym_rooted_nl->get_or_insert (idx_sym, &existed);
+ if (existed)
+ nl_copy->next = entry;
+ entry = nl_copy;
+ }
+ }
+
+ gfc_omp_namelist **n2_sym = sym_rooted_nl->get (use_sym);
+
+ if (n2_sym)
+ {
+ *sym_based = true;
+ return *n2_sym;
+ }
+
+ return NULL;
+}
+
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false,
if (clauses == NULL)
return NULL_TREE;
+ hash_map<gfc_symbol *, gfc_omp_namelist *> *sym_rooted_nl = NULL;
+
for (list = 0; list < OMP_LIST_NUM; list++)
{
gfc_omp_namelist *n = clauses->lists[list];
goto finalize_map_clause;
}
+ gfc_omp_namelist *n2
+ = openacc ? NULL : clauses->lists[OMP_LIST_MAP];
+
+ bool sym_based;
+ n2 = get_symbol_rooted_namelist (sym_rooted_nl, n,
+ n2, &sym_based);
+
+ /* If the last reference is a pointer to a derived
+ type ("foo%dt_ptr"), check if any subcomponents
+ of the same derived type member are being mapped
+ elsewhere in the clause list ("foo%dt_ptr%x",
+ etc.). If we have such subcomponent mappings,
+ we only create an ALLOC node for the pointer
+ itself, and inhibit mapping the whole derived
+ type. */
+
+ 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;
+
+ while (ref1->next && ref2->next)
+ {
+ ref1 = ref1->next;
+ ref2 = ref2->next;
+ }
+
+ if (ref2->next)
+ {
+ inner = build_fold_addr_expr (inner);
+ OMP_CLAUSE_SET_MAP_KIND (node,
+ GOMP_MAP_ALLOC);
+ OMP_CLAUSE_DECL (node) = inner;
+ OMP_CLAUSE_SIZE (node)
+ = TYPE_SIZE_UNIT (TREE_TYPE (inner));
+ goto finalize_map_clause;
+ }
+ }
+
tree data, size;
if (lastref->u.c.component->ts.type == BT_CLASS)
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
{
gomp_map_kind map_kind;
- tree desc_node;
tree type = TREE_TYPE (inner);
tree ptr = gfc_conv_descriptor_data_get (inner);
ptr = build_fold_indirect_ref (ptr);
else if (n->u.map_op == OMP_MAP_RELEASE
|| n->u.map_op == OMP_MAP_DELETE)
;
- else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+ else if (op == EXEC_OMP_TARGET_EXIT_DATA
+ || op == EXEC_OACC_EXIT_DATA)
map_kind = GOMP_MAP_RELEASE;
else
map_kind = GOMP_MAP_ALLOC;
OMP_CLAUSE_SIZE (node)
= fold_build2 (MULT_EXPR, gfc_array_index_type,
OMP_CLAUSE_SIZE (node), elemsz);
- desc_node = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- if (openacc)
- OMP_CLAUSE_SET_MAP_KIND (desc_node,
- 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;
+ node2 = build_omp_clause (input_location,
+ 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)
{
- node2 = node;
- node = desc_node; /* Put first. */
+ 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;
}
- if (op == EXEC_OMP_TARGET_EXIT_DATA)
- goto finalize_map_clause;
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3,
}
}
+ /* Free hashmap if we built it. */
+ if (sym_rooted_nl)
+ {
+ typedef hash_map<gfc_symbol *, gfc_omp_namelist *>::iterator hti;
+ for (hti it = sym_rooted_nl->begin (); it != sym_rooted_nl->end (); ++it)
+ {
+ gfc_omp_namelist *&nl = (*it).second;
+ while (nl)
+ {
+ gfc_omp_namelist *next = nl->next;
+ free (nl);
+ nl = next;
+ }
+ }
+ delete sym_rooted_nl;
+ }
+
if (clauses->if_expr)
{
tree if_var;
gfc_start_block (&block);
oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
- code->loc, false, true);
+ code->loc, false, true, code->op);
stmt = build1_loc (input_location, construct_code, void_type_node,
oacc_clauses);
gfc_add_expr_to_block (&block, stmt);
return 1;
}
+/* True if mapping node C maps, or unmaps, a (Fortran) array descriptor. */
+
+static bool
+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;
+
+ return false;
+}
+
/* For a set of mappings describing an array section pointed to by a struct
(or derived type, etc.) component, create an "alloc" or "release" node to
insert into a list following a GOMP_MAP_STRUCT node. For some types of
if (OMP_CLAUSE_CHAIN (grp_start) != grp_end)
grp_mid = OMP_CLAUSE_CHAIN (grp_start);
- if (grp_mid
- && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET)
+ if (grp_mid && omp_map_clause_descriptor_p (grp_mid))
OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (grp_mid);
else
OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
return NULL_TREE;
node = OMP_CLAUSE_CHAIN (node);
- if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+ if (node && omp_map_clause_descriptor_p (node))
{
gcc_assert (node != grp->grp_end);
node = OMP_CLAUSE_CHAIN (node);
== GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
- || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET))
+ || omp_map_clause_descriptor_p (nc)))
{
tree nc2 = OMP_CLAUSE_CHAIN (nc);
if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH)
return node;
node = OMP_CLAUSE_CHAIN (node);
- if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+ if (!node)
+ internal_error ("unexpected mapping node");
+ if (omp_map_clause_descriptor_p (node))
{
if (node == grp->grp_end)
return *grp->grp_start;
node = OMP_CLAUSE_CHAIN (node);
}
- if (node)
- switch (OMP_CLAUSE_MAP_KIND (node))
- {
- case GOMP_MAP_POINTER:
- case GOMP_MAP_FIRSTPRIVATE_POINTER:
- case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
- case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
- *firstprivate = OMP_CLAUSE_DECL (node);
- return *grp->grp_start;
+ switch (OMP_CLAUSE_MAP_KIND (node))
+ {
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ *firstprivate = OMP_CLAUSE_DECL (node);
+ return *grp->grp_start;
- case GOMP_MAP_ALWAYS_POINTER:
- case GOMP_MAP_ATTACH_DETACH:
- case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
- case GOMP_MAP_DETACH:
- return *grp->grp_start;
+ case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ case GOMP_MAP_DETACH:
+ return *grp->grp_start;
- default:
- internal_error ("unexpected mapping node");
- }
- else
- internal_error ("unexpected mapping node");
+ default:
+ internal_error ("unexpected mapping node");
+ }
return error_mark_node;
case GOMP_MAP_TO_PSET:
static omp_mapping_group *
omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
hash_map<tree_operand_hash_no_se, omp_mapping_group *>
- *grpmap)
+ *grpmap,
+ bool enter_exit_data)
{
omp_mapping_group *grp, *outlist = NULL, **cursor;
unsigned int i;
+ bool saw_runtime_implicit = false;
cursor = &outlist;
FOR_EACH_VEC_ELT (*groups, i, grp)
{
if (grp->mark != PERMANENT)
- if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
- return NULL;
+ {
+ if (OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start))
+ {
+ saw_runtime_implicit = true;
+ continue;
+ }
+ if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+ return NULL;
+ }
+ }
+
+ if (!saw_runtime_implicit)
+ return outlist;
+
+ FOR_EACH_VEC_ELT (*groups, i, grp)
+ {
+ if (grp->mark != PERMANENT
+ && OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start))
+ {
+ /* Clear the flag for enter/exit data because it is currently
+ meaningless for those operations in libgomp. */
+ if (enter_exit_data)
+ OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start) = 0;
+
+ if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+ return NULL;
+ }
}
return outlist;
mapping. However, if we have a reference to pointer, make other appropriate
adjustments to the mapping nodes instead.
+ If we have an ATTACH_DETACH node with a Fortran pointer-set (array
+ descriptor) mapping for a derived-type component, and we're also mapping the
+ whole of the derived-type variable on another clause, the pointer-set
+ mapping is removed.
+
If we have a component access but we're also mapping the whole of the
containing struct, drop the former access.
GOMP_MAP_ATTACH_ZLAS for it. */
if (!base_mapped_to && referenced_ptr_node)
OMP_CLAUSE_SET_MAP_KIND (referenced_ptr_node, zlas_kind);
+
+ omp_mapping_group *struct_group;
+ tree desc;
+ if ((desc = OMP_CLAUSE_CHAIN (*grp->grp_start))
+ && omp_map_clause_descriptor_p (desc)
+ && omp_mapped_by_containing_struct (grpmap, decl,
+ &struct_group))
+ /* If we have a pointer set but we're mapping (or unmapping)
+ the whole of the containing struct, we can remove the
+ pointer set mapping. */
+ OMP_CLAUSE_CHAIN (*grp->grp_start) = OMP_CLAUSE_CHAIN (desc);
}
else if (TREE_CODE (TREE_TYPE (base_ptr)) == REFERENCE_TYPE
&& (TREE_CODE (TREE_TYPE (TREE_TYPE (base_ptr)))
for the purposes of gathering sibling lists, etc. */
/* gcc_assert (base == addr_tokens[base_token]->expr); */
- bool ptr = (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_ALWAYS_POINTER);
bool attach_detach = ((OMP_CLAUSE_MAP_KIND (grp_end)
== GOMP_MAP_ATTACH_DETACH)
|| (OMP_CLAUSE_MAP_KIND (grp_end)
== GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION));
+ bool has_descriptor = false;
+ if (OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
+ {
+ tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p);
+ if (grp_mid && omp_map_clause_descriptor_p (grp_mid))
+ has_descriptor = true;
+ }
if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
{
GOMP_MAP_STRUCT into the middle of the old one. */
tree *insert_node_pos = reprocessing_struct ? *added_tail : grp_start_p;
- if (ptr || attach_detach)
+ if (has_descriptor)
+ {
+ tree desc = OMP_CLAUSE_CHAIN (*grp_start_p);
+ if (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)
+ OMP_CLAUSE_SET_MAP_KIND (desc, GOMP_MAP_RELEASE);
+ tree sc = *insert_node_pos;
+ OMP_CLAUSE_CHAIN (l) = desc;
+ OMP_CLAUSE_CHAIN (*grp_start_p) = OMP_CLAUSE_CHAIN (desc);
+ OMP_CLAUSE_CHAIN (desc) = sc;
+ *insert_node_pos = l;
+ }
+ else if (attach_detach)
{
tree extra_node;
tree alloc_node
|| OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_ATTACH_DETACH)
sc = &OMP_CLAUSE_CHAIN (*sc);
for (i = 0; i < elems; i++, sc = &OMP_CLAUSE_CHAIN (*sc))
- if ((ptr || attach_detach) && sc == grp_start_p)
+ if (attach_detach && sc == grp_start_p)
break;
else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
&& TREE_CODE (OMP_CLAUSE_DECL (*sc)) != INDIRECT_REF
|| (known_eq (coffset, offset)
&& maybe_lt (cbitpos, bitpos)))
{
- if (ptr || attach_detach)
+ if (attach_detach)
scp = sc;
else
break;
the list manipulation below. We only need to handle the (pointer
or reference) attach/detach case. */
tree extra_node, alloc_node;
- if (attach_detach)
+ if (has_descriptor)
+ gcc_unreachable ();
+ else if (attach_detach)
alloc_node = build_omp_struct_comp_nodes (code, *grp_start_p,
grp_end, &extra_node);
else
return NULL;
}
- if (ptr || attach_detach)
+ if (has_descriptor)
+ {
+ tree desc = OMP_CLAUSE_CHAIN (*grp_start_p);
+ if (code == OMP_TARGET_EXIT_DATA
+ || code == OACC_EXIT_DATA)
+ OMP_CLAUSE_SET_MAP_KIND (desc, GOMP_MAP_RELEASE);
+ omp_siblist_move_node_after (desc,
+ &OMP_CLAUSE_CHAIN (*grp_start_p),
+ scp ? scp : sc);
+ }
+ else if (attach_detach)
{
tree cl = NULL_TREE, extra_node;
tree alloc_node = build_omp_struct_comp_nodes (code, *grp_start_p,
as a struct (the GOMP_MAP_POINTER following will have the form
"var.data", but such mappings are handled specially). */
tree grpmid = OMP_CLAUSE_CHAIN (*grp_start_p);
- if (OMP_CLAUSE_CODE (grpmid) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (grpmid) == GOMP_MAP_TO_PSET
+ if (omp_map_clause_descriptor_p (grpmid)
&& DECL_P (OMP_CLAUSE_DECL (grpmid)))
continue;
}
list_p);
omp_mapping_group *outlist = NULL;
+ bool enter_exit = (code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA);
/* Topological sorting may fail if we have duplicate nodes, which
we should have detected and shown an error for already. Skip
groups = omp_gather_mapping_groups (list_p);
grpmap = omp_index_mapping_groups (groups);
- outlist = omp_tsort_mapping_groups (groups, grpmap);
+ outlist = omp_tsort_mapping_groups (groups, grpmap, enter_exit);
outlist = omp_segregate_mapping_groups (outlist);
list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
--- /dev/null
+! { dg-additional-options "-fdump-tree-original" }
+
+type t
+integer, pointer :: arr(:)
+end type t
+
+type(t) :: var
+
+allocate (var%arr(1:100))
+
+!$acc enter data copyin(var%arr(10:20))
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc enter data map\(to:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\);$} 1 "original" } }
+
+!$acc exit data delete(var%arr(10:20))
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\);$} 1 "original" } }
+
+
+!$acc enter data create(var%arr(20:30))
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc enter data map\(alloc:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\);$} 1 "original" } }
+
+!$acc exit data finalize delete(var%arr(20:30))
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\) finalize;$} 1 "original" } }
+
+
+!$acc enter data copyin(var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc enter data map\(to:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: 0\]\);$} 1 "original" } }
+
+!$acc exit data delete(var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: 0\]\);$} 1 "original" } }
+
+
+!$acc enter data create(var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc enter data map\(alloc:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: 0\]\);$} 1 "original" } }
+
+!$acc exit data finalize delete(var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma acc exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:var\.arr\.data \[bias: 0\]\) finalize;$} 1 "original" } }
+
+end
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(release:del_f_p \\\[len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(delete:del_f_p \\\[len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(release:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(delete:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_r)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(release:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(delete:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
END SUBROUTINE f
! PR fortran/108545
-! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(always,to:x\.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(attach:x\.a\.data \\\[bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x\.a \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(attach:x\.a\.data \\\[bias: 0\\\]\\)" "omplower" } }
program p
type t
--- /dev/null
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+type T
+integer, pointer :: arr1(:)
+integer, pointer :: arr2(:)
+integer, pointer :: arr3(:)
+integer, pointer :: arr4(:)
+end type T
+
+type(T) :: tv
+integer, allocatable, target, dimension(:) :: arr
+
+allocate(arr(1:20))
+
+tv%arr1 => arr
+tv%arr2 => arr
+tv%arr3 => arr
+tv%arr4 => arr
+
+!$omp target enter data map(to: tv%arr1)
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr1 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr1)
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(release:tv\.arr1 \[pointer set, len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv%arr2) map(to: tv%arr2(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr2 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr2) map(from: tv%arr2(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(release:tv\.arr2 \[pointer set, len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv, tv%arr3(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(to:tv \[len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr3\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv, tv%arr3(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(from:tv \[len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)[_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr3\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv%arr4(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr4 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr4\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr4(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(release:tv\.arr4 \[pointer set, len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr4\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+end
+
--- /dev/null
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+type T
+integer, pointer :: arr1(:)
+integer, pointer :: arr2(:)
+end type T
+
+type(T) :: tv
+integer, allocatable, target, dimension(:) :: arr
+
+allocate(arr(1:20))
+
+tv%arr1 => arr
+tv%arr2 => arr
+
+!$omp target map(tv%arr1)
+tv%arr1(1) = tv%arr1(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target.* map\(struct:tv \[len: 1\]\) map\(to:tv\.arr1 \[pointer set, len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\[implicit\]\) map\(attach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+!$omp target map(tv%arr2) map(tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+!$omp target map(tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target.* map\(struct:tv \[len: 1\]\) map\(to:tv\.arr2 \[pointer set, len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} 2 "gimple" } }
+
+!$omp target map(tv, tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target.* map\(tofrom:tv \[len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+end
+
--- /dev/null
+! { dg-additional-options "-fdump-tree-original" }
+
+type t
+integer, pointer :: arr(:)
+end type t
+
+type(t) :: var
+
+allocate (var%arr(1:100))
+
+!$omp target enter data map(to: var%arr(10:20))
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target enter data map\(to:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\)$} 1 "original" } }
+
+!$omp target exit data map(release: var%arr(10:20))
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\)$} 1 "original" } }
+
+
+!$omp target enter data map(alloc: var%arr(20:30))
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target enter data map\(alloc:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\)$} 1 "original" } }
+
+!$omp target exit data map(delete: var%arr(20:30))
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target exit data map\(delete:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) parm\.[0-9]+\.data \[len: D\.[0-9]+ \* [0-9]+\]\) map\(delete:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: \(integer\(kind=[0-9]+\)\) parm\.[0-9]+\.data - \(integer\(kind=[0-9]+\)\) var\.arr\.data\]\)$} 1 "original" } }
+
+
+!$omp target enter data map(to: var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target enter data map\(to:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\[implicit\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: 0\]\)$} 1 "original" } }
+
+!$omp target exit data map(release: var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target exit data map\(release:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\[implicit\]\) map\(release:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: 0\]\)$} 1 "original" } }
+
+
+!$omp target enter data map(alloc: var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target enter data map\(alloc:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\[implicit\]\) map\(to:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: 0\]\)$} 1 "original" } }
+
+!$omp target exit data map(delete: var%arr)
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target exit data map\(delete:\*\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[len: D\.[0-9]+ \* [0-9]+\]\[implicit\]\) map\(delete:var\.arr \[pointer set, len: [0-9]+\]\) map\(attach_detach:\(integer\(kind=[0-9]+\)\[0:\] \*\) var\.arr\.data \[bias: 0\]\)$} 1 "original" } }
+
+
+end
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, " [bias: ");
break;
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
+ if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_RELEASE_DESCRIPTOR (clause))
+ {
+ pp_string (pp, " [pointer set, len: ");
+ break;
+ }
+ /* Fallthrough. */
default:
pp_string (pp, " [len: ");
break;
same directive. */
#define OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED(NODE) \
TREE_STATIC (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* Nonzero if this is a release/delete node which refers to a (Fortran) array
+ descriptor. */
+#define OMP_CLAUSE_RELEASE_DESCRIPTOR(NODE) \
+ TREE_NOTHROW (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
/* Flag that 'OMP_CLAUSE_DECL (NODE)' is to be made addressable during OMP
lowering. */
--- /dev/null
+! { dg-do run }
+
+program myprog
+type u
+ integer, dimension (:), pointer :: tarr1
+ integer, dimension (:), pointer :: tarr2
+ integer, dimension (:), pointer :: tarr3
+end type u
+
+type(u) :: myu1, myu2, myu3
+
+integer, dimension (12), target :: myarray1
+integer, dimension (12), target :: myarray2
+integer, dimension (12), target :: myarray3
+integer, dimension (12), target :: myarray4
+integer, dimension (12), target :: myarray5
+integer, dimension (12), target :: myarray6
+integer, dimension (12), target :: myarray7
+integer, dimension (12), target :: myarray8
+integer, dimension (12), target :: myarray9
+
+myu1%tarr1 => myarray1
+myu1%tarr2 => myarray2
+myu1%tarr3 => myarray3
+myu2%tarr1 => myarray4
+myu2%tarr2 => myarray5
+myu2%tarr3 => myarray6
+myu3%tarr1 => myarray7
+myu3%tarr2 => myarray8
+myu3%tarr3 => myarray9
+
+myu1%tarr1 = 0
+myu1%tarr2 = 0
+myu1%tarr3 = 0
+myu2%tarr1 = 0
+myu2%tarr2 = 0
+myu2%tarr3 = 0
+myu3%tarr1 = 0
+myu3%tarr2 = 0
+myu3%tarr3 = 0
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(:)) &
+!$omp& map(to:myu1%tarr2) map(tofrom:myu1%tarr2(:)) &
+!$omp& map(to:myu1%tarr3) map(tofrom:myu1%tarr3(:)) &
+!$omp& map(to:myu2%tarr1) map(tofrom:myu2%tarr1(:)) &
+!$omp& map(to:myu2%tarr2) map(tofrom:myu2%tarr2(:)) &
+!$omp& map(to:myu2%tarr3) map(tofrom:myu2%tarr3(:)) &
+!$omp& map(to:myu3%tarr1) map(tofrom:myu3%tarr1(:)) &
+!$omp& map(to:myu3%tarr2) map(tofrom:myu3%tarr2(:)) &
+!$omp& map(to:myu3%tarr3) map(tofrom:myu3%tarr3(:))
+myu1%tarr1(1) = myu1%tarr1(1) + 1
+myu2%tarr1(1) = myu2%tarr1(1) + 1
+myu3%tarr1(1) = myu3%tarr1(1) + 1
+!$omp end target
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(1:2)) &
+!$omp& map(to:myu1%tarr2) map(tofrom:myu1%tarr2(1:2)) &
+!$omp& map(to:myu1%tarr3) map(tofrom:myu1%tarr3(1:2)) &
+!$omp& map(to:myu2%tarr1) map(tofrom:myu2%tarr1(1:2)) &
+!$omp& map(to:myu2%tarr2) map(tofrom:myu2%tarr2(1:2)) &
+!$omp& map(to:myu2%tarr3) map(tofrom:myu2%tarr3(1:2)) &
+!$omp& map(to:myu3%tarr1) map(tofrom:myu3%tarr1(1:2)) &
+!$omp& map(to:myu3%tarr2) map(tofrom:myu3%tarr2(1:2)) &
+!$omp& map(to:myu3%tarr3) map(tofrom:myu3%tarr3(1:2))
+myu1%tarr2(1) = myu1%tarr2(1) + 1
+myu2%tarr2(1) = myu2%tarr2(1) + 1
+myu3%tarr2(1) = myu3%tarr2(1) + 1
+!$omp end target
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(1)) &
+!$omp& map(to:myu1%tarr2) map(tofrom:myu1%tarr2(1)) &
+!$omp& map(to:myu1%tarr3) map(tofrom:myu1%tarr3(1)) &
+!$omp& map(to:myu2%tarr1) map(tofrom:myu2%tarr1(1)) &
+!$omp& map(to:myu2%tarr2) map(tofrom:myu2%tarr2(1)) &
+!$omp& map(to:myu2%tarr3) map(tofrom:myu2%tarr3(1)) &
+!$omp& map(to:myu3%tarr1) map(tofrom:myu3%tarr1(1)) &
+!$omp& map(to:myu3%tarr2) map(tofrom:myu3%tarr2(1)) &
+!$omp& map(to:myu3%tarr3) map(tofrom:myu3%tarr3(1))
+myu1%tarr3(1) = myu1%tarr3(1) + 1
+myu2%tarr3(1) = myu2%tarr3(1) + 1
+myu3%tarr3(1) = myu3%tarr3(1) + 1
+!$omp end target
+
+!$omp target map(tofrom:myu1%tarr1) &
+!$omp& map(tofrom:myu1%tarr2) &
+!$omp& map(tofrom:myu1%tarr3) &
+!$omp& map(tofrom:myu2%tarr1) &
+!$omp& map(tofrom:myu2%tarr2) &
+!$omp& map(tofrom:myu2%tarr3) &
+!$omp& map(tofrom:myu3%tarr1) &
+!$omp& map(tofrom:myu3%tarr2) &
+!$omp& map(tofrom:myu3%tarr3)
+myu1%tarr2(1) = myu1%tarr2(1) + 1
+myu2%tarr2(1) = myu2%tarr2(1) + 1
+myu3%tarr2(1) = myu3%tarr2(1) + 1
+!$omp end target
+
+if (myu1%tarr1(1).ne.1) stop 1
+if (myu2%tarr1(1).ne.1) stop 2
+if (myu3%tarr1(1).ne.1) stop 3
+if (myu1%tarr2(1).ne.2) stop 4
+if (myu2%tarr2(1).ne.2) stop 5
+if (myu3%tarr2(1).ne.2) stop 6
+if (myu1%tarr3(1).ne.1) stop 7
+if (myu2%tarr3(1).ne.1) stop 8
+if (myu3%tarr3(1).ne.1) stop 9
+
+end program myprog
--- /dev/null
+! { dg-do run }
+
+module mymod
+type G
+integer :: x, y
+integer, pointer :: arr(:)
+integer :: z
+end type G
+end module mymod
+
+program myprog
+use mymod
+
+integer, target :: arr1(10)
+integer, target :: arr2(10)
+integer, target :: arr3(10)
+type(G), dimension(3) :: gvar
+
+integer :: i, j
+
+gvar(1)%arr => arr1
+gvar(2)%arr => arr2
+gvar(3)%arr => arr3
+
+gvar(1)%arr = 0
+gvar(2)%arr = 0
+gvar(3)%arr = 0
+
+i = 1
+j = 1
+
+! Here 'gvar(i)' and 'gvar(j)' are the same element, so this should work.
+! This generates a whole-array mapping for gvar(i)%arr, but with the
+! "runtime implicit" bit set so the smaller subarray gvar(j)%arr(1:5) takes
+! precedence.
+
+!$omp target map(gvar(i)%arr, gvar(j)%arr(1:5))
+gvar(i)%arr(1) = gvar(i)%arr(1) + 1
+gvar(j)%arr(1) = gvar(j)%arr(1) + 2
+!$omp end target
+
+!$omp target map(gvar(i)%arr(1:5), gvar(j)%arr)
+gvar(i)%arr(1) = gvar(i)%arr(1) + 3
+gvar(j)%arr(1) = gvar(j)%arr(1) + 4
+!$omp end target
+
+! For these ones, we know the array index is the same, so we can just
+! drop the whole-array mapping.
+
+!$omp target map(gvar(i)%arr, gvar(i)%arr(1:5))
+gvar(i)%arr(1) = gvar(i)%arr(1) + 1
+gvar(i)%arr(1) = gvar(j)%arr(1) + 2
+!$omp end target
+
+!$omp target map(gvar(i)%arr(1:5), gvar(i)%arr)
+gvar(i)%arr(1) = gvar(i)%arr(1) + 3
+gvar(i)%arr(1) = gvar(j)%arr(1) + 4
+!$omp end target
+
+if (gvar(1)%arr(1).ne.20) stop 1
+
+end program myprog
--- /dev/null
+! { dg-do run }
+
+type t
+ integer, pointer :: p(:)
+end type t
+
+type(t) :: var(2)
+
+allocate (var(1)%p, source=[1,2,3,5])
+allocate (var(2)%p, source=[2,3,5])
+
+!$omp target map(var(1)%p, var(2)%p)
+var(1)%p(1) = 5
+var(2)%p(2) = 7
+!$omp end target
+
+!$omp target map(var(1)%p(1:3), var(1)%p, var(2)%p)
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+!$omp target map(var(1)%p, var(2)%p, var(2)%p(1:3))
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+!$omp target map(var(1)%p, var(1)%p(1:3), var(2)%p, var(2)%p(2))
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+if (var(1)%p(1).ne.8) stop 1
+if (var(2)%p(2).ne.10) stop 2
+
+end
--- /dev/null
+! { dg-do run }
+
+type t
+ integer, pointer :: p(:)
+ integer, pointer :: p2(:)
+end type t
+
+type(t) :: var
+integer, target :: tgt(5), tgt2(1000)
+var%p => tgt
+var%p2 => tgt2
+
+p = 0
+p2 = 0
+
+!$omp target map(tgt, tgt2(4:6), var)
+ var%p(1) = 5
+ var%p2(5) = 7
+!$omp end target
+
+if (var%p(1).ne.5) stop 1
+if (var%p2(5).ne.7) stop 2
+
+end
+
+! { dg-shouldfail "" { offload_device_nonshared_as } }
--- /dev/null
+type t
+integer, pointer :: p2(:)
+end type t
+
+integer, target :: A(5)
+integer, pointer :: p(:), p2(:)
+type(t) :: var
+
+allocate(p2(1:20))
+p => A
+var%p2 => p2
+
+A = 0
+p2 = 0
+
+! These arrays "share original storage", so are unsupported. This will
+! (correctly) fail with a non-shared address space.
+
+!$omp target map(A(3:4), p2(4:8), p, var%p2)
+A(3) = A(3) + 1
+p2(4) = p2(4) + 2
+!$omp end target
+
+if (A(3).ne.1) stop 1
+if (p2(4).ne.2) stop 2
+
+end program
+
+! { dg-shouldfail "" { offload_device_nonshared_as } }
--- /dev/null
+! { dg-do run }
+
+type F
+integer, pointer :: mem(:)
+end type F
+
+type(F) :: fv
+integer, allocatable, target :: arr(:)
+
+allocate(arr(1:20))
+
+fv%mem => arr
+fv%mem = 0
+
+!$omp target enter data map(to: fv%mem(1:10))
+!$omp target map(alloc: fv%mem)
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem(1:10))
+
+if (fv%mem(1).ne.1) stop 1
+
+!$omp target enter data map(to: fv, fv%mem(1:10))
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv, fv%mem(1:10))
+
+if (fv%mem(1).ne.2) stop 2
+
+!$omp target enter data map(to: fv%mem, fv%mem(1:10))
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem, fv%mem(1:10))
+
+if (fv%mem(1).ne.3) stop 3
+
+!$omp target enter data map(to: fv%mem)
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem)
+
+if (fv%mem(1).ne.4) stop 4
+
+end
--- /dev/null
+! { dg-do run }
+
+program myprog
+type u
+ integer, dimension (:), pointer :: tarr
+end type u
+
+type(u) :: myu
+integer, dimension (12), target :: myarray
+
+myu%tarr => myarray
+
+myu%tarr = 0
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(:))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(1:2))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(1))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(tofrom:myu%tarr)
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+if (myu%tarr(1).ne.4) stop 1
+
+end program myprog
--- /dev/null
+! { dg-do run }
+
+module mymod
+type F
+integer :: a, b, c
+integer, dimension(10) :: d
+end type F
+
+type G
+integer :: x, y
+type(F), pointer :: myf
+integer :: z
+end type G
+end module mymod
+
+program myprog
+use mymod
+
+type(F), target :: ftmp
+type(G) :: gvar
+
+gvar%myf => ftmp
+
+gvar%myf%d = 0
+
+!$omp target map(to:gvar%myf) map(tofrom: gvar%myf%b, gvar%myf%d)
+gvar%myf%d(1) = gvar%myf%d(1) + 1
+!$omp end target
+
+if (gvar%myf%d(1).ne.1) stop 1
+
+end program myprog
call six ()
call seven ()
call eight ()
+ call nine ()
+ call ten ()
+ call eleven ()
+ call twelve ()
contains
! Implicitly mapped – but no pointers are mapped
!$omp end target
end subroutine eight
-end program main
+ ! This is "subroutine four" but with explicit base-pointer mappings
+ ! (var%f, etc.).
+ subroutine nine()
+ type(t2) :: var
+
+ print '(g0)', '==== TESTCASE "nine" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+ uni1 = 4_"abcde", &
+ uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+ allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+! !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3)) &
+! !$omp& map(tofrom: var%str4(2:2), var%uni2(2:3), var%uni4(2:2))
+ !$omp target map(to: var%f) map(tofrom: var%d(4:7), var%f(2:3), &
+ !$omp& var%str2(2:3), var%uni2(2:3))
+ if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+
+ if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 19
+! if (.not. associated (var%uni4)) stop 20
+! if (len (var%uni4) /= 5) stop 21
+! if (size (var%uni4) /= 2) stop 22
+! if (var%uni4(2) /= "Go!!!") stop 23
+ !$omp end target
+
+ deallocate(var%f, var%str4)
+ end subroutine nine
+
+ ! This is "subroutine five" but with explicit base-pointer mappings.
+ subroutine ten()
+ type(t2) :: var
+
+ print '(g0)', '==== TESTCASE "ten" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+ uni1 = 4_"abcde", &
+ uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+ !$omp target map(tofrom: var%d(4:7))
+ if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+ !$omp end target
+ !$omp target map(tofrom: var%str2(2:3))
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+ !$omp end target
+
+ !$omp target map(to: var%f) map(tofrom: var%f(2:3))
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+ !$omp end target
+! !$omp target map(tofrom: var%str4(2:2))
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+! !$omp end target
+! !$omp target map(tofrom: var%uni4(2:2))
+! if (.not. associated (var%uni4)) stop 15
+! if (len (var%uni4) /= 5) stop 16
+! if (size (var%uni4) /= 2) stop 17
+! if (var%uni4(2) /= 4_"Go!!!") stop 18
+! !$omp end target
+
+ deallocate(var%f, var%str4)
+ end subroutine ten
+
+ ! This is "subroutine six" but with explicit base pointer mappings.
+ subroutine eleven()
+ type(t2) :: var
-! Fixed by the "Fortran pointers and member mappings" patch
-! { dg-xfail-run-if TODO { offload_device_nonshared_as } }
+ print '(g0)', '==== TESTCASE "eleven" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+ uni1 = 4_"abcde", &
+ uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+ allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+! !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), &
+! !$omp var%str4(2), var%uni2(3), var%uni4(2))
+ !$omp target map(to: var%f) map(tofrom: var%d(5), var%f(3), &
+ !$omp& var%str2(3), var%uni2(3))
+ if (var%d(5) /= -3*5) stop 4
+ if (var%str2(3) /= "ABCDE") stop 6
+ if (var%uni2(3) /= 4_"ABCDE") stop 7
+
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (var%f(3) /= 44) stop 11
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+! if (.not. associated (var%uni4)) stop 19
+! if (len (var%uni4) /= 5) stop 20
+! if (size (var%uni4) /= 2) stop 21
+! if (var%uni4(2) /= 4_"Go!!!") stop 22
+ !$omp end target
+
+ deallocate(var%f, var%str4, var%uni4)
+ end subroutine eleven
+
+ ! This is "subroutine seven" but with explicit base-pointer mappings.
+ subroutine twelve()
+ type(t2) :: var
+
+ print '(g0)', '==== TESTCASE "twelve" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+ uni1 = 4_"abcde", &
+ uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+ allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+ !$omp target map(tofrom: var%d(5))
+ if (var%d(5) /= (-3*5)) stop 4
+ !$omp end target
+ !$omp target map(tofrom: var%str2(2:3))
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+ !$omp end target
+ !$omp target map(tofrom: var%uni2(2:3))
+ if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 7
+ !$omp end target
+
+ !$omp target map(to: var%f) map(tofrom: var%f(2:3))
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+ !$omp end target
+! !$omp target map(tofrom: var%str4(2:2))
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+! !$omp end target
+! !$omp target map(tofrom: var%uni4(2:2))
+! if (.not. associated (var%uni4)) stop 15
+! if (len (var%uni4) /= 5) stop 16
+! if (size (var%uni4) /= 2) stop 17
+! if (var%uni4(2) /= 4_"Go!!!") stop 18
+! !$omp end target
+
+ deallocate(var%f, var%str4, var%uni4)
+ end subroutine twelve
+
+end program main
projects / thirdparty / gcc.git / commitdiff
