@@ -2334,3 +2334,131 @@ gfc_dep_resolver (gfc_ref *lref, gfc_ref *rref, gfc_reverse *reverse,
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;
+}
@@ -40,5 +40,6 @@ int gfc_expr_is_one (gfc_expr *, int);
int gfc_dep_resolver (gfc_ref *, gfc_ref *, gfc_reverse *,
bool identical = false);
int 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 *);
@@ -1362,6 +1362,7 @@ typedef struct gfc_omp_namelist
{
struct gfc_omp_namelist_udr *udr;
gfc_namespace *ns;
+ struct gfc_omp_namelist *duplicate_of;
} u2;
struct gfc_omp_namelist *next;
locus where;
@@ -40,6 +40,7 @@ along with GCC; see the file COPYING3. If not see
#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__
@@ -2471,28 +2472,15 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
}
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 (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. */
- }
- else
- {
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
- node2 = desc_node;
- }
- node3 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
+ node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_DECL (node2) = decl;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ 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);
+ OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
/* This purposely does not include GOMP_MAP_ALWAYS_POINTER. The extra
cast prevents gimplify.cc from recognising it as being part of the
struct - and adding an 'alloc: for the 'desc.data' pointer, which
@@ -2593,6 +2581,73 @@ handle_iterator (gfc_namespace *ns, stmtblock_t *iter_block, tree block)
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,
@@ -2610,6 +2665,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
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];
@@ -3455,6 +3512,54 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
{
if (pointer || (openacc && allocatable))
{
+ 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)
@@ -3502,7 +3607,6 @@ 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 desc_node;
tree type = TREE_TYPE (inner);
tree ptr = gfc_conv_descriptor_data_get (inner);
ptr = build_fold_indirect_ref (ptr);
@@ -3541,21 +3645,69 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
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;
- else
+ node2 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ 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;
}
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
@@ -3707,6 +3859,23 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
}
}
+ /* 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;
@@ -9651,18 +9651,45 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
static omp_mapping_group *
omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
hash_map<tree_operand_hash, 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;
@@ -10060,6 +10087,11 @@ omp_check_mapping_compatibility (location_t loc,
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.
@@ -10239,6 +10271,18 @@ omp_resolve_clause_dependencies (enum tree_code code,
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_CLAUSE_CODE (desc) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_TO_PSET
+ && omp_mapped_by_containing_struct (grpmap, decl,
+ &struct_group))
+ /* If we have a pointer set but we're mapping 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)))
@@ -10673,11 +10717,19 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
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_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET)
+ has_descriptor = true;
+ }
if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
{
@@ -10700,7 +10752,18 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
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
@@ -10931,7 +10994,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
|| 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
@@ -10987,7 +11050,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
|| (known_eq (coffset, offset)
&& maybe_lt (cbitpos, bitpos)))
{
- if (ptr || attach_detach)
+ if (attach_detach)
scp = sc;
else
break;
@@ -11003,7 +11066,9 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
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
@@ -11036,7 +11101,17 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
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,
@@ -11457,6 +11532,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
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
@@ -11471,7 +11548,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
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);
@@ -11865,27 +11942,45 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
explicit-shape array to a function expecting an assumed-shape
argument. Performing "enter data" inside the called function
would thus be problematic. */
- if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
- OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*prev_list_p)
- == GOMP_MAP_DELETE
- ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
- else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE)
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
- remove = true;
- else if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
- && OMP_CLAUSE_CHAIN (c)
- && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
- == OMP_CLAUSE_MAP)
- && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ATTACH_DETACH)
- || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
- && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
- (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
+
+ tree desc;
+ if (DECL_P (decl)
+ || (prev_list_p
+ && (desc = OMP_CLAUSE_CHAIN (*prev_list_p))
+ && OMP_CLAUSE_CODE (desc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_RELEASE
+ || OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_DELETE)
+ && DECL_P (OMP_CLAUSE_DECL (desc))))
+ {
+ /* We only do the transformations here for non-component
+ mappings. Array descriptors for component mappings are
+ handled in omp_accumulate_sibling_list. */
+ if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
+ {
+ enum gomp_map_kind k
+ = (OMP_CLAUSE_MAP_KIND (*prev_list_p) == GOMP_MAP_DELETE
+ ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
+ OMP_CLAUSE_SET_MAP_KIND (c, k);
+ }
+ else if ((code == OMP_TARGET_EXIT_DATA
+ || code == OMP_TARGET_UPDATE)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
+ remove = true;
+ else if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ && OMP_CLAUSE_CHAIN (c)
+ && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
+ == OMP_CLAUSE_MAP)
+ && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ATTACH_DETACH)
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
+ && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
+ (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
+ }
if (remove)
break;
new file mode 100644
@@ -0,0 +1,57 @@
+! { 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 \[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 \[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 \[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
+
new file mode 100644
@@ -0,0 +1,40 @@
+! { 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
+
new file mode 100644
@@ -0,0 +1,108 @@
+! { 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
new file mode 100644
@@ -0,0 +1,62 @@
+! { 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
new file mode 100644
@@ -0,0 +1,35 @@
+! { 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
new file mode 100644
@@ -0,0 +1,26 @@
+! { 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 } }
new file mode 100644
@@ -0,0 +1,29 @@
+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 } }
new file mode 100644
@@ -0,0 +1,47 @@
+! { 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
new file mode 100644
@@ -0,0 +1,33 @@
+! { 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
new file mode 100644
@@ -0,0 +1,32 @@
+! { 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
@@ -36,6 +36,10 @@ program main
call six ()
call seven ()
call eight ()
+ call nine ()
+ call ten ()
+ call eleven ()
+ call twelve ()
contains
! Implicitly mapped – but no pointers are mapped
@@ -408,7 +412,180 @@ contains
!$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
-! Fixed by the "Fortran pointers and member mappings" patch
-! { dg-xfail-run-if TODO { offload_device_nonshared_as } }
+ 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
+
+ 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