This patch lifts struct sibling-list handling out of the main loop in
gimplify_scan_omp_clauses. The reasons for this are several: first,
it means that we can subject created sibling list groups to topological
sorting (see previous patch) so base-pointer data dependencies are
handled correctly.
Secondly, it means that in the first pass gathering up sibling lists
from parsed OpenMP/OpenACC clauses, we don't need to worry about
gimplifying: that means we can see struct bases & components we need
to sort sibling lists properly, even when we're using a non-DECL_P
struct base. Gimplification proper still happens
Thirdly, because we use more than one pass through the clause list and
gather appropriate data, we can tell if we're mapping a whole struct
in a different node, and avoid building struct sibling lists for that
struct appropriately.
Fourthly, we can re-use the node grouping functions from the
previous patch, and thus mostly avoid the "prev_list_p" handling in
gimplify_scan_omp_clauses that tracks the first node in such groups
at present.
Some redundant code has been removed and code paths for OpenACC/OpenMP are
now shared where appropriate, though OpenACC doesn't do the topological
sorting of nodes (yet?).
OK for mainline?
Thanks,
Julian
2021-09-29 Julian Brown <jul...@codesourcery.com>
gcc/
* gimplify.c (gimplify_omp_var_data): Remove GOVD_MAP_HAS_ATTACHMENTS.
(extract_base_bit_offset): Remove OFFSETP parameter.
(strip_components_and_deref): Extend with POINTER_PLUS_EXPR and
COMPOUND_EXPR handling.
(aggregate_base_p): Remove.
(omp_group_last, omp_group_base): Add GOMP_MAP_STRUCT handling.
(build_struct_group): Remove CTX, DECL, PD, COMPONENT_REF_P, FLAGS,
STRUCT_SEEN_CLAUSE, PRE_P, CONT parameters. Replace PREV_LIST_P and C
parameters with GRP_START_P and GRP_END. Add INNER. Update calls to
extract_base_bit_offset. Remove gimplification of clauses for OpenMP.
Rework inner struct handling for OpenACC. Don't use context's
variables splay tree.
(omp_build_struct_sibling_lists): New function, extracted from
gimplify_scan_omp_clauses and refactored.
(gimplify_scan_omp_clauses): Call above function to handle struct
sibling lists. Remove STRUCT_MAP_TO_CLAUSE, STRUCT_SEEN_CLAUSE,
STRUCT_DEREF_SET. Rework flag handling, adding decl for struct
variables.
(gimplify_adjust_omp_clauses_1): Remove GOVD_MAP_HAS_ATTACHMENTS
handling, unused now.
gcc/testsuite/
* g++.dg/goacc/member-array-acc.C: Update expected output.
* g++.dg/gomp/target-3.C: Likewise.
* g++.dg/gomp/target-lambda-1.C: Likewise.
* g++.dg/gomp/target-this-2.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
---
gcc/gimplify.c | 943 ++++++++----------
gcc/testsuite/g++.dg/goacc/member-array-acc.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-3.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-this-2.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +-
6 files changed, 410 insertions(+), 547 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index c10a3e8842a..31e2e4d9fe7 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -125,10 +125,6 @@ enum gimplify_omp_var_data
/* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause. */
GOVD_REDUCTION_INSCAN = 0x2000000,
- /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for
- fields. */
- GOVD_MAP_HAS_ATTACHMENTS = 0x4000000,
-
/* Flag for GOVD_FIRSTPRIVATE: OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT. */
GOVD_FIRSTPRIVATE_IMPLICIT = 0x8000000,
@@ -8642,7 +8638,7 @@ build_struct_comp_nodes (enum tree_code code, tree
grp_start, tree grp_end,
static tree
extract_base_bit_offset (tree base, poly_int64 *bitposp,
- poly_offset_int *poffsetp, tree *offsetp)
+ poly_offset_int *poffsetp)
{
tree offset;
poly_int64 bitsize, bitpos;
@@ -8670,7 +8666,6 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp,
*bitposp = bitpos;
*poffsetp = poffset;
- *offsetp = offset;
return base;
}
@@ -8683,8 +8678,15 @@ strip_components_and_deref (tree expr)
while (TREE_CODE (expr) == COMPONENT_REF
|| TREE_CODE (expr) == INDIRECT_REF
|| (TREE_CODE (expr) == MEM_REF
- && integer_zerop (TREE_OPERAND (expr, 1))))
- expr = TREE_OPERAND (expr, 0);
+ && integer_zerop (TREE_OPERAND (expr, 1)))
+ || TREE_CODE (expr) == POINTER_PLUS_EXPR
+ || TREE_CODE (expr) == COMPOUND_EXPR)
+ if (TREE_CODE (expr) == COMPOUND_EXPR)
+ expr = TREE_OPERAND (expr, 1);
+ else
+ expr = TREE_OPERAND (expr, 0);
+
+ STRIP_NOPS (expr);
return expr;
}
@@ -8700,34 +8702,6 @@ strip_indirections (tree expr)
return expr;
}
-/* Return TRUE if EXPR is something we will use as the base of an aggregate
- access, either:
-
- - a DECL_P.
- - a struct component with no indirection ("a.b.c").
- - a struct component with indirection ("a->b->c").
-*/
-
-static bool
-aggregate_base_p (tree expr)
-{
- while (TREE_CODE (expr) == COMPONENT_REF
- && (DECL_P (TREE_OPERAND (expr, 0))
- || (TREE_CODE (TREE_OPERAND (expr, 0)) == COMPONENT_REF)))
- expr = TREE_OPERAND (expr, 0);
-
- if (DECL_P (expr))
- return true;
-
- if (TREE_CODE (expr) == COMPONENT_REF
- && (TREE_CODE (TREE_OPERAND (expr, 0)) == INDIRECT_REF
- || (TREE_CODE (TREE_OPERAND (expr, 0)) == MEM_REF
- && integer_zerop (TREE_OPERAND (TREE_OPERAND (expr, 0), 1)))))
- return true;
-
- return false;
-}
-
enum omp_tsort_mark {
UNVISITED,
TEMPORARY,
@@ -8956,6 +8930,18 @@ omp_group_last (tree *start_p)
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH))
grp_last_p = &OMP_CLAUSE_CHAIN (c);
break;
+
+ case GOMP_MAP_STRUCT:
+ {
+ unsigned HOST_WIDE_INT num_mappings
+ = tree_to_uhwi (OMP_CLAUSE_SIZE (c));
+ if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ grp_last_p = &OMP_CLAUSE_CHAIN (*grp_last_p);
+ for (unsigned i = 0; i < num_mappings; i++)
+ grp_last_p = &OMP_CLAUSE_CHAIN (*grp_last_p);
+ }
+ break;
}
return grp_last_p;
@@ -9089,6 +9075,21 @@ omp_group_base (omp_mapping_group *grp, unsigned int
*chained,
internal_error ("unexpected mapping node");
return error_mark_node;
+ case GOMP_MAP_STRUCT:
+ {
+ unsigned HOST_WIDE_INT num_mappings
+ = tree_to_uhwi (OMP_CLAUSE_SIZE (node));
+ node = OMP_CLAUSE_CHAIN (node);
+ if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ {
+ *firstprivate = OMP_CLAUSE_DECL (node);
+ node = OMP_CLAUSE_CHAIN (node);
+ }
+ *chained = num_mappings;
+ return node;
+ }
+
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
case GOMP_MAP_LINK:
@@ -9751,21 +9752,16 @@ move_concat_nodes_after (tree first_new, tree
*last_new_tail, tree *first_ptr,
next node. PREV_LIST_P and LIST_P may be modified by the function when a
list rearrangement has taken place. */
-static tree
-build_struct_group (struct gimplify_omp_ctx *ctx,
- enum omp_region_type region_type, enum tree_code code,
- tree decl, tree *pd, bool component_ref_p,
- unsigned int *flags, tree c,
+static tree *
+build_struct_group (enum omp_region_type region_type, enum tree_code code,
hash_map<tree_operand_hash, tree> *&struct_map_to_clause,
- hash_map<tree_operand_hash, tree *> *&struct_seen_clause,
- tree *&prev_list_p, tree *&list_p, gimple_seq *pre_p,
- bool *cont)
+ tree *grp_start_p, tree grp_end, tree *inner)
{
poly_offset_int coffset;
poly_int64 cbitpos;
- tree tree_coffset;
- tree ocd = OMP_CLAUSE_DECL (c);
+ tree ocd = OMP_CLAUSE_DECL (grp_end);
bool openmp = !(region_type & ORT_ACC);
+ tree *continue_at = NULL;
while (TREE_CODE (ocd) == ARRAY_REF)
ocd = TREE_OPERAND (ocd, 0);
@@ -9773,90 +9769,31 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
if (TREE_CODE (ocd) == INDIRECT_REF)
ocd = TREE_OPERAND (ocd, 0);
- tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset, &tree_coffset);
- tree sbase;
+ tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset);
- if (openmp)
- {
- if (TREE_CODE (base) == INDIRECT_REF
- && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE)
- sbase = strip_indirections (base);
- else
- sbase = base;
- }
- else
- {
- sbase = strip_indirections (base);
+ 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 attach = (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_DETACH);
- STRIP_NOPS (sbase);
- }
-
- bool do_map_struct = (sbase == decl && !tree_coffset);
-
- /* Here, DECL is usually a DECL_P, unless we have chained indirect member
- accesses, e.g. mystruct->a->b. In that case it'll be the "mystruct->a"
- part. */
- splay_tree_node n
- = (DECL_P (decl)
- ? splay_tree_lookup (ctx->variables, (splay_tree_key) decl)
- : NULL);
- bool ptr = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER);
- bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH);
- bool attach = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH);
- bool has_attachments = false;
-
- /* For OpenACC, pointers in structs should trigger an attach action. */
- if (attach_detach
- && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA))
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA))
- {
- /* Turn a GOMP_MAP_ATTACH_DETACH clause into a GOMP_MAP_ATTACH or
- GOMP_MAP_DETACH clause after we have detected a case that needs a
- GOMP_MAP_STRUCT mapping added. */
- 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);
- has_attachments = true;
- }
-
- /* We currently don't handle non-constant offset accesses wrt to
- GOMP_MAP_STRUCT elements. */
- if (!do_map_struct)
- return NULL_TREE;
-
- /* Nor for attach_detach for OpenMP. */
+ /* FIXME: If we're not mapping the base pointer in some other clause on this
+ directive, I think we want to create ALLOC/RELEASE here -- i.e. not
+ early-exit. */
if (openmp && attach_detach)
- {
- if (DECL_P (decl))
- {
- if (struct_seen_clause == NULL)
- struct_seen_clause = new hash_map<tree_operand_hash, tree *>;
- if (!struct_seen_clause->get (decl))
- struct_seen_clause->put (decl, list_p);
- }
+ return NULL;
- return NULL_TREE;
- }
-
- if ((DECL_P (decl) && (n == NULL || (n->value & GOVD_MAP) == 0))
- || (!DECL_P (decl)
- && (!struct_map_to_clause
- || struct_map_to_clause->get (decl) == NULL)))
+ if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
{
- tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ tree l = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
OMP_CLAUSE_MAP);
gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT;
OMP_CLAUSE_SET_MAP_KIND (l, k);
OMP_CLAUSE_DECL (l) = unshare_expr (base);
- if (openmp
- && !DECL_P (OMP_CLAUSE_DECL (l))
- && (gimplify_expr (&OMP_CLAUSE_DECL (l), pre_p, NULL,
- is_gimple_lvalue, fb_lvalue) == GS_ERROR))
- return error_mark_node;
+
OMP_CLAUSE_SIZE (l)
= (!attach ? size_int (1)
: (DECL_P (OMP_CLAUSE_DECL (l))
@@ -9864,19 +9801,17 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
: TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l)))));
if (struct_map_to_clause == NULL)
struct_map_to_clause = new hash_map<tree_operand_hash, tree>;
- struct_map_to_clause->put (decl, l);
+ struct_map_to_clause->put (base, l);
if (ptr || attach_detach)
{
tree extra_node;
tree alloc_node
- = build_struct_comp_nodes (code, *prev_list_p, c, &extra_node);
+ = build_struct_comp_nodes (code, *grp_start_p, grp_end,
+ &extra_node);
OMP_CLAUSE_CHAIN (l) = alloc_node;
- tree **sc = (struct_seen_clause
- ? struct_seen_clause->get (decl)
- : NULL);
- tree *insert_node_pos = sc ? *sc : prev_list_p;
+ tree *insert_node_pos = grp_start_p;
if (extra_node)
{
@@ -9887,131 +9822,89 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
OMP_CLAUSE_CHAIN (alloc_node) = *insert_node_pos;
*insert_node_pos = l;
- prev_list_p = NULL;
}
else
- list_p = insert_node_after (l, list_p);
-
- bool base_ref
- = (TREE_CODE (base) == INDIRECT_REF
- && ((TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0)))
- == REFERENCE_TYPE)
- || ((TREE_CODE (TREE_OPERAND (base, 0)) == INDIRECT_REF)
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND
- (TREE_OPERAND (base, 0), 0)))
- == REFERENCE_TYPE))));
- bool base_ind = ((TREE_CODE (base) == INDIRECT_REF
- || (TREE_CODE (base) == MEM_REF
- && integer_zerop (TREE_OPERAND (base, 1))))
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0)))
- == POINTER_TYPE));
-
- /* Handle pointers to structs and references to structs: these cases
- have an additional GOMP_MAP_FIRSTPRIVATE_{REFERENCE,POINTER} node
- inserted after the GOMP_MAP_STRUCT node. References to pointers
- use GOMP_MAP_FIRSTPRIVATE_REFERENCE. */
- if (base_ref && code == OMP_TARGET)
{
- tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- enum gomp_map_kind mkind
- = GOMP_MAP_FIRSTPRIVATE_REFERENCE;
- OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
- OMP_CLAUSE_DECL (c2) = decl;
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
- OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
- OMP_CLAUSE_CHAIN (l) = c2;
+ gcc_assert (*grp_start_p == grp_end);
+ grp_start_p = insert_node_after (l, grp_start_p);
}
- else if (!openmp
- && (base_ind || base_ref)
- && (region_type & ORT_TARGET))
+
+ tree noind = strip_indirections (base);
+
+ if (!openmp
+ && (region_type & ORT_TARGET)
+ && TREE_CODE (noind) == COMPONENT_REF)
{
- tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ /* The base for this component access is a struct component access
+ itself. Insert a node to be processed on the next iteration of
+ our caller's loop, which will subsequently be turned into a new,
+ inner GOMP_MAP_STRUCT mapping.
+
+ We need to do this else the non-DECL_P base won't be
+ rewritten correctly in the offloaded region. */
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT);
+ OMP_CLAUSE_DECL (c2) = unshare_expr (noind);
+ OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (TREE_TYPE (noind));
+ *inner = c2;
+ return NULL;
+ }
+
+ tree sdecl = strip_components_and_deref (base);
+
+ if (POINTER_TYPE_P (TREE_TYPE (sdecl)) && (region_type & ORT_TARGET))
+ {
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end),
+ OMP_CLAUSE_MAP);
+ bool base_ref
+ = (TREE_CODE (base) == INDIRECT_REF
+ && ((TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0)))
+ == REFERENCE_TYPE)
+ || ((TREE_CODE (TREE_OPERAND (base, 0))
+ == INDIRECT_REF)
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND
+ (TREE_OPERAND (base, 0), 0)))
+ == REFERENCE_TYPE))));
enum gomp_map_kind mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE
: GOMP_MAP_FIRSTPRIVATE_POINTER;
OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
- tree sdecl = strip_components_and_deref (decl);
- if (DECL_P (decl)
- && (POINTER_TYPE_P (TREE_TYPE (sdecl))
- || TREE_CODE (TREE_TYPE (sdecl)) == REFERENCE_TYPE))
- {
- /* Insert after struct node. */
- OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
- OMP_CLAUSE_DECL (c2) = decl;
- OMP_CLAUSE_CHAIN (l) = c2;
- }
- else
- {
- /* If the ultimate base for this component access is not a
- pointer or reference, that means it is a struct component
- access itself. Insert a node to be processed on the next
- iteration of our caller's loop, which will subsequently be
- turned into a new GOMP_MAP_STRUCT mapping itself.
-
- We need to do this else the non-DECL_P base won't be
- rewritten correctly in the offloaded region. */
- tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT);
- OMP_CLAUSE_DECL (c2) = unshare_expr (decl);
- OMP_CLAUSE_SIZE (c2) = (DECL_P (decl)
- ? DECL_SIZE_UNIT (decl)
- : TYPE_SIZE_UNIT (TREE_TYPE (decl)));
- tree *next_node = &OMP_CLAUSE_CHAIN (*list_p);
- OMP_CLAUSE_CHAIN (c2) = *next_node;
- *next_node = c2;
- return NULL_TREE;
- }
- }
- *flags = GOVD_MAP | GOVD_EXPLICIT;
- if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach)
- *flags |= GOVD_SEEN;
- if (has_attachments)
- *flags |= GOVD_MAP_HAS_ATTACHMENTS;
-
- /* If this is a *pointer-to-struct expression, make sure a
- firstprivate map of the base-pointer exists. */
- if (openmp
- && component_ref_p
- && ((TREE_CODE (decl) == MEM_REF
- && integer_zerop (TREE_OPERAND (decl, 1)))
- || INDIRECT_REF_P (decl))
- && DECL_P (TREE_OPERAND (decl, 0))
- && !splay_tree_lookup (ctx->variables,
- ((splay_tree_key) TREE_OPERAND (decl, 0))))
- {
- decl = TREE_OPERAND (decl, 0);
- tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- enum gomp_map_kind mkind = GOMP_MAP_FIRSTPRIVATE_POINTER;
- OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
- OMP_CLAUSE_DECL (c2) = decl;
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
- OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = c2;
+ OMP_CLAUSE_DECL (c2) = sdecl;
+ tree baddr = build_fold_addr_expr (base);
+ baddr = fold_convert_loc (OMP_CLAUSE_LOCATION (grp_end),
+ ptrdiff_type_node, baddr);
+ /* This isn't going to be good enough when we add support for more
+ complicated lvalue expressions. FIXME. */
+ if (TREE_CODE (TREE_TYPE (sdecl)) == REFERENCE_TYPE
+ && TREE_CODE (TREE_TYPE (TREE_TYPE (sdecl))) == POINTER_TYPE)
+ sdecl = build_simple_mem_ref (sdecl);
+ tree decladdr = fold_convert_loc (OMP_CLAUSE_LOCATION (grp_end),
+ ptrdiff_type_node, sdecl);
+ OMP_CLAUSE_SIZE (c2)
+ = fold_build2_loc (OMP_CLAUSE_LOCATION (grp_end), MINUS_EXPR,
+ ptrdiff_type_node, baddr, decladdr);
+ /* Insert after struct node. */
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
+ OMP_CLAUSE_CHAIN (l) = c2;
}
- return decl;
+ return NULL;
}
else if (struct_map_to_clause)
{
- tree *osc = struct_map_to_clause->get (decl);
+ tree *osc = struct_map_to_clause->get (base);
tree *sc = NULL, *scp = NULL;
- if (n && (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
- || ptr
- || attach_detach))
- n->value |= GOVD_SEEN;
sc = &OMP_CLAUSE_CHAIN (*osc);
/* The struct mapping might be immediately followed by a
FIRSTPRIVATE_POINTER and/or FIRSTPRIVATE_REFERENCE -- if it's an
indirect access or a reference, or both. (This added node is removed
in omp-low.c after it has been processed there.) */
- if (*sc != c
+ if (*sc != grp_end
&& (OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
sc = &OMP_CLAUSE_CHAIN (*sc);
- for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
- if ((ptr || attach_detach) && sc == prev_list_p)
+ for (; *sc != grp_end; sc = &OMP_CLAUSE_CHAIN (*sc))
+ if ((ptr || 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
@@ -10022,7 +9915,6 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
tree sc_decl = OMP_CLAUSE_DECL (*sc);
poly_offset_int offset;
poly_int64 bitpos;
- tree tree_offset;
if (TREE_CODE (sc_decl) == ARRAY_REF)
{
@@ -10038,8 +9930,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
== REFERENCE_TYPE))
sc_decl = TREE_OPERAND (sc_decl, 0);
- tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset,
- &tree_offset);
+ tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset);
if (!base2 || !operand_equal_p (base2, base, 0))
break;
if (scp)
@@ -10049,7 +9940,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
/* This duplicate checking code is currently only enabled for
OpenACC. */
tree d1 = OMP_CLAUSE_DECL (*sc);
- tree d2 = OMP_CLAUSE_DECL (c);
+ tree d2 = OMP_CLAUSE_DECL (grp_end);
while (TREE_CODE (d1) == ARRAY_REF)
d1 = TREE_OPERAND (d1, 0);
while (TREE_CODE (d2) == ARRAY_REF)
@@ -10069,10 +9960,10 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
break;
if (d1 == d2)
{
- error_at (OMP_CLAUSE_LOCATION (c),
+ error_at (OMP_CLAUSE_LOCATION (grp_end),
"%qE appears more than once in map clauses",
- OMP_CLAUSE_DECL (c));
- return error_mark_node;
+ OMP_CLAUSE_DECL (grp_end));
+ return NULL;
}
}
if (maybe_lt (coffset, offset)
@@ -10092,15 +9983,15 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
if (ptr || attach_detach)
{
tree cl = NULL_TREE, extra_node;
- tree alloc_node = build_struct_comp_nodes (code, *prev_list_p, c,
- &extra_node);
+ tree alloc_node = build_struct_comp_nodes (code, *grp_start_p,
+ grp_end, &extra_node);
tree *tail_chain = NULL;
/* Here, we have:
- c : the currently-processed node.
- prev_list_p : pointer to the first node in a pointer mapping group
- up to and including C.
+ grp_end : the last (or only) node in this group.
+ grp_start_p : pointer to the first node in a pointer mapping group
+ up to and including GRP_END.
sc : pointer to the chain for the end of the struct component
list.
scp : pointer to the chain for the sorted position at which we
@@ -10111,7 +10002,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
(the end of the struct component list).
extra_node : a newly-synthesized node for an additional indirect
pointer mapping or a Fortran pointer set, if needed.
- cl : first node to prepend before prev_list_p.
+ cl : first node to prepend before grp_start_p.
tail_chain : pointer to chain of last prepended node.
The general idea is we move the nodes for this struct mapping
@@ -10147,32 +10038,180 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
tail_chain = &OMP_CLAUSE_CHAIN (alloc_node);
}
- tree *continue_at
- = cl ? move_concat_nodes_after (cl, tail_chain, prev_list_p, c, sc)
- : move_nodes_after (prev_list_p, c, sc);
-
- prev_list_p = NULL;
-
- if (continue_at)
- {
- list_p = continue_at;
- *cont = true;
- }
+ continue_at
+ = cl ? move_concat_nodes_after (cl, tail_chain, grp_start_p,
+ grp_end, sc)
+ : move_nodes_after (grp_start_p, grp_end, sc);
}
- else if (*sc != c)
+ else if (*sc != grp_end)
{
- if (openmp
- && (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
- == GS_ERROR))
- return error_mark_node;
+ gcc_assert (*grp_start_p == grp_end);
+
+ /* We are moving the current node back to a previous struct node:
+ the node that used to point to the current node will now point to
+ the next node. */
+ continue_at = grp_start_p;
/* In the non-pointer case, the mapping clause itself is moved into
the correct position in the struct component list, which in this
case is just SC. */
- move_node_after (c, list_p, sc);
- *cont = true;
+ move_node_after (grp_end, grp_start_p, sc);
}
}
- return NULL_TREE;
+ return continue_at;
+}
+
+/* Scan through GROUPS, and create sorted structure sibling lists without
+ gimplifying. */
+
+static bool
+omp_build_struct_sibling_lists (enum tree_code code,
+ enum omp_region_type region_type,
+ vec<omp_mapping_group> *groups,
+ hash_map<tree_operand_hash, omp_mapping_group *>
+ *grpmap)
+{
+ unsigned i;
+ omp_mapping_group *grp;
+ hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
+ bool success = true;
+ tree *new_next = NULL;
+ tree *tail = &OMP_CLAUSE_CHAIN ((*groups)[groups->length () - 1].grp_end);
+
+ FOR_EACH_VEC_ELT (*groups, i, grp)
+ {
+ tree c = grp->grp_end;
+ tree decl = OMP_CLAUSE_DECL (c);
+ tree *grp_start_p = new_next ? new_next : grp->grp_start;
+ tree grp_end = grp->grp_end;
+
+ new_next = NULL;
+
+ if (DECL_P (decl))
+ continue;
+
+ if (OMP_CLAUSE_CHAIN (*grp_start_p)
+ && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
+ {
+ /* Don't process an array descriptor that isn't inside a derived type
+ 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
+ && DECL_P (OMP_CLAUSE_DECL (grpmid)))
+ continue;
+ }
+
+ tree d = decl;
+ 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;
+ }
+ 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))
+ decl = TREE_OPERAND (decl, 0);
+
+ STRIP_NOPS (decl);
+
+ if (TREE_CODE (decl) != COMPONENT_REF)
+ continue;
+
+ omp_mapping_group **wholestruct = NULL;
+ tree wsdecl = omp_containing_struct (OMP_CLAUSE_DECL (c));
+
+ if (!(region_type & ORT_ACC) && wsdecl != OMP_CLAUSE_DECL (c))
+ {
+ wholestruct = grpmap->get (wsdecl);
+ if (!wholestruct
+ && TREE_CODE (wsdecl) == MEM_REF
+ && integer_zerop (TREE_OPERAND (wsdecl, 1)))
+ {
+ tree deref = TREE_OPERAND (wsdecl, 0);
+ deref = build1 (INDIRECT_REF, TREE_TYPE (wsdecl), deref);
+ wholestruct = grpmap->get (deref);
+ }
+ }
+
+ if (wholestruct)
+ {
+
+ if (*grp_start_p == grp_end)
+ {
+ /* Remove the whole of this mapping -- redundant. */
+ new_next = grp_start_p;
+ *grp_start_p = OMP_CLAUSE_CHAIN (grp_end);
+ }
+
+ continue;
+ }
+
+ if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
+ && code != OACC_UPDATE
+ && code != OMP_TARGET_UPDATE)
+ {
+ if (error_operand_p (decl))
+ {
+ success = false;
+ goto error_out;
+ }
+
+ tree stype = TREE_TYPE (decl);
+ if (TREE_CODE (stype) == REFERENCE_TYPE)
+ stype = TREE_TYPE (stype);
+ if (TYPE_SIZE_UNIT (stype) == NULL
+ || TREE_CODE (TYPE_SIZE_UNIT (stype)) != INTEGER_CST)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "mapping field %qE of variable length "
+ "structure", OMP_CLAUSE_DECL (c));
+ success = false;
+ goto error_out;
+ }
+
+ tree inner = NULL_TREE;
+
+ new_next = build_struct_group (region_type, code,
+ struct_map_to_clause, grp_start_p,
+ grp_end, &inner);
+
+ if (inner)
+ {
+ if (new_next && *new_next == NULL_TREE)
+ *new_next = inner;
+ else
+ *tail = inner;
+
+ OMP_CLAUSE_CHAIN (inner) = NULL_TREE;
+
+ omp_mapping_group newgrp;
+ newgrp.grp_start = new_next ? new_next : tail;
+ newgrp.grp_end = inner;
+ newgrp.mark = UNVISITED;
+ newgrp.sibling = NULL;
+ newgrp.next = NULL;
+ groups->safe_push (newgrp);
+
+ tail = &OMP_CLAUSE_CHAIN (inner);
+ }
+ }
+ }
+
+error_out:
+ if (struct_map_to_clause)
+ delete struct_map_to_clause;
+
+ return success;
}
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
@@ -10185,9 +10224,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
{
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
- hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
- hash_map<tree_operand_hash, tree *> *struct_seen_clause = NULL;
- hash_set<tree> *struct_deref_set = NULL;
tree *prev_list_p = NULL, *orig_list_p = list_p;
int handled_depend_iterators = -1;
int nowait = -1;
@@ -10219,14 +10255,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
break;
}
- /* 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 ()
- && (code == OMP_TARGET
- || code == OMP_TARGET_DATA
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA))
+ 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);
@@ -10234,12 +10266,46 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
{
hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
grpmap = omp_index_mapping_groups (groups);
- omp_mapping_group *outlist
- = omp_tsort_mapping_groups (groups, grpmap);
- outlist = omp_segregate_mapping_groups (outlist);
- list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
+
+ omp_build_struct_sibling_lists (code, region_type, groups, grpmap);
+
+ omp_mapping_group *outlist = NULL;
+
+ /* 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);
+ outlist = omp_segregate_mapping_groups (outlist);
+ list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
+
+ failure:
+ delete grpmap;
+ delete groups;
+ }
+ }
+ 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, omp_mapping_group *> *grpmap;
+ grpmap = omp_index_mapping_groups (groups);
+
+ omp_build_struct_sibling_lists (code, region_type, groups, grpmap);
+
+ delete groups;
+ delete grpmap;
}
}
@@ -10648,6 +10714,28 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
GOVD_FIRSTPRIVATE | GOVD_SEEN);
}
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+ {
+ tree base = strip_components_and_deref (decl);
+ if (DECL_P (base))
+ {
+ decl = base;
+ splay_tree_node n
+ = splay_tree_lookup (ctx->variables,
+ (splay_tree_key) decl);
+ if (seen_error ()
+ && n
+ && (n->value & (GOVD_MAP | GOVD_FIRSTPRIVATE)) != 0)
+ {
+ remove = true;
+ break;
+ }
+ flags = GOVD_MAP | GOVD_EXPLICIT;
+
+ goto do_add_decl;
+ }
+ }
+
if (TREE_CODE (decl) == TARGET_EXPR)
{
if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
@@ -10678,143 +10766,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
pd = &TREE_OPERAND (decl, 0);
decl = TREE_OPERAND (decl, 0);
}
- bool indir_p = false;
- bool component_ref_p = false;
- tree indir_base = NULL_TREE;
- tree orig_decl = decl;
- tree decl_ref = NULL_TREE;
- if ((region_type & ORT_ACC) && TREE_CODE (decl) == COMPONENT_REF)
- {
- /* Strip off component refs from RHS of e.g. "a->b->c.d.e"
- (which would leave "a->b" in that case). This is intended
- to be equivalent to the base finding done by
- get_inner_reference. */
- while (TREE_CODE (decl) == COMPONENT_REF
- && (DECL_P (TREE_OPERAND (decl, 0))
- || (TREE_CODE (TREE_OPERAND (decl, 0))
- == COMPONENT_REF)))
- decl = TREE_OPERAND (decl, 0);
-
- if (TREE_CODE (decl) == COMPONENT_REF)
- decl = TREE_OPERAND (decl, 0);
-
- /* Strip off RHS from "a->b". */
- if ((TREE_CODE (decl) == INDIRECT_REF
- || (TREE_CODE (decl) == MEM_REF
- && integer_zerop (TREE_OPERAND (decl, 1))))
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == POINTER_TYPE))
- decl = TREE_OPERAND (decl, 0);
-
- /* Strip off RHS from "a_ref.b" (where a_ref is
- reference-typed). */
- if (TREE_CODE (decl) == INDIRECT_REF
- && DECL_P (TREE_OPERAND (decl, 0))
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == REFERENCE_TYPE))
- decl = TREE_OPERAND (decl, 0);
-
- STRIP_NOPS (decl);
- }
- else if ((region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
- && TREE_CODE (*pd) == COMPONENT_REF
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
- && code != OACC_UPDATE)
- {
- while (TREE_CODE (decl) == COMPONENT_REF)
- {
- decl = TREE_OPERAND (decl, 0);
- component_ref_p = true;
- if (((TREE_CODE (decl) == MEM_REF
- && integer_zerop (TREE_OPERAND (decl, 1)))
- || INDIRECT_REF_P (decl))
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == POINTER_TYPE))
- {
- indir_p = true;
- indir_base = decl;
- decl = TREE_OPERAND (decl, 0);
- STRIP_NOPS (decl);
- }
- if (TREE_CODE (decl) == INDIRECT_REF
- && DECL_P (TREE_OPERAND (decl, 0))
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == REFERENCE_TYPE))
- {
- decl_ref = decl;
- decl = TREE_OPERAND (decl, 0);
- }
- }
- }
- else if (TREE_CODE (decl) == COMPONENT_REF
- && (OMP_CLAUSE_MAP_KIND (c)
- != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
- {
- component_ref_p = true;
- while (TREE_CODE (decl) == COMPONENT_REF)
- decl = TREE_OPERAND (decl, 0);
- if (TREE_CODE (decl) == INDIRECT_REF
- && DECL_P (TREE_OPERAND (decl, 0))
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == REFERENCE_TYPE))
- decl = TREE_OPERAND (decl, 0);
- }
- if (decl != orig_decl && DECL_P (decl) && indir_p)
- {
- gomp_map_kind k
- = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
- ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
- /* We have a dereference of a struct member. Make this an
- attach/detach operation, and ensure the base pointer is
- mapped as a FIRSTPRIVATE_POINTER. */
- OMP_CLAUSE_SET_MAP_KIND (c, k);
- flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT;
- tree next_clause = OMP_CLAUSE_CHAIN (c);
- if (k == GOMP_MAP_ATTACH
- && code != OACC_ENTER_DATA
- && code != OMP_TARGET_ENTER_DATA
- && (!next_clause
- || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
- || (OMP_CLAUSE_MAP_KIND (next_clause)
- != GOMP_MAP_POINTER)
- || OMP_CLAUSE_DECL (next_clause) != decl)
- && (!struct_deref_set
- || !struct_deref_set->contains (decl))
- && (!struct_map_to_clause
- || !struct_map_to_clause->get (indir_base)))
- {
- if (!struct_deref_set)
- struct_deref_set = new hash_set<tree> ();
- /* As well as the attach, we also need a
- FIRSTPRIVATE_POINTER clause to properly map the
- pointer to the struct base. */
- tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
- OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2)
- = 1;
- tree charptr_zero
- = build_int_cst (build_pointer_type (char_type_node),
- 0);
- OMP_CLAUSE_DECL (c2)
- = build2 (MEM_REF, char_type_node,
- decl_ref ? decl_ref : decl, charptr_zero);
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
- tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c3,
- GOMP_MAP_FIRSTPRIVATE_POINTER);
- OMP_CLAUSE_DECL (c3) = decl;
- OMP_CLAUSE_SIZE (c3) = size_zero_node;
- tree mapgrp = *prev_list_p;
- *prev_list_p = c2;
- OMP_CLAUSE_CHAIN (c3) = mapgrp;
- OMP_CLAUSE_CHAIN (c2) = c3;
-
- struct_deref_set->add (decl);
- }
- goto do_add_decl;
- }
/* 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
@@ -10822,91 +10773,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
if (code == OACC_UPDATE
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
- if ((((region_type & ORT_ACC) && aggregate_base_p (decl))
- || (!(region_type & ORT_ACC)
- && (DECL_P (decl)
- || (component_ref_p
- && (INDIRECT_REF_P (decl)
- || TREE_CODE (decl) == MEM_REF
- || TREE_CODE (decl) == ARRAY_REF)))))
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
- && code != OACC_UPDATE
- && code != OMP_TARGET_UPDATE)
+
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
{
- if (error_operand_p (decl))
+ if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c)))
+ == ARRAY_TYPE)
+ remove = true;
+ else
{
- remove = true;
- break;
+ 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 stype = TREE_TYPE (decl);
- if (TREE_CODE (stype) == REFERENCE_TYPE)
- stype = TREE_TYPE (stype);
- if (TYPE_SIZE_UNIT (stype) == NULL
- || TREE_CODE (TYPE_SIZE_UNIT (stype)) != INTEGER_CST)
- {
- error_at (OMP_CLAUSE_LOCATION (c),
- "mapping field %qE of variable length "
- "structure", OMP_CLAUSE_DECL (c));
- remove = true;
- break;
- }
-
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- {
- /* Error recovery. */
- if (prev_list_p == NULL)
- {
- remove = true;
- break;
- }
-
- /* The below prev_list_p based error recovery code is
- currently no longer valid for OpenMP. */
- if (code != OMP_TARGET
- && code != OMP_TARGET_DATA
- && code != OMP_TARGET_UPDATE
- && code != OMP_TARGET_ENTER_DATA
- && code != OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_CHAIN (*prev_list_p) != c)
- {
- tree ch = OMP_CLAUSE_CHAIN (*prev_list_p);
- if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c)
- {
- remove = true;
- break;
- }
- }
- }
- bool cont = false;
- tree add_decl
- = build_struct_group (ctx, region_type, code, decl, pd,
- component_ref_p, &flags, c,
- struct_map_to_clause,
- struct_seen_clause, prev_list_p,
- list_p, pre_p, &cont);
- if (add_decl == error_mark_node)
- {
- remove = true;
- break;
- }
- else if (add_decl && DECL_P (add_decl))
- {
- decl = add_decl;
- goto do_add_decl;
- }
- if (cont)
- continue;
}
- else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+
+ 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)
{
- 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 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;
+ }
}
if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
@@ -11024,24 +10933,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
break;
}
- /* If this was of the form map(*pointer_to_struct), then the
- 'pointer_to_struct' DECL should be considered deref'ed. */
- if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
- || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c))
- || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c)))
- && INDIRECT_REF_P (orig_decl)
- && DECL_P (TREE_OPERAND (orig_decl, 0))
- && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE)
- {
- tree ptr = TREE_OPERAND (orig_decl, 0);
- if (!struct_deref_set || !struct_deref_set->contains (ptr))
- {
- if (!struct_deref_set)
- struct_deref_set = new hash_set<tree> ();
- struct_deref_set->add (ptr);
- }
- }
-
if (!remove
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
@@ -11058,28 +10949,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
break;
}
- else
- {
- /* DECL_P (decl) == true */
- tree *sc;
- if (struct_map_to_clause
- && (sc = struct_map_to_clause->get (decl)) != NULL
- && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT
- && decl == OMP_CLAUSE_DECL (*sc))
- {
- /* We have found a map of the whole structure after a
- leading GOMP_MAP_STRUCT has been created, so refill the
- leading clause into a map of the whole structure
- variable, and remove the current one.
- TODO: we should be able to remove some maps of the
- following structure element maps if they are of
- compatible TO/FROM/ALLOC type. */
- OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c));
- OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c));
- remove = true;
- break;
- }
- }
flags = GOVD_MAP | GOVD_EXPLICIT;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
@@ -11721,12 +11590,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
ctx->clauses = *orig_list_p;
gimplify_omp_ctxp = ctx;
- if (struct_seen_clause)
- delete struct_seen_clause;
- if (struct_map_to_clause)
- delete struct_map_to_clause;
- if (struct_deref_set)
- delete struct_deref_set;
}
/* Return true if DECL is a candidate for shared to firstprivate
@@ -11875,8 +11738,6 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void
*data)
return 0;
if ((flags & GOVD_SEEN) == 0)
return 0;
- if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0)
- return 0;
if (flags & GOVD_DEBUG_PRIVATE)
{
gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED);
diff --git a/gcc/testsuite/g++.dg/goacc/member-array-acc.C
b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
index e0c11570f5d..9993768ef20 100644
--- a/gcc/testsuite/g++.dg/goacc/member-array-acc.C
+++ b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
@@ -10,4 +10,4 @@ struct Foo {
};
int main() { Foo x; x.init(1024); }
-/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\)
map\(alloc:\(\(struct Foo \*\) this\)->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+
\[len: _[0-9]+\]\) map\(attach:\(\(struct Foo \*\) this\)->a \[bias: 0\]\)}
"gimple" } } */
+/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\)
map\(alloc:this->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\)
map\(attach:this->a \[bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C
b/gcc/testsuite/g++.dg/gomp/target-3.C
index f4d40ec8e4b..432f02614d8 100644
--- a/gcc/testsuite/g++.dg/gomp/target-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-3.C
@@ -33,4 +33,6 @@ T<N>::bar (int x)
template struct T<0>;
-/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\)
map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len:
\[0-9\]+\\\]\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct S \\*\\) this
\\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\)
map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct T \\*\\) this
\\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\)
map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index 150f286e312..bff7fa7c669 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,7 +87,7 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr
\[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\)
map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len:
[0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\)
map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len:
[0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\)
map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)
map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\)
map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } }
*/
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr
\[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\)
map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len:
[0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\)
map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this
\[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\)
map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\)
map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } }
*/
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\)
map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } }
*/
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C
b/gcc/testsuite/g++.dg/gomp/target-this-2.C
index 679c85a54dd..cc08e7e8693 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -46,4 +46,4 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\)
map\(firstprivate:__closure \[pointer assign, bias: 0\]\)
map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len:
[0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\)
map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" }
} */
+/* { dg-final { scan-tree-dump {map\(alloc:MEM\[\(char \*\)_[0-9]+\] \[len:
[0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)
firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\)
map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+
\[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias:
0\]\) map\(from:v \[len: [0-9]+\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C
b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index 8166f43ad42..9ade3cc0b2b 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -102,6 +102,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this
\[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\)
map\(firstprivate:__closure \[pointer assign, bias: 0\]\)
map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len:
[0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\)
map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)
map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } }
*/
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this
\[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\)
map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+
\[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias:
0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } }
*/
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this
\[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\)
map\(firstprivate:__closure \[pointer assign, bias: 0\]\)
map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len:
[0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\)
map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)
map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\)
map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this
\[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\)
map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+
\[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias:
0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)
map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\)
map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */
--
2.29.2
