diff options
22 files changed, 1220 insertions, 88 deletions
diff --git a/gcc/fortran/dependency.cc b/gcc/fortran/dependency.cc index 632b398..9464703 100644 --- a/gcc/fortran/dependency.cc +++ b/gcc/fortran/dependency.cc @@ -2337,3 +2337,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; +} diff --git a/gcc/fortran/dependency.h b/gcc/fortran/dependency.h index fbbede8..a2458f7 100644 --- a/gcc/fortran/dependency.h +++ b/gcc/fortran/dependency.h @@ -40,5 +40,6 @@ int gfc_expr_is_one (gfc_expr *, int); 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 *); diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index c86a025..b5e1b4c 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1380,6 +1380,7 @@ typedef struct gfc_omp_namelist 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; diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index c6061ce..9c51ac4 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -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__ @@ -2491,36 +2492,24 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op, } 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); @@ -2624,6 +2613,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, @@ -2641,6 +2697,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]; @@ -3664,6 +3722,54 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, 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) @@ -3719,7 +3825,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); @@ -3738,7 +3843,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, 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; @@ -3764,24 +3870,78 @@ 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; + 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, @@ -3945,6 +4105,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; @@ -4787,7 +4964,7 @@ gfc_trans_oacc_executable_directive (gfc_code *code) 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); diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 26df5b0..9665746 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9151,6 +9151,25 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) 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 @@ -9186,9 +9205,7 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, 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); @@ -9374,7 +9391,7 @@ omp_get_attachment (omp_mapping_group *grp) 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); @@ -9469,7 +9486,7 @@ omp_group_last (tree *start_p) == 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) @@ -9636,33 +9653,32 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained, 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: @@ -10010,18 +10026,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_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; @@ -10424,6 +10467,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. @@ -10603,6 +10651,17 @@ 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_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))) @@ -11001,11 +11060,17 @@ 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_map_clause_descriptor_p (grp_mid)) + has_descriptor = true; + } if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL) { @@ -11028,7 +11093,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 @@ -11259,7 +11335,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 @@ -11315,7 +11391,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; @@ -11331,7 +11407,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 @@ -11364,7 +11442,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, @@ -11509,8 +11597,7 @@ omp_build_struct_sibling_lists (enum tree_code code, 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; } @@ -11786,6 +11873,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 @@ -11800,7 +11889,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); diff --git a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data-2.f90 b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data-2.f90 new file mode 100644 index 0000000..6a16c8a --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data-2.f90 @@ -0,0 +1,38 @@ +! { 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 diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f index 23f0ffc..63beb47 100644 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -20,8 +20,8 @@ ! { 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" } } @@ -32,6 +32,6 @@ ! { 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 diff --git a/gcc/testsuite/gfortran.dg/gomp/map-9.f90 b/gcc/testsuite/gfortran.dg/gomp/map-9.f90 index f930a49..8c8d4f7 100644 --- a/gcc/testsuite/gfortran.dg/gomp/map-9.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/map-9.f90 @@ -2,7 +2,7 @@ ! 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 diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90 new file mode 100644 index 0000000..26e113f --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90 @@ -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 \[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 + diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90 new file mode 100644 index 0000000..197888a --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90 @@ -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 + diff --git a/gcc/testsuite/gfortran.dg/gomp/target-enter-exit-data.f90 b/gcc/testsuite/gfortran.dg/gomp/target-enter-exit-data.f90 new file mode 100644 index 0000000..c14a11d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-enter-exit-data.f90 @@ -0,0 +1,39 @@ +! { 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 diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index cab99f9..e6d2ce0 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1050,6 +1050,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) 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; @@ -1831,6 +1831,10 @@ class auto_suppress_location_wrappers 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. */ diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90 new file mode 100644 index 0000000..02f08c5 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90 new file mode 100644 index 0000000..318e77e --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90 new file mode 100644 index 0000000..5d15808 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90 new file mode 100644 index 0000000..9f0edf7 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90 @@ -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 } } diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90 new file mode 100644 index 0000000..42da729 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90 @@ -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 } } diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90 new file mode 100644 index 0000000..a47360e --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray.f90 new file mode 100644 index 0000000..85f5af3 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90 b/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90 new file mode 100644 index 0000000..c7f9013 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 index 7f3d817..b1d6966 100644 --- a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 @@ -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 + + 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 |