@@ -9779,10 +9779,15 @@ omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
return outlist;
}
-/* Split INLIST into two parts, moving groups corresponding to
- ALLOC/RELEASE/DELETE mappings to one list, and other mappings to another.
- The former list is then appended to the latter. Each sub-list retains the
- order of the original list.
+/* Split INLIST into four parts:
+
+ - "present" to/from groups
+ - "present" alloc groups
+ - other to/from groups
+ - other alloc/release/delete groups
+
+ These sub-lists are then concatenated together to form the final list.
+ Each sub-list retains the order of the original list.
Note that ATTACH nodes are later moved to the end of the list in
gimplify_adjust_omp_clauses, for target regions. */
@@ -9790,7 +9795,9 @@ static omp_mapping_group *
omp_segregate_mapping_groups (omp_mapping_group *inlist)
{
omp_mapping_group *ard_groups = NULL, *tf_groups = NULL;
+ omp_mapping_group *pa_groups = NULL, *ptf_groups = NULL;
omp_mapping_group **ard_tail = &ard_groups, **tf_tail = &tf_groups;
+ omp_mapping_group **pa_tail = &pa_groups, **ptf_tail = &ptf_groups;
for (omp_mapping_group *w = inlist; w;)
{
@@ -9809,6 +9816,20 @@ omp_segregate_mapping_groups (omp_mapping_group *inlist)
ard_tail = &w->next;
break;
+ case GOMP_MAP_PRESENT_ALLOC:
+ *pa_tail = w;
+ w->next = NULL;
+ pa_tail = &w->next;
+ break;
+
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_TOFROM:
+ *ptf_tail = w;
+ w->next = NULL;
+ ptf_tail = &w->next;
+ break;
+
default:
*tf_tail = w;
w->next = NULL;
@@ -9820,8 +9841,10 @@ omp_segregate_mapping_groups (omp_mapping_group *inlist)
/* Now splice the lists together... */
*tf_tail = ard_groups;
+ *pa_tail = tf_groups;
+ *ptf_tail = pa_groups;
- return tf_groups;
+ return ptf_groups;
}
/* Given a list LIST_P containing groups of mappings given by GROUPS, reorder
@@ -11673,119 +11696,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
- if (code == OMP_TARGET
- || code == OMP_TARGET_DATA
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA)
- {
- vec<omp_mapping_group> *groups;
- groups = omp_gather_mapping_groups (list_p);
- if (groups)
- {
- hash_map<tree_operand_hash_no_se, omp_mapping_group *> *grpmap;
- grpmap = omp_index_mapping_groups (groups);
+ vec<omp_mapping_group> *groups = omp_gather_mapping_groups (list_p);
+ hash_map<tree_operand_hash_no_se, omp_mapping_group *> *grpmap = NULL;
+ unsigned grpnum = 0;
+ tree *grp_start_p = NULL, grp_end = NULL_TREE;
- omp_resolve_clause_dependencies (code, groups, grpmap);
- omp_build_struct_sibling_lists (code, region_type, groups, &grpmap,
- 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
- sorting in that case. */
- if (seen_error ())
- goto failure;
-
- delete grpmap;
- delete groups;
-
- /* Rebuild now we have struct sibling lists. */
- groups = omp_gather_mapping_groups (list_p);
- grpmap = omp_index_mapping_groups (groups);
-
- 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);
-
- failure:
- delete grpmap;
- delete groups;
- }
-
- /* OpenMP map clauses with 'present' need to go in front of those
- without. */
- tree present_map_head = NULL;
- tree *present_map_tail_p = &present_map_head;
- tree *first_map_clause_p = NULL;
-
- for (tree *c_p = list_p; *c_p; )
- {
- tree c = *c_p;
- tree *next_c_p = &OMP_CLAUSE_CHAIN (c);
-
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
- {
- if (!first_map_clause_p)
- first_map_clause_p = c_p;
- switch (OMP_CLAUSE_MAP_KIND (c))
- {
- case GOMP_MAP_PRESENT_ALLOC:
- case GOMP_MAP_PRESENT_FROM:
- case GOMP_MAP_PRESENT_TO:
- case GOMP_MAP_PRESENT_TOFROM:
- next_c_p = c_p;
- *c_p = OMP_CLAUSE_CHAIN (c);
-
- OMP_CLAUSE_CHAIN (c) = NULL;
- *present_map_tail_p = c;
- present_map_tail_p = &OMP_CLAUSE_CHAIN (c);
-
- break;
-
- default:
- break;
- }
- }
-
- c_p = next_c_p;
- }
- if (first_map_clause_p && present_map_head)
- {
- tree next = *first_map_clause_p;
- *first_map_clause_p = present_map_head;
- *present_map_tail_p = next;
- }
- }
- else if (region_type & ORT_ACC)
- {
- vec<omp_mapping_group> *groups;
- groups = omp_gather_mapping_groups (list_p);
- if (groups)
- {
- hash_map<tree_operand_hash_no_se, omp_mapping_group *> *grpmap;
- grpmap = omp_index_mapping_groups (groups);
-
- oacc_resolve_clause_dependencies (groups, grpmap);
- omp_build_struct_sibling_lists (code, region_type, groups, &grpmap,
- list_p);
-
- delete groups;
- delete grpmap;
- }
- }
+ if (groups)
+ grpmap = omp_index_mapping_groups (groups);
while ((c = *list_p) != NULL)
{
bool remove = false;
bool notice_outer = true;
+ bool map_descriptor;
const char *check_non_private = NULL;
unsigned int flags;
tree decl;
auto_vec<omp_addr_token *, 10> addr_tokens;
+ if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end))
+ {
+ grp_start_p = NULL;
+ grp_end = NULL_TREE;
+ }
+
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_PRIVATE:
@@ -12090,45 +12024,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
goto do_add;
case OMP_CLAUSE_MAP:
+ if (!grp_start_p)
+ {
+ grp_start_p = list_p;
+ grp_end = (*groups)[grpnum].grp_end;
+ grpnum++;
+ }
decl = OMP_CLAUSE_DECL (c);
+ if (error_operand_p (decl))
+ {
+ remove = true;
+ break;
+ }
+
if (!omp_parse_expr (addr_tokens, decl))
{
remove = true;
break;
}
- if (error_operand_p (decl))
- remove = true;
- switch (code)
- {
- case OMP_TARGET:
- break;
- case OACC_DATA:
- if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
- break;
- goto check_firstprivate;
- case OACC_ENTER_DATA:
- case OACC_EXIT_DATA:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
- && addr_tokens[0]->type == ARRAY_BASE)
- remove = true;
- /* FALLTHRU */
- case OMP_TARGET_DATA:
- case OMP_TARGET_ENTER_DATA:
- case OMP_TARGET_EXIT_DATA:
- case OACC_HOST_DATA:
- check_firstprivate:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
- /* For target {,enter ,exit }data only the array slice is
- mapped, but not the pointer to it. */
- remove = true;
- break;
- default:
- break;
- }
if (remove)
break;
if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC))
@@ -12147,49 +12062,55 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
DECL_NAME (decl));
}
}
- if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
- OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
- : TYPE_SIZE_UNIT (TREE_TYPE (decl));
- if (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
- {
- gcc_assert (OMP_CLAUSE_SIZE (c)
- && TREE_CODE (OMP_CLAUSE_SIZE (c)) == TREE_LIST);
- /* For non-contiguous array maps, OMP_CLAUSE_SIZE is a TREE_LIST
- of the individual array dimensions, which gimplify_expr doesn't
- handle, so skip the call to gimplify_expr here. */
- }
- else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
- NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
- {
- remove = true;
- break;
- }
- else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
- {
- OMP_CLAUSE_SIZE (c)
- = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
- false);
- if ((region_type & ORT_TARGET) != 0)
- omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
- GOVD_FIRSTPRIVATE | GOVD_SEEN);
- }
- if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT_UNORD)
- && (addr_tokens[0]->type == STRUCTURE_BASE
- || addr_tokens[0]->type == ARRAY_BASE)
- && addr_tokens[0]->u.structure_base_kind == BASE_DECL)
+ map_descriptor = false;
+
+ /* This condition checks if we're mapping an array descriptor that
+ isn't inside a derived type -- these have special handling, and
+ are not handled as structs in omp_build_struct_sibling_lists.
+ See that function for further details. */
+ if (*grp_start_p != grp_end
+ && OMP_CLAUSE_CHAIN (*grp_start_p)
+ && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
+ {
+ tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p);
+ if (omp_map_clause_descriptor_p (grp_mid)
+ && DECL_P (OMP_CLAUSE_DECL (grp_mid)))
+ map_descriptor = true;
+ }
+ else if (OMP_CLAUSE_CODE (grp_end) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_RELEASE
+ || OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_DELETE)
+ && OMP_CLAUSE_RELEASE_DESCRIPTOR (grp_end))
+ map_descriptor = true;
+
+ /* Adding the decl for a struct access: we haven't created
+ GOMP_MAP_STRUCT nodes yet, so this statement needs to predict
+ whether they will be created in gimplify_adjust_omp_clauses. */
+ if (c == grp_end
+ && addr_tokens[0]->type == STRUCTURE_BASE
+ && addr_tokens[0]->u.structure_base_kind == BASE_DECL
+ && !map_descriptor)
{
gcc_assert (addr_tokens[1]->type == ACCESS_METHOD);
/* If we got to this struct via a chain of pointers, maybe we
want to map it implicitly instead. */
if (omp_access_chain_p (addr_tokens, 1))
break;
+ omp_mapping_group *wholestruct;
+ if (!(region_type & ORT_ACC)
+ && omp_mapped_by_containing_struct (grpmap,
+ OMP_CLAUSE_DECL (c),
+ &wholestruct))
+ break;
decl = addr_tokens[1]->expr;
+ if (splay_tree_lookup (ctx->variables, (splay_tree_key) decl))
+ break;
+ /* Standalone attach or detach clauses for a struct element
+ should not inhibit implicit mapping of the whole struct. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ break;
flags = GOVD_MAP | GOVD_EXPLICIT;
gcc_assert (addr_tokens[1]->u.access_kind != ACCESS_DIRECT
@@ -12197,14 +12118,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
goto do_add_decl;
}
- if (TREE_CODE (decl) == TARGET_EXPR)
- {
- if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
- is_gimple_lvalue, fb_lvalue)
- == GS_ERROR)
- remove = true;
- }
- else if (!DECL_P (decl))
+ if (!DECL_P (decl))
{
tree d = decl, *pd;
if (TREE_CODE (d) == ARRAY_REF)
@@ -12227,56 +12141,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
pd = &TREE_OPERAND (decl, 0);
decl = TREE_OPERAND (decl, 0);
}
- /* An "attach/detach" operation on an update directive should
- behave as a GOMP_MAP_ALWAYS_POINTER. Beware that
- unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER
- depends on the previous mapping. */
- if (code == OACC_UPDATE
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ if (addr_tokens[0]->type == STRUCTURE_BASE
+ && addr_tokens[0]->u.structure_base_kind == BASE_DECL
+ && addr_tokens[1]->type == ACCESS_METHOD
+ && (addr_tokens[1]->u.access_kind == ACCESS_POINTER
+ || (addr_tokens[1]->u.access_kind
+ == ACCESS_POINTER_OFFSET))
+ && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)))
{
- if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c)))
- == ARRAY_TYPE)
- remove = true;
- else
- {
- gomp_map_kind k = ((code == OACC_EXIT_DATA
- || code == OMP_TARGET_EXIT_DATA)
- ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
- OMP_CLAUSE_SET_MAP_KIND (c, k);
- }
- }
-
- tree cref = decl;
-
- while (TREE_CODE (cref) == ARRAY_REF)
- cref = TREE_OPERAND (cref, 0);
-
- if (TREE_CODE (cref) == INDIRECT_REF)
- cref = TREE_OPERAND (cref, 0);
-
- if (TREE_CODE (cref) == COMPONENT_REF)
- {
- tree base = cref;
- while (base && !DECL_P (base))
- {
- tree innerbase = omp_get_base_pointer (base);
- if (!innerbase)
- break;
- base = innerbase;
- }
- if (base
- && DECL_P (base)
- && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
- && POINTER_TYPE_P (TREE_TYPE (base)))
- {
- splay_tree_node n
- = splay_tree_lookup (ctx->variables,
- (splay_tree_key) base);
- n->value |= GOVD_SEEN;
- }
+ tree base = addr_tokens[1]->expr;
+ splay_tree_node n
+ = splay_tree_lookup (ctx->variables,
+ (splay_tree_key) base);
+ n->value |= GOVD_SEEN;
}
if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
@@ -12387,13 +12265,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
}
}
- else if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
- fb_lvalue) == GS_ERROR)
- {
- remove = true;
- break;
- }
-
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -12402,44 +12273,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
flags |= GOVD_MAP_ALWAYS_TO;
else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
flags |= GOVD_DEVICEPTR;
-
- if ((code == OMP_TARGET
- || code == OMP_TARGET_DATA
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA)
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- {
- for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
- octx = octx->outer_context)
- {
- splay_tree_node n
- = splay_tree_lookup (octx->variables,
- (splay_tree_key) OMP_CLAUSE_DECL (c));
- /* If this is contained in an outer OpenMP region as a
- firstprivate value, remove the attach/detach. */
- if (n && (n->value & GOVD_FIRSTPRIVATE))
- {
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
- goto do_add;
- }
- }
-
- enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
- ? GOMP_MAP_DETACH
- : GOMP_MAP_ATTACH);
- OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
- }
- else if ((code == OACC_ENTER_DATA
- || code == OACC_EXIT_DATA
- || code == OACC_PARALLEL)
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- {
- enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA
- ? GOMP_MAP_DETACH
- : GOMP_MAP_ATTACH);
- OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
- }
-
goto do_add;
case OMP_CLAUSE_AFFINITY:
@@ -13131,6 +12964,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
list_p = &OMP_CLAUSE_CHAIN (c);
}
+ if (groups)
+ {
+ delete grpmap;
+ delete groups;
+ }
+
ctx->clauses = *orig_list_p;
gimplify_omp_ctxp = ctx;
}
@@ -13605,15 +13444,75 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
+ if (code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ {
+ vec<omp_mapping_group> *groups;
+ groups = omp_gather_mapping_groups (list_p);
+ hash_map<tree_operand_hash_no_se, omp_mapping_group *> *grpmap = NULL;
+
+ if (groups)
+ {
+ grpmap = omp_index_mapping_groups (groups);
+
+ omp_resolve_clause_dependencies (code, groups, grpmap);
+ omp_build_struct_sibling_lists (code, ctx->region_type, groups,
+ &grpmap, list_p);
+
+ omp_mapping_group *outlist = NULL;
+
+ delete grpmap;
+ delete groups;
+
+ /* Rebuild now we have struct sibling lists. */
+ groups = omp_gather_mapping_groups (list_p);
+ grpmap = omp_index_mapping_groups (groups);
+
+ bool enter_exit = (code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA);
+
+ 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);
+
+ delete grpmap;
+ delete groups;
+ }
+ }
+ else if (ctx->region_type & ORT_ACC)
+ {
+ vec<omp_mapping_group> *groups;
+ groups = omp_gather_mapping_groups (list_p);
+ if (groups)
+ {
+ hash_map<tree_operand_hash_no_se, omp_mapping_group *> *grpmap;
+ grpmap = omp_index_mapping_groups (groups);
+
+ oacc_resolve_clause_dependencies (groups, grpmap);
+ omp_build_struct_sibling_lists (code, ctx->region_type, groups,
+ &grpmap, list_p);
+
+ delete groups;
+ delete grpmap;
+ }
+ }
+
tree attach_list = NULL_TREE;
tree *attach_tail = &attach_list;
+ tree *grp_start_p = NULL, grp_end = NULL_TREE;
+
while ((c = *list_p) != NULL)
{
splay_tree_node n;
bool remove = false;
bool move_attach = false;
+ if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end))
+ grp_end = NULL_TREE;
+
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_FIRSTPRIVATE:
@@ -13768,6 +13667,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
break;
case OMP_CLAUSE_MAP:
+ decl = OMP_CLAUSE_DECL (c);
+ if (!grp_end)
+ {
+ grp_start_p = list_p;
+ grp_end = *omp_group_last (grp_start_p);
+ }
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_PRESENT_ALLOC:
@@ -13779,26 +13684,70 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
default:
break;
}
- if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+ switch (code)
{
+ case OMP_TARGET:
+ break;
+ case OACC_DATA:
+ if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
+ break;
+ goto check_firstprivate;
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ case OMP_TARGET_DATA:
+ case OMP_TARGET_ENTER_DATA:
+ case OMP_TARGET_EXIT_DATA:
+ case OACC_HOST_DATA:
+ check_firstprivate:
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ /* For target {,enter ,exit }data only the array slice is
+ mapped, but not the pointer to it. */
+ remove = true;
+ if (code == OMP_TARGET_EXIT_DATA
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER))
+ remove = true;
+ break;
+ default:
+ break;
+ }
+ if (remove)
+ break;
+ if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+ OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
+ : TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ gimplify_omp_ctxp = ctx->outer_context;
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ {
+ gcc_assert (OMP_CLAUSE_SIZE (c)
+ && TREE_CODE (OMP_CLAUSE_SIZE (c)) == TREE_LIST);
+ /* For non-contiguous array maps, OMP_CLAUSE_SIZE is a TREE_LIST
+ of the individual array dimensions, which gimplify_expr doesn't
+ handle, so skip the call to gimplify_expr here. */
+ }
+ else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL,
+ is_gimple_val, fb_rvalue) == GS_ERROR)
+ {
+ gimplify_omp_ctxp = ctx;
remove = true;
break;
}
- /* If we have a target region, we can push all the attaches to the
- end of the list (we may have standalone "attach" operations
- synthesized for GOMP_MAP_STRUCT nodes that must be processed after
- the attachment point AND the pointed-to block have been mapped).
- If we have something else, e.g. "enter data", we need to keep
- "attach" nodes together with the previous node they attach to so
- that separate "exit data" operations work properly (see
- libgomp/target.c). */
- if ((ctx->region_type & ORT_TARGET) != 0
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
- || (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
- move_attach = true;
- decl = OMP_CLAUSE_DECL (c);
+ else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
+ {
+ OMP_CLAUSE_SIZE (c)
+ = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
+ false);
+ if ((ctx->region_type & ORT_TARGET) != 0)
+ omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+ GOVD_FIRSTPRIVATE | GOVD_SEEN);
+ }
+ gimplify_omp_ctxp = ctx;
/* Data clauses associated with reductions must be
compatible with present_or_copy. Warn and adjust the clause
if that is not the case. */
@@ -13835,7 +13784,13 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
remove = true;
break;
}
- if (!DECL_P (decl))
+ if (TREE_CODE (decl) == TARGET_EXPR)
+ {
+ if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
+ is_gimple_lvalue, fb_lvalue) == GS_ERROR)
+ remove = true;
+ }
+ else if (!DECL_P (decl))
{
if ((ctx->region_type & ORT_TARGET) != 0
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
@@ -13858,8 +13813,122 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
}
- break;
- }
+
+ tree d = decl, *pd;
+ if (TREE_CODE (d) == ARRAY_REF)
+ {
+ while (TREE_CODE (d) == ARRAY_REF)
+ d = TREE_OPERAND (d, 0);
+ if (TREE_CODE (d) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
+ decl = d;
+ }
+ pd = &OMP_CLAUSE_DECL (c);
+ if (d == decl
+ && TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION))
+ {
+ pd = &TREE_OPERAND (decl, 0);
+ decl = TREE_OPERAND (decl, 0);
+ }
+
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ switch (code)
+ {
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c)))
+ == ARRAY_TYPE)
+ remove = true;
+ else if (code == OACC_ENTER_DATA)
+ goto change_to_attach;
+ /* Fallthrough. */
+ case OMP_TARGET_EXIT_DATA:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DETACH);
+ break;
+ case OACC_UPDATE:
+ /* An "attach/detach" operation on an update directive
+ should behave as a GOMP_MAP_ALWAYS_POINTER. Note that
+ both GOMP_MAP_ATTACH_DETACH and GOMP_MAP_ALWAYS_POINTER
+ kinds depend on the previous mapping (for non-TARGET
+ regions). */
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
+ break;
+ default:
+ change_to_attach:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH);
+ if ((ctx->region_type & ORT_TARGET) != 0)
+ move_attach = true;
+ }
+ else if ((ctx->region_type & ORT_TARGET) != 0
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
+ move_attach = true;
+
+ /* If we have e.g. map(struct: *var), don't gimplify the
+ argument since omp-low.cc wants to see the decl itself. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+ break;
+
+ /* We've already partly gimplified this in
+ gimplify_scan_omp_clauses. Don't do any more. */
+ if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
+ break;
+
+ gimplify_omp_ctxp = ctx->outer_context;
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+ fb_lvalue) == GS_ERROR)
+ remove = true;
+ gimplify_omp_ctxp = ctx;
+ break;
+ }
+
+ if ((code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ {
+ bool firstprivatize = false;
+
+ for (struct gimplify_omp_ctx *octx = ctx->outer_context; octx;
+ octx = octx->outer_context)
+ {
+ splay_tree_node n
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) OMP_CLAUSE_DECL (c));
+ /* If this is contained in an outer OpenMP region as a
+ firstprivate value, remove the attach/detach. */
+ if (n && (n->value & GOVD_FIRSTPRIVATE))
+ {
+ firstprivatize = true;
+ break;
+ }
+ }
+
+ enum gomp_map_kind map_kind;
+ if (firstprivatize)
+ map_kind = GOMP_MAP_FIRSTPRIVATE_POINTER;
+ else if (code == OMP_TARGET_EXIT_DATA)
+ map_kind = GOMP_MAP_DETACH;
+ else
+ map_kind = GOMP_MAP_ATTACH;
+ OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+ }
+ else if ((ctx->region_type & ORT_ACC) != 0
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ {
+ enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA
+ ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH);
+ OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+ }
+
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if ((ctx->region_type & ORT_TARGET) != 0
&& !(n->value & GOVD_SEEN)
@@ -13934,6 +14003,21 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
|| ((n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE))
== 0));
}
+
+ /* If we have a target region, we can push all the attaches to the
+ end of the list (we may have standalone "attach" operations
+ synthesized for GOMP_MAP_STRUCT nodes that must be processed after
+ the attachment point AND the pointed-to block have been mapped).
+ If we have something else, e.g. "enter data", we need to keep
+ "attach" nodes together with the previous node they attach to so
+ that separate "exit data" operations work properly (see
+ libgomp/target.c). */
+ if ((ctx->region_type & ORT_TARGET) != 0
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
+ move_attach = true;
+
break;
case OMP_CLAUSE_TO:
@@ -60,7 +60,7 @@ end subroutine
! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:b \\\[len: 4\\\]\\) map\\(to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(always,to:b \\\[len: 4\\\]\\) map\\(always,to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:a \\\[len: 4\\\]\\) map\\(force_present:b \\\[len: 4\\\]\\) map\\(force_present:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:b \\\[len: 4\\\]\\) map\\(force_present:b1 \\\[len: 4\\\]\\) map\\(force_present:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:a \\\[len: 4\\\]\\) map\\(always,present,to:b \\\[len: 4\\\]\\) map\\(always,present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }