public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-13] OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc
@ 2023-06-19 22:17 Julian Brown
0 siblings, 0 replies; only message in thread
From: Julian Brown @ 2023-06-19 22:17 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:626e2f4f9dd57d6800a8951c25068306a15bed41
commit 626e2f4f9dd57d6800a8951c25068306a15bed41
Author: Julian Brown <julian@codesourcery.com>
Date: Tue May 23 09:37:00 2023 +0000
OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc
This patch has been separated out from the C++ "declare mapper"
support patch. It contains just the gimplify.cc rearrangement
work, mostly moving gimplification from gimplify_scan_omp_clauses
to gimplify_adjust_omp_clauses for map clauses.
The motivation for doing this was that we don't know if we need to
instantiate mappers implicitly until the body of an offload region has
been scanned, i.e. in gimplify_adjust_omp_clauses, but we also need the
un-gimplified form of clauses to sort by base-pointer dependencies after
mapper instantiation has taken place.
The patch also reimplements the "present" clause sorting code to avoid
another sorting pass on mapping nodes.
2023-06-16 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.cc (omp_segregate_mapping_groups): Handle "present" groups.
(gimplify_scan_omp_clauses): Use mapping group functionality to
iterate through mapping nodes. Remove most gimplification of
OMP_CLAUSE_MAP nodes from here, but still populate ctx->variables
splay tree.
(gimplify_adjust_omp_clauses): Move most gimplification of
OMP_CLAUSE_MAP nodes here.
gcc/testsuite/
* gfortran.dg/gomp/map-12.f90: Adjust scan output.
Diff:
---
gcc/ChangeLog.omp | 10 +
gcc/gimplify.cc | 664 +++++++++++++++++-------------
gcc/testsuite/ChangeLog.omp | 4 +
gcc/testsuite/gfortran.dg/gomp/map-12.f90 | 2 +-
4 files changed, 389 insertions(+), 291 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index f8300b3c370..f4dfc6cc6cd 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,13 @@
+2023-06-19 Julian Brown <julian@codesourcery.com>
+
+ * gimplify.cc (omp_segregate_mapping_groups): Handle "present" groups.
+ (gimplify_scan_omp_clauses): Use mapping group functionality to
+ iterate through mapping nodes. Remove most gimplification of
+ OMP_CLAUSE_MAP nodes from here, but still populate ctx->variables
+ splay tree.
+ (gimplify_adjust_omp_clauses): Move most gimplification of
+ OMP_CLAUSE_MAP nodes here.
+
2023-06-19 Julian Brown <julian@codesourcery.com>
* gimplify.cc (extract_base_bit_offset): Add VARIABLE_OFFSET parameter.
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 9ce1f5b983a..e21e9d99cc9 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -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);
-
- 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;
+ 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;
- 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 (!omp_parse_expr (addr_tokens, decl))
+ if (error_operand_p (decl))
{
remove = true;
break;
}
- if (error_operand_p (decl))
- remove = true;
- switch (code)
+ if (!omp_parse_expr (addr_tokens, decl))
{
- 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:
+ remove = true;
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)
+
+ 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)
{
- 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);
+ 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;
}
-
- 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)
+ 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 (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)
+ 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)))
{
- 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:
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index f7ecfef7097..0d89832a180 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,7 @@
+2023-06-19 Julian Brown <julian@codesourcery.com>
+
+ * gfortran.dg/gomp/map-12.f90: Adjust scan output.
+
2023-06-19 Julian Brown <julian@codesourcery.com>
* c-c++-comon/gomp/target-implicit-map-2.c: Update expected scan output.
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-12.f90 b/gcc/testsuite/gfortran.dg/gomp/map-12.f90
index ac9a0f8aae0..433bf98911c 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-12.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-12.f90
@@ -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" } }
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2023-06-19 22:17 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-06-19 22:17 [gcc/devel/omp/gcc-13] OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc Julian Brown
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).