* [PATCH 00/11] OpenMP: Deep struct dereferences
@ 2021-10-01 17:07 Julian Brown
2021-10-01 17:07 ` [PATCH 01/11] libgomp: Release device lock on cbuf error path Julian Brown
` (10 more replies)
0 siblings, 11 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
This is a series of patches to support deep struct dereferences for
OpenMP 5.0 (i.e. with multiple arrow operators,
"a->b[foo]->c[lo:hi]"). Apart from a couple of general bug fixes, the
main parts of this comprise:
1. Topological sorting of OMP clauses by base pointer dependencies.
2. Hoisting of struct sibling list handling out of
gimplify_scan_omp_clauses.
These patches replace and continue from the last part of the
previously-posted series:
https://gcc.gnu.org/pipermail/gcc-patches/2021-August/577219.html
and (still) depend on the parts of 1 through 7 of that series.
The patches have been bootstrapped & regression tested individually with
offloading to NVPTX (though some prior to the most recent rebase).
OK?
Thanks,
Julian
Julian Brown (11):
libgomp: Release device lock on cbuf error path
Remove base_ind/base_ref handling from extract_base_bit_offset
OpenMP 5.0: Clause ordering for OpenMP 5.0 (topological sorting by
base pointer)
Remove omp_target_reorder_clauses
OpenMP/OpenACC: Hoist struct sibling list handling in gimplification
OpenMP: Allow array ref components for C & C++
OpenMP: Fix non-zero attach/detach bias for struct dereferences
Not for committing: noisy topological sorting output
Not for committing: noisy sibling-list handling output
Not for committing: noisy mapping-group taxonomy
OpenMP/OpenACC: [WIP] Add gcc_unreachable to apparently-dead path in
build_struct_comp_nodes
gcc/c-family/c-common.h | 1 +
gcc/c-family/c-omp.c | 42 +
gcc/c/c-typeck.c | 15 +-
gcc/cp/semantics.c | 17 +-
gcc/gimplify.c | 2479 ++++++++++++-----
gcc/omp-low.c | 7 +-
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 | 6 +-
gcc/testsuite/g++.dg/gomp/target-this-2.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-this-3.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +-
libgomp/target.c | 5 +-
libgomp/testsuite/libgomp.c++/baseptrs-3.C | 182 ++
.../libgomp.c-c++-common/baseptrs-1.c | 50 +
.../libgomp.c-c++-common/baseptrs-2.c | 70 +
16 files changed, 2151 insertions(+), 739 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c++/baseptrs-3.C
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 01/11] libgomp: Release device lock on cbuf error path
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
@ 2021-10-01 17:07 ` Julian Brown
2021-10-12 9:23 ` Jakub Jelinek
2021-10-01 17:07 ` [PATCH 02/11] Remove base_ind/base_ref handling from extract_base_bit_offset Julian Brown
` (9 subsequent siblings)
10 siblings, 1 reply; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
This patch releases the device lock on a sanity-checking error path in
transfer combining (cbuf) handling in libgomp:target.c. This shouldn't
happen when handling well-formed mapping clauses, but erroneous clauses
can currently cause a hang if the condition triggers.
Tested with offloading to NVPTX. OK?
2021-09-29 Julian Brown <julian@codesourcery.com>
libgomp/
* target.c (gomp_copy_host2dev): Release device lock on cbuf
error path.
---
libgomp/target.c | 5 ++++-
1 file changed, 4 insertions(+), 1 deletion(-)
diff --git a/libgomp/target.c b/libgomp/target.c
index 65bb40100e5..84c6fdf2c47 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -385,7 +385,10 @@ gomp_copy_host2dev (struct gomp_device_descr *devicep,
else if (cbuf->chunks[middle].start <= doff)
{
if (doff + sz > cbuf->chunks[middle].end)
- gomp_fatal ("internal libgomp cbuf error");
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("internal libgomp cbuf error");
+ }
memcpy ((char *) cbuf->buf + (doff - cbuf->chunks[0].start),
h, sz);
return;
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 02/11] Remove base_ind/base_ref handling from extract_base_bit_offset
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
2021-10-01 17:07 ` [PATCH 01/11] libgomp: Release device lock on cbuf error path Julian Brown
@ 2021-10-01 17:07 ` Julian Brown
2021-10-12 9:27 ` Jakub Jelinek
2021-10-01 17:07 ` [PATCH 03/11] OpenMP 5.0: Clause ordering for OpenMP 5.0 (topological sorting by base pointer) Julian Brown
` (8 subsequent siblings)
10 siblings, 1 reply; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
In preparation for follow-up patches extending struct dereference
handling for OpenMP, this patch removes base_ind/base_ref handling from
gimplify.c:extract_base_bit_offset. This arguably simplifies some of the
code around the callers of the function also, though subsequent patches
modify those parts further.
OK for mainline?
Thanks,
Julian
2021-09-29 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (extract_base_bit_offset): Remove BASE_IND, BASE_REF and
OPENMP parameters.
(strip_indirections): New function.
(build_struct_group): Update calls to extract_base_bit_offset.
Rearrange indirect/reference handling accordingly. Use extracted base
instead of passed-in decl when grouping component accesses together.
---
gcc/gimplify.c | 109 ++++++++++++++++++++++++++-----------------------
1 file changed, 57 insertions(+), 52 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 92f8a7b4073..ece22b7a4ae 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8641,9 +8641,8 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
has array type, else return NULL. */
static tree
-extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref,
- poly_int64 *bitposp, poly_offset_int *poffsetp,
- tree *offsetp, bool openmp)
+extract_base_bit_offset (tree base, poly_int64 *bitposp,
+ poly_offset_int *poffsetp, tree *offsetp)
{
tree offset;
poly_int64 bitsize, bitpos;
@@ -8651,38 +8650,12 @@ extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref,
int unsignedp, reversep, volatilep = 0;
poly_offset_int poffset;
- if (base_ind)
- *base_ind = NULL_TREE;
-
- if (base_ref)
- *base_ref = NULL_TREE;
+ STRIP_NOPS (base);
base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode,
&unsignedp, &reversep, &volatilep);
- if (!openmp
- && (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)
- {
- if (base_ind)
- *base_ind = base;
- base = TREE_OPERAND (base, 0);
- }
- if ((TREE_CODE (base) == INDIRECT_REF
- || (TREE_CODE (base) == MEM_REF
- && integer_zerop (TREE_OPERAND (base, 1))))
- && DECL_P (TREE_OPERAND (base, 0))
- && TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE)
- {
- if (base_ref)
- *base_ref = base;
- base = TREE_OPERAND (base, 0);
- }
-
- if (!openmp)
- STRIP_NOPS (base);
+ STRIP_NOPS (base);
if (offset && poly_int_tree_p (offset))
{
@@ -8739,6 +8712,17 @@ strip_components_and_deref (tree expr)
return expr;
}
+static tree
+strip_indirections (tree expr)
+{
+ while (TREE_CODE (expr) == INDIRECT_REF
+ || (TREE_CODE (expr) == MEM_REF
+ && integer_zerop (TREE_OPERAND (expr, 1))))
+ expr = TREE_OPERAND (expr, 0);
+
+ return expr;
+}
+
/* Return TRUE if EXPR is something we will use as the base of an aggregate
access, either:
@@ -9232,7 +9216,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
{
poly_offset_int coffset;
poly_int64 cbitpos;
- tree base_ind, base_ref, tree_coffset;
+ tree tree_coffset;
tree ocd = OMP_CLAUSE_DECL (c);
bool openmp = !(region_type & ORT_ACC);
@@ -9242,10 +9226,25 @@ 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, &base_ind, &base_ref, &cbitpos,
- &coffset, &tree_coffset, openmp);
+ tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset, &tree_coffset);
+ tree sbase;
- bool do_map_struct = (base == decl && !tree_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);
+
+ 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"
@@ -9305,19 +9304,12 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
OMP_CLAUSE_SET_MAP_KIND (l, k);
- if (!openmp && base_ind)
- OMP_CLAUSE_DECL (l) = unshare_expr (base_ind);
- else if (base_ref)
- OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
- else
- {
- OMP_CLAUSE_DECL (l) = unshare_expr (decl);
- 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_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))
@@ -9353,6 +9345,20 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
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
@@ -9485,10 +9491,9 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
== REFERENCE_TYPE))
sc_decl = TREE_OPERAND (sc_decl, 0);
- tree base = extract_base_bit_offset (sc_decl, NULL, NULL,
- &bitpos, &offset,
- &tree_offset, openmp);
- if (!base || !operand_equal_p (base, decl, 0))
+ tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset,
+ &tree_offset);
+ if (!base2 || !operand_equal_p (base2, base, 0))
break;
if (scp)
continue;
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 03/11] OpenMP 5.0: Clause ordering for OpenMP 5.0 (topological sorting by base pointer)
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
2021-10-01 17:07 ` [PATCH 01/11] libgomp: Release device lock on cbuf error path Julian Brown
2021-10-01 17:07 ` [PATCH 02/11] Remove base_ind/base_ref handling from extract_base_bit_offset Julian Brown
@ 2021-10-01 17:07 ` Julian Brown
2021-10-01 17:07 ` [PATCH 04/11] Remove omp_target_reorder_clauses Julian Brown
` (7 subsequent siblings)
10 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
This patch reimplements the omp_target_reorder_clauses function in
anticipation of supporting "deeper" struct mappings (that is, with
several structure dereference operators, or similar).
The idea is that in place of the (possibly quadratic) algorithm in
omp_target_reorder_clauses that greedily moves clauses containing
addresses that are subexpressions of other addresses before those other
addresses, we employ a topological sort algorithm to calculate a proper
order for map clauses. This should run in linear time, and hopefully
handles degenerate cases where multiple "levels" of indirect accesses
are present on a given directive.
The new method also takes care to keep clause groups together, addressing
the concerns raised in:
https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570501.html
To figure out if some given clause depends on a base pointer in another
clause, we strip off the outer layers of the address expression, and check
(via a tree_operand_hash hash table we have built) if the result is a
"base pointer" as defined in OpenMP 5.0 (1.2.6 Data Terminology). There
are some subtleties involved, however:
- We must treat MEM_REF with zero offset the same as INDIRECT_REF.
This should probably be fixed in the front ends instead so we always
use a canonical form (probably INDIRECT_REF). The following patch
shows one instance of the problem, but there may be others:
https://gcc.gnu.org/pipermail/gcc-patches/2021-May/571382.html
- Mapping a whole struct implies mapping each of that struct's
elements, which may be base pointers. Because those base pointers
aren't necessarily explicitly referenced in the directive in question,
we treat the whole-struct mapping as a dependency instead.
This version of the patch is significantly improved over the version
posted previously in order to support the subsequent patches in this
series.
OK for mainline?
Thanks,
Julian
2021-09-29 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (is_or_contains_p, omp_target_reorder_clauses): Delete
functions.
(omp_tsort_mark): Add enum.
(omp_mapping_group): Add struct.
(omp_get_base_pointer, omp_get_attachment, omp_group_last,
omp_gather_mapping_groups, omp_group_base,
omp_index_mapping_groups, omp_containing_struct,
omp_tsort_mapping_groups_1,
omp_tsort_mapping_groups, omp_segregate_mapping_groups,
omp_reorder_mapping_groups): New functions.
(gimplify_scan_omp_clauses): Call above functions instead of
omp_target_reorder_clauses, unless we've seen an error.
* omp-low.c (scan_sharing_clauses): Avoid strict test if we haven't
sorted mapping groups.
gcc/testsuite/
* g++.dg/gomp/target-lambda-1.C: Adjust expected output.
* g++.dg/gomp/target-this-3.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
---
gcc/gimplify.c | 804 +++++++++++++++++++-
gcc/omp-low.c | 7 +-
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 6 +-
gcc/testsuite/g++.dg/gomp/target-this-3.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +-
5 files changed, 788 insertions(+), 37 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index ece22b7a4ae..d3346fc8d35 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8675,29 +8675,6 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp,
return base;
}
-/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR. */
-
-static bool
-is_or_contains_p (tree expr, tree base_ptr)
-{
- if ((TREE_CODE (expr) == INDIRECT_REF && TREE_CODE (base_ptr) == MEM_REF)
- || (TREE_CODE (expr) == MEM_REF && TREE_CODE (base_ptr) == INDIRECT_REF))
- return operand_equal_p (TREE_OPERAND (expr, 0),
- TREE_OPERAND (base_ptr, 0));
- while (!operand_equal_p (expr, base_ptr))
- {
- if (TREE_CODE (base_ptr) == COMPOUND_EXPR)
- base_ptr = TREE_OPERAND (base_ptr, 1);
- if (TREE_CODE (base_ptr) == COMPONENT_REF
- || TREE_CODE (base_ptr) == POINTER_PLUS_EXPR
- || TREE_CODE (base_ptr) == SAVE_EXPR)
- base_ptr = TREE_OPERAND (base_ptr, 0);
- else
- break;
- }
- return operand_equal_p (expr, base_ptr);
-}
-
/* Remove COMPONENT_REFS and indirections from EXPR. */
static tree
@@ -8751,6 +8728,7 @@ aggregate_base_p (tree expr)
return false;
}
+#if 0
/* Implement OpenMP 5.x map ordering rules for target directives. There are
several rules, and with some level of ambiguity, hopefully we can at least
collect the complexity here in one place. */
@@ -8930,6 +8908,758 @@ omp_target_reorder_clauses (tree *list_p)
}
}
}
+#endif
+
+
+enum omp_tsort_mark {
+ UNVISITED,
+ TEMPORARY,
+ PERMANENT
+};
+
+struct omp_mapping_group {
+ tree *grp_start;
+ tree grp_end;
+ omp_tsort_mark mark;
+ struct omp_mapping_group *sibling;
+ struct omp_mapping_group *next;
+};
+
+__attribute__((used)) static void
+debug_mapping_group (omp_mapping_group *grp)
+{
+ tree tmp = OMP_CLAUSE_CHAIN (grp->grp_end);
+ OMP_CLAUSE_CHAIN (grp->grp_end) = NULL;
+ debug_generic_expr (*grp->grp_start);
+ OMP_CLAUSE_CHAIN (grp->grp_end) = tmp;
+}
+
+/* Return the OpenMP "base pointer" of an expression EXPR, or NULL if there
+ isn't one. This needs improvement. */
+
+static tree
+omp_get_base_pointer (tree expr)
+{
+ while (TREE_CODE (expr) == ARRAY_REF)
+ expr = TREE_OPERAND (expr, 0);
+
+ while (TREE_CODE (expr) == COMPONENT_REF
+ && (DECL_P (TREE_OPERAND (expr, 0))
+ || (TREE_CODE (TREE_OPERAND (expr, 0)) == 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)))
+ || TREE_CODE (TREE_OPERAND (expr, 0)) == ARRAY_REF))
+ {
+ expr = TREE_OPERAND (expr, 0);
+
+ while (TREE_CODE (expr) == ARRAY_REF)
+ expr = TREE_OPERAND (expr, 0);
+
+ if (TREE_CODE (expr) == INDIRECT_REF || TREE_CODE (expr) == MEM_REF)
+ break;
+ }
+
+ if (DECL_P (expr))
+ return NULL_TREE;
+
+ if (TREE_CODE (expr) == INDIRECT_REF
+ || TREE_CODE (expr) == MEM_REF)
+ {
+ expr = TREE_OPERAND (expr, 0);
+ while (TREE_CODE (expr) == COMPOUND_EXPR)
+ expr = TREE_OPERAND (expr, 1);
+ if (TREE_CODE (expr) == POINTER_PLUS_EXPR)
+ expr = TREE_OPERAND (expr, 0);
+ if (TREE_CODE (expr) == SAVE_EXPR)
+ expr = TREE_OPERAND (expr, 0);
+ STRIP_NOPS (expr);
+ return expr;
+ }
+
+ return NULL_TREE;
+}
+
+/* An attach or detach operation depends directly on the address being
+ attached/detached. Return that address, or none if there are no
+ attachments/detachments. */
+
+static tree
+omp_get_attachment (omp_mapping_group *grp)
+{
+ tree node = *grp->grp_start;
+
+ switch (OMP_CLAUSE_MAP_KIND (node))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_TOFROM:
+ case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
+ case GOMP_MAP_FORCE_ALLOC:
+ if (node == grp->grp_end)
+ return NULL_TREE;
+
+ node = OMP_CLAUSE_CHAIN (node);
+ if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+ {
+ gcc_assert (node != grp->grp_end);
+ node = OMP_CLAUSE_CHAIN (node);
+ }
+ if (node)
+ switch (OMP_CLAUSE_MAP_KIND (node))
+ {
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ return NULL_TREE;
+
+ case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ return OMP_CLAUSE_DECL (node);
+
+ default:
+ internal_error ("unexpected mapping node");
+ }
+ return error_mark_node;
+
+ case GOMP_MAP_TO_PSET:
+ gcc_assert (node != grp->grp_end);
+ node = OMP_CLAUSE_CHAIN (node);
+ if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH)
+ return OMP_CLAUSE_DECL (node);
+ else
+ internal_error ("unexpected mapping node");
+ return error_mark_node;
+
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ node = OMP_CLAUSE_CHAIN (node);
+ if (!node || *grp->grp_start == grp->grp_end)
+ return OMP_CLAUSE_DECL (*grp->grp_start);
+ if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ return OMP_CLAUSE_DECL (*grp->grp_start);
+ else
+ internal_error ("unexpected mapping node");
+ return error_mark_node;
+
+ case GOMP_MAP_STRUCT:
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_LINK:
+ case GOMP_MAP_IF_PRESENT:
+ case GOMP_MAP_FIRSTPRIVATE:
+ case GOMP_MAP_FIRSTPRIVATE_INT:
+ case GOMP_MAP_USE_DEVICE_PTR:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ return NULL_TREE;
+
+ default:
+ internal_error ("unexpected mapping node");
+ }
+
+ return error_mark_node;
+}
+
+/* Given a pointer START_P to the start of a group of related (e.g. pointer)
+ mappings, return the chain pointer to the end of that group in the list. */
+
+static tree *
+omp_group_last (tree *start_p)
+{
+ tree c = *start_p, nc, *grp_last_p = start_p;
+
+ gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP);
+
+ nc = OMP_CLAUSE_CHAIN (c);
+
+ if (!nc || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP)
+ return grp_last_p;
+
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ default:
+ while (nc
+ && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH_DETACH
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_POINTER
+ || (OMP_CLAUSE_MAP_KIND (nc)
+ == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
+ || (OMP_CLAUSE_MAP_KIND (nc)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET))
+ {
+ grp_last_p = &OMP_CLAUSE_CHAIN (c);
+ c = nc;
+ tree nc2 = OMP_CLAUSE_CHAIN (nc);
+ if (nc2
+ && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc)
+ == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
+ && OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_ATTACH)
+ {
+ grp_last_p = &OMP_CLAUSE_CHAIN (nc);
+ c = nc2;
+ nc2 = OMP_CLAUSE_CHAIN (nc2);
+ }
+ nc = nc2;
+ }
+ break;
+
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ /* This is a weird artifact of how directives are parsed: bare attach or
+ detach clauses get a subsequent (meaningless) FIRSTPRIVATE_POINTER or
+ FIRSTPRIVATE_REFERENCE node. FIXME. */
+ if (nc
+ && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER))
+ grp_last_p = &OMP_CLAUSE_CHAIN (c);
+ break;
+
+ case GOMP_MAP_TO_PSET:
+ if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH))
+ grp_last_p = &OMP_CLAUSE_CHAIN (c);
+ break;
+ }
+
+ return grp_last_p;
+}
+
+/* Walk through LIST_P, and return a list of groups of mappings found (e.g.
+ OMP_CLAUSE_MAP with GOMP_MAP_{TO/FROM/TOFROM} followed by one or two
+ associated GOMP_MAP_POINTER mappings). Return a vector of omp_mapping_group
+ if we have more than one such group, else return NULL. */
+
+static vec<omp_mapping_group> *
+omp_gather_mapping_groups (tree *list_p)
+{
+ vec<omp_mapping_group> *groups = new vec<omp_mapping_group> ();
+
+ for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+ {
+ if (OMP_CLAUSE_CODE (*cp) != OMP_CLAUSE_MAP)
+ continue;
+
+ tree *grp_last_p = omp_group_last (cp);
+ omp_mapping_group grp;
+
+ grp.grp_start = cp;
+ grp.grp_end = *grp_last_p;
+ grp.mark = UNVISITED;
+ grp.sibling = NULL;
+ grp.next = NULL;
+ groups->safe_push (grp);
+
+ cp = grp_last_p;
+ }
+
+ if (groups->length () > 0)
+ return groups;
+ else
+ {
+ delete groups;
+ return NULL;
+ }
+}
+
+/* A pointer mapping group GRP may define a block of memory starting at some
+ base address, and maybe also define a firstprivate pointer or firstprivate
+ reference that points to that block. The return value is a node containing
+ the former, and the *FIRSTPRIVATE pointer is set if we have the latter.
+ If we define several base pointers, i.e. for a GOMP_MAP_STRUCT mapping,
+ return the number of consecutive chained nodes in CHAINED. */
+
+static tree
+omp_group_base (omp_mapping_group *grp, unsigned int *chained,
+ tree *firstprivate)
+{
+ tree node = *grp->grp_start;
+
+ *firstprivate = NULL_TREE;
+ *chained = 1;
+
+ switch (OMP_CLAUSE_MAP_KIND (node))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_TOFROM:
+ case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
+ case GOMP_MAP_FORCE_ALLOC:
+ if (node == grp->grp_end)
+ return node;
+
+ node = OMP_CLAUSE_CHAIN (node);
+ if (node && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_TO_PSET)
+ {
+ gcc_assert (node != grp->grp_end);
+ node = OMP_CLAUSE_CHAIN (node);
+ }
+ if (node)
+ switch (OMP_CLAUSE_MAP_KIND (node))
+ {
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ *firstprivate = OMP_CLAUSE_DECL (node);
+ return *grp->grp_start;
+
+ case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ return *grp->grp_start;
+
+ default:
+ internal_error ("unexpected mapping node");
+ }
+ else
+ internal_error ("unexpected mapping node");
+ return error_mark_node;
+
+ case GOMP_MAP_TO_PSET:
+ gcc_assert (node != grp->grp_end);
+ node = OMP_CLAUSE_CHAIN (node);
+ if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH)
+ return NULL_TREE;
+ else
+ internal_error ("unexpected mapping node");
+ return error_mark_node;
+
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ node = OMP_CLAUSE_CHAIN (node);
+ if (!node || *grp->grp_start == grp->grp_end)
+ return NULL_TREE;
+ if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ {
+ /* We're mapping the base pointer itself in a bare attach or detach
+ node. This is a side effect of how parsing works, and the mapping
+ will be removed anyway (at least for enter/exit data directives).
+ We should ignore the mapping here. FIXME. */
+ return NULL_TREE;
+ }
+ else
+ internal_error ("unexpected mapping node");
+ return error_mark_node;
+
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_LINK:
+ case GOMP_MAP_IF_PRESENT:
+ case GOMP_MAP_FIRSTPRIVATE:
+ case GOMP_MAP_FIRSTPRIVATE_INT:
+ case GOMP_MAP_USE_DEVICE_PTR:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ return NULL_TREE;
+
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ /* These shouldn't appear by themselves. */
+ if (!seen_error ())
+ internal_error ("unexpected pointer mapping node");
+ return error_mark_node;
+
+ default:
+ gcc_unreachable ();
+ }
+
+ return error_mark_node;
+}
+
+/* Given a vector of omp_mapping_groups, build a hash table so we can look up
+ nodes by tree_operand_hash. */
+
+static hash_map<tree_operand_hash, omp_mapping_group *> *
+omp_index_mapping_groups (vec<omp_mapping_group> *groups)
+{
+ hash_map<tree_operand_hash, omp_mapping_group *> *grpmap
+ = new hash_map<tree_operand_hash, omp_mapping_group *>;
+
+ omp_mapping_group *grp;
+ unsigned int i;
+
+ FOR_EACH_VEC_ELT (*groups, i, grp)
+ {
+ tree fpp;
+ unsigned int chained;
+ tree node = omp_group_base (grp, &chained, &fpp);
+
+ if (node == error_mark_node || (!node && !fpp))
+ continue;
+
+ for (unsigned j = 0;
+ node && j < chained;
+ node = OMP_CLAUSE_CHAIN (node), j++)
+ {
+ tree decl = OMP_CLAUSE_DECL (node);
+
+ /* Sometimes we see zero-offset MEM_REF instead of INDIRECT_REF,
+ meaning node-hash lookups don't work. This is a workaround for
+ that, but ideally we should just create the INDIRECT_REF at
+ source instead. FIXME. */
+ if (TREE_CODE (decl) == MEM_REF
+ && integer_zerop (TREE_OPERAND (decl, 1)))
+ decl = build1 (INDIRECT_REF, TREE_TYPE (decl),
+ TREE_OPERAND (decl, 0));
+
+ omp_mapping_group **prev = grpmap->get (decl);
+
+ if (prev && *prev == grp)
+ /* Empty. */;
+ else if (prev)
+ {
+ /* Mapping the same thing twice is normally diagnosed as an error,
+ but can happen under some circumstances, e.g. in pr99928-16.c,
+ the directive:
+
+ #pragma omp target simd reduction(+:a[:3]) \
+ map(always, tofrom: a[:6])
+ ...
+
+ will result in two "a[0]" mappings (of different sizes). */
+
+ grp->sibling = (*prev)->sibling;
+ (*prev)->sibling = grp;
+ }
+ else
+ grpmap->put (decl, grp);
+ }
+
+ if (!fpp)
+ continue;
+
+ omp_mapping_group **prev = grpmap->get (fpp);
+ if (prev)
+ {
+ grp->sibling = (*prev)->sibling;
+ (*prev)->sibling = grp;
+ }
+ else
+ grpmap->put (fpp, grp);
+ }
+ return grpmap;
+}
+
+/* Find the immediately-containing struct for a component ref (etc.)
+ expression EXPR. */
+
+static tree
+omp_containing_struct (tree expr)
+{
+ tree expr0 = expr;
+
+ STRIP_NOPS (expr);
+
+ tree expr1 = expr;
+
+ /* FIXME: other types of accessors? */
+ while (TREE_CODE (expr) == ARRAY_REF)
+ expr = TREE_OPERAND (expr, 0);
+
+ if (TREE_CODE (expr) == COMPONENT_REF)
+ {
+ if (DECL_P (TREE_OPERAND (expr, 0))
+ || TREE_CODE (TREE_OPERAND (expr, 0)) == 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)))
+ || TREE_CODE (TREE_OPERAND (expr, 0)) == ARRAY_REF)
+ expr = TREE_OPERAND (expr, 0);
+ else
+ internal_error ("unhandled component");
+ }
+
+ return (expr == expr1) ? expr0 : expr;
+}
+
+/* Helper function for omp_tsort_mapping_groups. Returns TRUE on success, or
+ FALSE on error. */
+
+static bool
+omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
+ vec<omp_mapping_group> *groups,
+ hash_map<tree_operand_hash, omp_mapping_group *>
+ *grpmap,
+ omp_mapping_group *grp)
+{
+ if (grp->mark == PERMANENT)
+ return true;
+ if (grp->mark == TEMPORARY)
+ {
+ fprintf (stderr, "when processing group:\n");
+ debug_mapping_group (grp);
+ internal_error ("base pointer cycle detected");
+ return false;
+ }
+ grp->mark = TEMPORARY;
+
+ tree attaches_to = omp_get_attachment (grp);
+
+ if (attaches_to)
+ {
+ omp_mapping_group **basep = grpmap->get (attaches_to);
+
+ if (basep)
+ {
+ gcc_assert (*basep != grp);
+ for (omp_mapping_group *w = *basep; w; w = w->sibling)
+ if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w))
+ return false;
+ }
+ }
+
+ tree decl = OMP_CLAUSE_DECL (*grp->grp_start);
+
+ while (decl)
+ {
+ tree base = omp_get_base_pointer (decl);
+
+ if (!base)
+ break;
+
+ omp_mapping_group **innerp = grpmap->get (base);
+
+ /* We should treat whole-structure mappings as if all (pointer, in this
+ case) members are mapped as individual list items. Check if we have
+ such a whole-structure mapping, if we don't have an explicit reference
+ to the pointer member itself. */
+ if (!innerp && TREE_CODE (base) == COMPONENT_REF)
+ {
+ base = omp_containing_struct (base);
+ innerp = grpmap->get (base);
+
+ if (!innerp
+ && TREE_CODE (base) == MEM_REF
+ && integer_zerop (TREE_OPERAND (base, 1)))
+ {
+ tree ind = TREE_OPERAND (base, 0);
+ ind = build1 (INDIRECT_REF, TREE_TYPE (base), ind);
+ innerp = grpmap->get (ind);
+ }
+ }
+
+ if (innerp && *innerp != grp)
+ {
+ for (omp_mapping_group *w = *innerp; w; w = w->sibling)
+ if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w))
+ return false;
+ break;
+ }
+
+ decl = base;
+ }
+
+ grp->mark = PERMANENT;
+
+ /* Emit grp to output list. */
+
+ **outlist = grp;
+ *outlist = &grp->next;
+
+ return true;
+}
+
+/* Topologically sort GROUPS, so that OMP 5.0-defined base pointers come
+ before mappings that use those pointers. This is an implementation of the
+ depth-first search algorithm, described e.g. at:
+
+ https://en.wikipedia.org/wiki/Topological_sorting
+*/
+
+static omp_mapping_group *
+omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
+ hash_map<tree_operand_hash, omp_mapping_group *>
+ *grpmap)
+{
+ omp_mapping_group *grp, *outlist = NULL, **cursor;
+ unsigned int i;
+
+ cursor = &outlist;
+
+ FOR_EACH_VEC_ELT (*groups, i, grp)
+ {
+ if (grp->mark != PERMANENT)
+ if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+ return NULL;
+ }
+
+ 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. */
+
+static omp_mapping_group *
+omp_segregate_mapping_groups (omp_mapping_group *inlist)
+{
+ omp_mapping_group *ard_groups = NULL, *tf_groups = NULL;
+ omp_mapping_group **ard_tail = &ard_groups, **tf_tail = &tf_groups;
+
+ for (omp_mapping_group *w = inlist; w;)
+ {
+ tree c = *w->grp_start;
+ omp_mapping_group *next = w->next;
+
+ gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP);
+
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
+ *ard_tail = w;
+ w->next = NULL;
+ ard_tail = &w->next;
+ break;
+
+ default:
+ *tf_tail = w;
+ w->next = NULL;
+ tf_tail = &w->next;
+ }
+
+ w = next;
+ }
+
+ /* Now splice the lists together... */
+ *tf_tail = ard_groups;
+
+ return tf_groups;
+}
+
+/* Given a list LIST_P containing groups of mappings given by GROUPS, reorder
+ those groups based on the output list of omp_tsort_mapping_groups --
+ singly-linked, threaded through each element's NEXT pointer starting at
+ HEAD. Each list element appears exactly once in that linked list.
+
+ Each element of GROUPS may correspond to one or several mapping nodes.
+ Node groups are kept together, and in the reordered list, the positions of
+ the original groups are reused for the positions of the reordered list.
+ Hence if we have e.g.
+
+ {to ptr ptr} firstprivate {tofrom ptr} ...
+ ^ ^ ^
+ first group non-"map" second group
+
+ and say the second group contains a base pointer for the first so must be
+ moved before it, the resulting list will contain:
+
+ {tofrom ptr} firstprivate {to ptr ptr} ...
+ ^ prev. second group ^ prev. first group
+*/
+
+static tree *
+omp_reorder_mapping_groups (vec<omp_mapping_group> *groups,
+ omp_mapping_group *head,
+ tree *list_p)
+{
+ omp_mapping_group *grp;
+ unsigned int i;
+ unsigned numgroups = groups->length ();
+ auto_vec<tree> old_heads (numgroups);
+ auto_vec<tree> new_heads (numgroups);
+ auto_vec<tree> old_succs (numgroups);
+ bool map_at_start = (list_p == (*groups)[0].grp_start);
+
+ tree *new_grp_tail = NULL;
+
+ /* Stash the start & end nodes of each mapping group before we start
+ modifying the list. */
+ FOR_EACH_VEC_ELT (*groups, i, grp)
+ {
+ old_heads.quick_push (*grp->grp_start);
+ old_succs.quick_push (OMP_CLAUSE_CHAIN (grp->grp_end));
+ }
+
+ /* And similarly, the heads of the groups in the order we want to rearrange
+ the list to. */
+ for (omp_mapping_group *w = head; w; w = w->next)
+ new_heads.quick_push (*w->grp_start);
+
+ FOR_EACH_VEC_ELT (*groups, i, grp)
+ {
+ gcc_assert (head);
+
+ if (new_grp_tail && old_succs[i - 1] == old_heads[i])
+ {
+ /* a {b c d} {e f g} h i j (original)
+ -->
+ a {k l m} {e f g} h i j (inserted new group on last iter)
+ -->
+ a {k l m} {n o p} h i j (this time, chain last group to new one)
+ ^new_grp_tail
+ */
+ *new_grp_tail = new_heads[i];
+ }
+ else if (new_grp_tail)
+ {
+ /* a {b c d} e {f g h} i j k (original)
+ -->
+ a {l m n} e {f g h} i j k (gap after last iter's group)
+ -->
+ a {l m n} e {o p q} h i j (chain last group to old successor)
+ ^new_grp_tail
+ */
+ *new_grp_tail = old_succs[i - 1];
+ }
+ else
+ {
+ /* The first inserted group -- point to new group, and leave end
+ open.
+ a {b c d} e f
+ -->
+ a {g h i...
+ */
+ *grp->grp_start = new_heads[i];
+ }
+
+ new_grp_tail = &OMP_CLAUSE_CHAIN (head->grp_end);
+
+ head = head->next;
+ }
+
+ if (new_grp_tail)
+ *new_grp_tail = old_succs[numgroups - 1];
+
+ gcc_assert (!head);
+
+ return map_at_start ? (*groups)[0].grp_start : list_p;
+}
/* DECL is supposed to have lastprivate semantics in the outer contexts
of combined/composite constructs, starting with OCTX.
@@ -9672,11 +10402,29 @@ 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)
- omp_target_reorder_clauses (list_p);
+ /* 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))
+ {
+ 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_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);
+ delete grpmap;
+ delete groups;
+ }
+ }
while ((c = *list_p) != NULL)
{
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index a63362c0511..d7d7f5310d0 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1530,8 +1530,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
/* If this is an offloaded region, an attach operation should
only exist when the pointer variable is mapped in a prior
- clause. */
- if (is_gimple_omp_offloaded (ctx->stmt))
+ clause.
+ If we had an error, we may not have attempted to sort clauses
+ properly, so avoid the test. */
+ if (is_gimple_omp_offloaded (ctx->stmt)
+ && !seen_error ())
gcc_assert
(maybe_lookup_decl (decl, ctx)
|| (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index e5a24d7abc4..150f286e312 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,8 +87,8 @@ 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\(attach_zero_length_array_section:__closure->__iptr \[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\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 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\(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\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 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" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 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\.__data2 \[bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
index 2755b4b58bd..bc2cc0b297d 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -100,6 +100,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate: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:this->refptr \[bias: 0\]\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index 3703762f45a..8166f43ad42 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\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 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: 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\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 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" } } */
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 04/11] Remove omp_target_reorder_clauses
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
` (2 preceding siblings ...)
2021-10-01 17:07 ` [PATCH 03/11] OpenMP 5.0: Clause ordering for OpenMP 5.0 (topological sorting by base pointer) Julian Brown
@ 2021-10-01 17:07 ` Julian Brown
2021-10-01 17:09 ` [PATCH 05/11] OpenMP/OpenACC: Hoist struct sibling list handling in gimplification Julian Brown
` (6 subsequent siblings)
10 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
This patch has been split out from the previous one to avoid a
confusingly-interleaved diff. The two patches should probably be
committed squashed together.
2021-10-01 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (omp_target_reorder_clauses): Delete.
---
gcc/gimplify.c | 183 -------------------------------------------------
1 file changed, 183 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index d3346fc8d35..c10a3e8842a 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8728,189 +8728,6 @@ aggregate_base_p (tree expr)
return false;
}
-#if 0
-/* Implement OpenMP 5.x map ordering rules for target directives. There are
- several rules, and with some level of ambiguity, hopefully we can at least
- collect the complexity here in one place. */
-
-static void
-omp_target_reorder_clauses (tree *list_p)
-{
- /* Collect refs to alloc/release/delete maps. */
- auto_vec<tree, 32> ard;
- tree *cp = list_p;
- while (*cp != NULL_TREE)
- if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC
- || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE
- || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE))
- {
- /* Unlink cp and push to ard. */
- tree c = *cp;
- tree nc = OMP_CLAUSE_CHAIN (c);
- *cp = nc;
- ard.safe_push (c);
-
- /* Any associated pointer type maps should also move along. */
- while (*cp != NULL_TREE
- && OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
- || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH
- || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER
- || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET))
- {
- c = *cp;
- nc = OMP_CLAUSE_CHAIN (c);
- *cp = nc;
- ard.safe_push (c);
- }
- }
- else
- cp = &OMP_CLAUSE_CHAIN (*cp);
-
- /* Link alloc/release/delete maps to the end of list. */
- for (unsigned int i = 0; i < ard.length (); i++)
- {
- *cp = ard[i];
- cp = &OMP_CLAUSE_CHAIN (ard[i]);
- }
- *cp = NULL_TREE;
-
- /* OpenMP 5.0 requires that pointer variables are mapped before
- its use as a base-pointer. */
- auto_vec<tree *, 32> atf;
- for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
- if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
- {
- /* Collect alloc, to, from, to/from clause tree pointers. */
- gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
- if (k == GOMP_MAP_ALLOC
- || k == GOMP_MAP_TO
- || k == GOMP_MAP_FROM
- || k == GOMP_MAP_TOFROM
- || k == GOMP_MAP_ALWAYS_TO
- || k == GOMP_MAP_ALWAYS_FROM
- || k == GOMP_MAP_ALWAYS_TOFROM)
- atf.safe_push (cp);
- }
-
- for (unsigned int i = 0; i < atf.length (); i++)
- if (atf[i])
- {
- tree *cp = atf[i];
- tree decl = OMP_CLAUSE_DECL (*cp);
- if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF)
- {
- tree base_ptr = TREE_OPERAND (decl, 0);
- STRIP_TYPE_NOPS (base_ptr);
- for (unsigned int j = i + 1; j < atf.length (); j++)
- if (atf[j])
- {
- tree *cp2 = atf[j];
- tree decl2 = OMP_CLAUSE_DECL (*cp2);
-
- decl2 = OMP_CLAUSE_DECL (*cp2);
- if (is_or_contains_p (decl2, base_ptr))
- {
- /* Move *cp2 to before *cp. */
- tree c = *cp2;
- *cp2 = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = *cp;
- *cp = c;
-
- if (*cp2 != NULL_TREE
- && OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (*cp2) == GOMP_MAP_ALWAYS_POINTER)
- {
- tree c2 = *cp2;
- *cp2 = OMP_CLAUSE_CHAIN (c2);
- OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = c2;
- }
-
- atf[j] = NULL;
- }
- }
- }
- }
-
- /* For attach_detach map clauses, if there is another map that maps the
- attached/detached pointer, make sure that map is ordered before the
- attach_detach. */
- atf.truncate (0);
- for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
- if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
- {
- /* Collect alloc, to, from, to/from clauses, and
- always_pointer/attach_detach clauses. */
- gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
- if (k == GOMP_MAP_ALLOC
- || k == GOMP_MAP_TO
- || k == GOMP_MAP_FROM
- || k == GOMP_MAP_TOFROM
- || k == GOMP_MAP_ALWAYS_TO
- || k == GOMP_MAP_ALWAYS_FROM
- || k == GOMP_MAP_ALWAYS_TOFROM
- || k == GOMP_MAP_ATTACH_DETACH
- || k == GOMP_MAP_ALWAYS_POINTER)
- atf.safe_push (cp);
- }
-
- for (unsigned int i = 0; i < atf.length (); i++)
- if (atf[i])
- {
- tree *cp = atf[i];
- tree ptr = OMP_CLAUSE_DECL (*cp);
- STRIP_TYPE_NOPS (ptr);
- if (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH)
- for (unsigned int j = i + 1; j < atf.length (); j++)
- {
- tree *cp2 = atf[j];
- tree decl2 = OMP_CLAUSE_DECL (*cp2);
- if (OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ATTACH_DETACH
- && OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ALWAYS_POINTER
- && is_or_contains_p (decl2, ptr))
- {
- /* Move *cp2 to before *cp. */
- tree c = *cp2;
- *cp2 = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = *cp;
- *cp = c;
- atf[j] = NULL;
-
- /* If decl2 is of the form '*decl2_opnd0', and followed by an
- ALWAYS_POINTER or ATTACH_DETACH of 'decl2_opnd0', move the
- pointer operation along with *cp2. This can happen for C++
- reference sequences. */
- if (j + 1 < atf.length ()
- && (TREE_CODE (decl2) == INDIRECT_REF
- || TREE_CODE (decl2) == MEM_REF))
- {
- tree *cp3 = atf[j + 1];
- tree decl3 = OMP_CLAUSE_DECL (*cp3);
- tree decl2_opnd0 = TREE_OPERAND (decl2, 0);
- if ((OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ALWAYS_POINTER
- || OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ATTACH_DETACH)
- && operand_equal_p (decl3, decl2_opnd0))
- {
- /* Also move *cp3 to before *cp. */
- c = *cp3;
- *cp2 = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = *cp;
- *cp = c;
- atf[j + 1] = NULL;
- j += 1;
- }
- }
- }
- }
- }
-}
-#endif
-
-
enum omp_tsort_mark {
UNVISITED,
TEMPORARY,
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 05/11] OpenMP/OpenACC: Hoist struct sibling list handling in gimplification
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
` (3 preceding siblings ...)
2021-10-01 17:07 ` [PATCH 04/11] Remove omp_target_reorder_clauses Julian Brown
@ 2021-10-01 17:09 ` Julian Brown
2021-10-01 17:16 ` Julian Brown
2021-10-01 17:09 ` [PATCH 06/11] OpenMP: Allow array ref components for C & C++ Julian Brown
` (5 subsequent siblings)
10 siblings, 1 reply; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:09 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
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 <julian@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
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 06/11] OpenMP: Allow array ref components for C & C++
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
` (4 preceding siblings ...)
2021-10-01 17:09 ` [PATCH 05/11] OpenMP/OpenACC: Hoist struct sibling list handling in gimplification Julian Brown
@ 2021-10-01 17:09 ` Julian Brown
2021-10-01 17:09 ` [PATCH 07/11] OpenMP: Fix non-zero attach/detach bias for struct dereferences Julian Brown
` (4 subsequent siblings)
10 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:09 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
This patch fixes parsing for struct components that are array references
in OMP clauses in both the C and C++ front ends.
OK for mainline?
Thanks,
Julian
2021-09-29 Julian Brown <julian@codesourcery.com>
gcc/c/
* c-typeck.c (c_finish_omp_clauses): Allow ARRAY_REF components.
gcc/cp/
* semantics.c (finish_omp_clauses): Allow ARRAY_REF components.
---
gcc/c/c-typeck.c | 3 ++-
gcc/cp/semantics.c | 3 ++-
2 files changed, 4 insertions(+), 2 deletions(-)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index e10e6aa8439..d0494cadf05 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14815,7 +14815,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF
- || TREE_CODE (t) == INDIRECT_REF)
+ || TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 6e954ca06a6..53bd8d236bb 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7849,7 +7849,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF
- || TREE_CODE (t) == INDIRECT_REF)
+ || TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 07/11] OpenMP: Fix non-zero attach/detach bias for struct dereferences
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
` (5 preceding siblings ...)
2021-10-01 17:09 ` [PATCH 06/11] OpenMP: Allow array ref components for C & C++ Julian Brown
@ 2021-10-01 17:09 ` Julian Brown
2021-10-11 14:49 ` Julian Brown
2021-10-01 17:09 ` [PATCH 08/11] Not for committing: noisy topological sorting output Julian Brown
` (3 subsequent siblings)
10 siblings, 1 reply; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:09 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
This patch fixes attach/detach operations for OpenMP that have a non-zero
bias: these can occur if we have a mapping such as:
#pragma omp target map(mystruct->a.b[idx].c[:arrsz])
i.e. where there is an offset between the attachment point ("mystruct"
here) and the pointed-to data. (The "b" and "c" members would be array
types here, not pointers themselves). In this example the difference
(thus bias encoded in the attach/detach node) will be something like:
(uintptr_t) &mystruct->a.b[idx].c[0] - (uintptr_t) &mystruct->a
OK for mainline?
Thanks,
Julian
2021-09-29 Julian Brown <julian@codesourcery.com>
gcc/c-family/
* c-common.h (c_omp_decompose_attachable_address): Add prototype.
* c-omp.c (c_omp_decompose_attachable_address): New function.
gcc/c/
* c-typeck.c (handle_omp_array_sections): Handle attach/detach for
struct dereferences with non-zero bias.
gcc/cp/
* semantics.c (handle_omp_array_section): Handle attach/detach for
struct dereferences with non-zero bias.
libgomp/
* testsuite/libgomp.c++/baseptrs-3.C: Add test (XFAILed for now).
* testsuite/libgomp.c-c++-common/baseptrs-1.c: Add test.
* testsuite/libgomp.c-c++-common/baseptrs-2.c: Add test.
---
gcc/c-family/c-common.h | 1 +
gcc/c-family/c-omp.c | 42 ++++
gcc/c/c-typeck.c | 12 +-
gcc/cp/semantics.c | 14 +-
libgomp/testsuite/libgomp.c++/baseptrs-3.C | 182 ++++++++++++++++++
.../libgomp.c-c++-common/baseptrs-1.c | 50 +++++
.../libgomp.c-c++-common/baseptrs-2.c | 70 +++++++
7 files changed, 364 insertions(+), 7 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c++/baseptrs-3.C
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 849cefab882..dab2dd33573 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1249,6 +1249,7 @@ extern tree c_omp_check_context_selector (location_t, tree);
extern void c_omp_mark_declare_variant (location_t, tree, tree);
extern const char *c_omp_map_clause_name (tree, bool);
extern void c_omp_adjust_map_clauses (tree, bool);
+extern tree c_omp_decompose_attachable_address (tree t, tree *virtbase);
enum c_omp_directive_kind {
C_OMP_DIR_STANDALONE,
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 1f07a0a454b..fc50f57e768 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -3119,6 +3119,48 @@ c_omp_adjust_map_clauses (tree clauses, bool is_target)
}
}
+tree
+c_omp_decompose_attachable_address (tree t, tree *virtbase)
+{
+ *virtbase = t;
+
+ /* It's already a pointer. Just use that. */
+ if (POINTER_TYPE_P (TREE_TYPE (t)))
+ return NULL_TREE;
+
+ /* Otherwise, look for a base pointer deeper within the expression. */
+
+ while (TREE_CODE (t) == COMPONENT_REF
+ && (TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF
+ || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
+ {
+ t = TREE_OPERAND (t, 0);
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ }
+
+
+ *virtbase = t;
+
+ if (TREE_CODE (t) != COMPONENT_REF)
+ return NULL_TREE;
+
+ t = TREE_OPERAND (t, 0);
+
+ tree attach_pt = NULL_TREE;
+
+ if ((TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == MEM_REF)
+ && TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == POINTER_TYPE)
+ {
+ attach_pt = TREE_OPERAND (t, 0);
+ if (TREE_CODE (attach_pt) == POINTER_PLUS_EXPR)
+ attach_pt = TREE_OPERAND (attach_pt, 0);
+ }
+
+ return attach_pt;
+}
+
static const struct c_omp_directive omp_directives[] = {
/* Keep this alphabetically sorted by the first word. Non-null second/third
if any should precede null ones. */
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index d0494cadf05..d1fd8be8e57 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13696,9 +13696,15 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
+ tree virtbase = t;
+ tree attach_pt
+ = ((ort != C_ORT_ACC)
+ ? c_omp_decompose_attachable_address (t, &virtbase)
+ : NULL_TREE);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (TREE_CODE (t) == COMPONENT_REF
- && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE
+ && !attach_pt))
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
switch (OMP_CLAUSE_MAP_KIND (c))
@@ -13731,10 +13737,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !c_mark_addressable (t))
return false;
- OMP_CLAUSE_DECL (c2) = t;
+ OMP_CLAUSE_DECL (c2) = attach_pt ? attach_pt : t;
t = build_fold_addr_expr (first);
t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t);
- tree ptr = OMP_CLAUSE_DECL (c2);
+ tree ptr = virtbase;
if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
ptr = build_fold_addr_expr (ptr);
t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 53bd8d236bb..a50ec0ad883 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5608,9 +5608,16 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_SIZE (c) = size;
if (TREE_CODE (t) == FIELD_DECL)
t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
+
+ tree virtbase = t;
+ tree attach_pt
+ = ((ort != C_ORT_ACC)
+ ? c_omp_decompose_attachable_address (t, &virtbase)
+ : NULL_TREE);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (TREE_CODE (t) == COMPONENT_REF
- && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE
+ && !attach_pt))
return false;
switch (OMP_CLAUSE_MAP_KIND (c))
{
@@ -5670,12 +5677,11 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !cxx_mark_addressable (t))
return false;
- OMP_CLAUSE_DECL (c2) = t;
+ OMP_CLAUSE_DECL (c2) = attach_pt ? attach_pt : t;
t = build_fold_addr_expr (first);
t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
ptrdiff_type_node, t);
- tree ptr = OMP_CLAUSE_DECL (c2);
- ptr = convert_from_reference (ptr);
+ tree ptr = convert_from_reference (virtbase);
if (!INDIRECT_TYPE_P (TREE_TYPE (ptr)))
ptr = build_fold_addr_expr (ptr);
t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
diff --git a/libgomp/testsuite/libgomp.c++/baseptrs-3.C b/libgomp/testsuite/libgomp.c++/baseptrs-3.C
new file mode 100644
index 00000000000..cabeb7c2b7a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/baseptrs-3.C
@@ -0,0 +1,182 @@
+/* { dg-xfail-if "fails to parse correctly" { *-*-* } } */
+
+#include <cstdlib>
+#include <cstring>
+#include <cassert>
+
+struct sa
+{
+ int *ptr;
+};
+
+struct sb
+{
+ int arr[10];
+};
+
+struct sc
+{
+ sa &a;
+ sb &b;
+ sc (sa &my_a, sb &my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foo ()
+{
+ sa my_a;
+ sb my_b;
+
+ my_a.ptr = (int *) malloc (sizeof (int) * 10);
+ sc my_c(my_a, my_b);
+
+ memset (my_c.a.ptr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_c.a.ptr, my_c.a.ptr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_c.a.ptr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_c.a.ptr[i] == i);
+
+ memset (my_c.b.arr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_c.b.arr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_c.b.arr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_c.b.arr[i] == i);
+
+ free (my_a.ptr);
+}
+
+void
+bar ()
+{
+ sa my_a;
+ sb my_b;
+
+ my_a.ptr = (int *) malloc (sizeof (int) * 10);
+ sc my_c(my_a, my_b);
+ sc &my_cref = my_c;
+
+ memset (my_cref.a.ptr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_cref.a.ptr, my_cref.a.ptr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_cref.a.ptr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_cref.a.ptr[i] == i);
+
+ memset (my_cref.b.arr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_cref.b.arr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_cref.b.arr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_cref.b.arr[i] == i);
+
+ free (my_a.ptr);
+}
+
+struct scp
+{
+ sa *&a;
+ sb *&b;
+ scp (sa *&my_a, sb *&my_b) : a(my_a), b(my_b) {}
+};
+
+void
+foop ()
+{
+ sa *my_a = new sa;
+ sb *my_b = new sb;
+
+ my_a->ptr = new int[10];
+ scp *my_c = new scp(my_a, my_b);
+
+ memset (my_c->a->ptr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_c->a->ptr, my_c->a->ptr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_c->a->ptr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_c->a->ptr[i] == i);
+
+ memset (my_c->b->arr, 0, sizeof (int) * 10);
+
+/* FIXME: This currently ICEs. */
+/* #pragma omp target map (my_c->b->arr[:10]) */
+ {
+ for (int i = 0; i < 10; i++)
+ my_c->b->arr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_c->b->arr[i] == i);
+
+ delete[] my_a->ptr;
+ delete my_a;
+ delete my_b;
+}
+
+void
+barp ()
+{
+ sa *my_a = new sa;
+ sb *my_b = new sb;
+
+ my_a->ptr = new int[10];
+ scp *my_c = new scp(my_a, my_b);
+ scp *&my_cref = my_c;
+
+ memset (my_cref->a->ptr, 0, sizeof (int) * 10);
+
+ #pragma omp target map (my_cref->a->ptr, my_cref->a->ptr[:10])
+ {
+ for (int i = 0; i < 10; i++)
+ my_cref->a->ptr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_cref->a->ptr[i] == i);
+
+ memset (my_cref->b->arr, 0, sizeof (int) * 10);
+
+/* FIXME: This currently ICEs. */
+/* #pragma omp target map (my_cref->b->arr[:10]) */
+ {
+ for (int i = 0; i < 10; i++)
+ my_cref->b->arr[i] = i;
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (my_cref->b->arr[i] == i);
+
+ delete my_a->ptr;
+ delete my_a;
+ delete my_b;
+}
+
+int main (int argc, char *argv[])
+{
+ foo ();
+ bar ();
+ foop ();
+ barp ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c
new file mode 100644
index 00000000000..073615625b7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c
@@ -0,0 +1,50 @@
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+#include <stdio.h>
+
+#define N 32
+
+typedef struct {
+ int x2[10][N];
+} x1type;
+
+typedef struct {
+ x1type x1[10];
+} p2type;
+
+typedef struct {
+ p2type *p2;
+} p1type;
+
+typedef struct {
+ p1type *p1;
+} x0type;
+
+typedef struct {
+ x0type x0[10];
+} p0type;
+
+int main(int argc, char *argv[])
+{
+ p0type *p0;
+ int k1 = 0, k2 = 0, k3 = 0, n = N;
+
+ p0 = (p0type *) malloc (sizeof *p0);
+ p0->x0[0].p1 = (p1type *) malloc (sizeof *p0->x0[0].p1);
+ p0->x0[0].p1->p2 = (p2type *) malloc (sizeof *p0->x0[0].p1->p2);
+ memset (p0->x0[0].p1->p2, 0, sizeof *p0->x0[0].p1->p2);
+
+#pragma omp target map(tofrom: p0->x0[k1].p1->p2[k2].x1[k3].x2[4][0:n]) \
+ map(to: p0->x0[k1].p1, p0->x0[k1].p1->p2) \
+ map(to: p0->x0[k1].p1[0])
+ {
+ for (int i = 0; i < n; i++)
+ p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i] = i;
+ }
+
+ for (int i = 0; i < n; i++)
+ assert (i == p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i]);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c
new file mode 100644
index 00000000000..e335d7da966
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c
@@ -0,0 +1,70 @@
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+
+#define N 32
+
+typedef struct {
+ int arr[N];
+ int *ptr;
+} sc;
+
+typedef struct {
+ sc *c;
+} sb;
+
+typedef struct {
+ sb *b;
+ sc *c;
+} sa;
+
+int main (int argc, char *argv[])
+{
+ sa *p;
+
+ p = (sa *) malloc (sizeof *p);
+ p->b = (sb *) malloc (sizeof *p->b);
+ p->b->c = (sc *) malloc (sizeof *p->b->c);
+ p->c = (sc *) malloc (sizeof *p->c);
+ p->b->c->ptr = (int *) malloc (N * sizeof (int));
+ p->c->ptr = (int *) malloc (N * sizeof (int));
+
+ for (int i = 0; i < N; i++)
+ {
+ p->b->c->ptr[i] = 0;
+ p->c->ptr[i] = 0;
+ p->b->c->arr[i] = 0;
+ p->c->arr[i] = 0;
+ }
+
+#pragma omp target map(to: p->b, p->b[0], p->c, p->c[0], p->b->c, p->b->c[0]) \
+ map(to: p->b->c->ptr, p->c->ptr) \
+ map(tofrom: p->b->c->ptr[:N], p->c->ptr[:N])
+ {
+ for (int i = 0; i < N; i++)
+ {
+ p->b->c->ptr[i] = i;
+ p->c->ptr[i] = i * 2;
+ }
+ }
+
+#pragma omp target map(to: p->b, p->b[0], p->b->c, p->c) \
+ map(tofrom: p->c[0], p->b->c[0])
+ {
+ for (int i = 0; i < N; i++)
+ {
+ p->b->c->arr[i] = i * 3;
+ p->c->arr[i] = i * 4;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (p->b->c->ptr[i] == i);
+ assert (p->c->ptr[i] == i * 2);
+ assert (p->b->c->arr[i] == i * 3);
+ assert (p->c->arr[i] == i * 4);
+ }
+
+ return 0;
+}
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 08/11] Not for committing: noisy topological sorting output
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
` (6 preceding siblings ...)
2021-10-01 17:09 ` [PATCH 07/11] OpenMP: Fix non-zero attach/detach bias for struct dereferences Julian Brown
@ 2021-10-01 17:09 ` Julian Brown
2021-10-01 17:09 ` [PATCH 09/11] Not for committing: noisy sibling-list handling output Julian Brown
` (2 subsequent siblings)
10 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:09 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
As a possible aid to review, this is my "printf-style" debugging cruft for
the topological sorting implementation.
We might want to rework this into something that emits scannable output
into the gimple dump in order to write tests to make sure base pointer
dependencies are being found properly, but that hasn't been done yet.
This is not for committing.
---
gcc/gimplify.c | 169 ++++++++++++++++++++++++++++++++++++++++++++++---
1 file changed, 161 insertions(+), 8 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 31e2e4d9fe7..2ec83bf273b 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -70,6 +70,8 @@ along with GCC; see the file COPYING3. If not see
#include "context.h"
#include "tree-nested.h"
+//#define NOISY_TOPOSORT
+
/* Hash set of poisoned variables in a bind expr. */
static hash_set<tree> *asan_poisoned_variables = NULL;
@@ -8957,6 +8959,10 @@ omp_gather_mapping_groups (tree *list_p)
{
vec<omp_mapping_group> *groups = new vec<omp_mapping_group> ();
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "GATHER MAPPING GROUPS\n");
+#endif
+
for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
{
if (OMP_CLAUSE_CODE (*cp) != OMP_CLAUSE_MAP)
@@ -8965,6 +8971,25 @@ omp_gather_mapping_groups (tree *list_p)
tree *grp_last_p = omp_group_last (cp);
omp_mapping_group grp;
+#ifdef NOISY_TOPOSORT
+ if (cp == grp_last_p)
+ {
+ tree tmp = OMP_CLAUSE_CHAIN (*cp);
+ OMP_CLAUSE_CHAIN (*cp) = NULL_TREE;
+ fprintf (stderr, "found singleton clause:\n");
+ debug_generic_expr (*cp);
+ OMP_CLAUSE_CHAIN (*cp) = tmp;
+ }
+ else
+ {
+ tree tmp = OMP_CLAUSE_CHAIN (*grp_last_p);
+ OMP_CLAUSE_CHAIN (*grp_last_p) = NULL_TREE;
+ fprintf (stderr, "found group:\n");
+ debug_generic_expr (*cp);
+ OMP_CLAUSE_CHAIN (*grp_last_p) = tmp;
+ }
+#endif
+
grp.grp_start = cp;
grp.grp_end = *grp_last_p;
grp.mark = UNVISITED;
@@ -9129,14 +9154,44 @@ omp_index_mapping_groups (vec<omp_mapping_group> *groups)
omp_mapping_group *grp;
unsigned int i;
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "INDEX MAPPING GROUPS\n");
+#endif
+
FOR_EACH_VEC_ELT (*groups, i, grp)
{
+#ifdef NOISY_TOPOSORT
+ debug_mapping_group (grp);
+#endif
+
tree fpp;
unsigned int chained;
tree node = omp_group_base (grp, &chained, &fpp);
if (node == error_mark_node || (!node && !fpp))
- continue;
+ {
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, " -- NULL base, not indexing.\n");
+#endif
+ continue;
+ }
+
+#ifdef NOISY_TOPOSORT
+ if (node)
+ {
+ fprintf (stderr, "base%s: ", chained > 1 ? " list" : "");
+
+ tree walk = node;
+ for (unsigned j = 0; j < chained; walk = OMP_CLAUSE_CHAIN (walk), j++)
+ debug_generic_expr (OMP_CLAUSE_DECL (walk));
+ }
+
+ if (fpp)
+ {
+ fprintf (stderr, "firstprivate pointer/reference: ");
+ debug_generic_expr (fpp);
+ }
+#endif
for (unsigned j = 0;
node && j < chained;
@@ -9156,7 +9211,11 @@ omp_index_mapping_groups (vec<omp_mapping_group> *groups)
omp_mapping_group **prev = grpmap->get (decl);
if (prev && *prev == grp)
- /* Empty. */;
+ {
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, " -- same node\n");
+#endif
+ }
else if (prev)
{
/* Mapping the same thing twice is normally diagnosed as an error,
@@ -9171,9 +9230,17 @@ omp_index_mapping_groups (vec<omp_mapping_group> *groups)
grp->sibling = (*prev)->sibling;
(*prev)->sibling = grp;
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, " -- index as sibling\n");
+#endif
}
else
- grpmap->put (decl, grp);
+ {
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, " -- index as new decl\n");
+#endif
+ grpmap->put (decl, grp);
+ }
}
if (!fpp)
@@ -9184,9 +9251,17 @@ omp_index_mapping_groups (vec<omp_mapping_group> *groups)
{
grp->sibling = (*prev)->sibling;
(*prev)->sibling = grp;
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, " -- index fpp as sibling\n");
+#endif
}
else
- grpmap->put (fpp, grp);
+ {
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, " -- index fpp as new decl\n");
+#endif
+ grpmap->put (fpp, grp);
+ }
}
return grpmap;
}
@@ -9233,6 +9308,11 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
*grpmap,
omp_mapping_group *grp)
{
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "processing node/group:\n");
+ debug_mapping_group (grp);
+#endif
+
if (grp->mark == PERMANENT)
return true;
if (grp->mark == TEMPORARY)
@@ -9253,11 +9333,26 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
if (basep)
{
gcc_assert (*basep != grp);
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "has attachment to:\n");
+ debug_mapping_group (*basep);
+#endif
for (omp_mapping_group *w = *basep; w; w = w->sibling)
if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w))
return false;
}
+#ifdef NOISY_TOPOSORT
+ else
+ {
+ fprintf (stderr, "can't find base for attachment, tried:\n");
+ debug_generic_expr (attaches_to);
+ }
+#endif
}
+#ifdef NOISY_TOPOSORT
+ else
+ fprintf (stderr, "doesn't attach to anything\n");
+#endif
tree decl = OMP_CLAUSE_DECL (*grp->grp_start);
@@ -9266,7 +9361,13 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
tree base = omp_get_base_pointer (decl);
if (!base)
- break;
+ {
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "no base pointer for decl:\n");
+ debug_generic_expr (decl);
+#endif
+ break;
+ }
omp_mapping_group **innerp = grpmap->get (base);
@@ -9283,6 +9384,9 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
&& TREE_CODE (base) == MEM_REF
&& integer_zerop (TREE_OPERAND (base, 1)))
{
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "we have a mem_ref, retry as indirect_ref\n");
+#endif
tree ind = TREE_OPERAND (base, 0);
ind = build1 (INDIRECT_REF, TREE_TYPE (base), ind);
innerp = grpmap->get (ind);
@@ -9291,12 +9395,22 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
if (innerp && *innerp != grp)
{
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "base pointer of node is mapped too:\n");
+ debug_generic_expr (base);
+#endif
for (omp_mapping_group *w = *innerp; w; w = w->sibling)
if (!omp_tsort_mapping_groups_1 (outlist, groups, grpmap, w))
return false;
break;
}
-
+#ifdef NOISY_TOPOSORT
+ else
+ {
+ fprintf (stderr, "base pointer of node is not mapped:\n");
+ debug_generic_expr (base);
+ }
+#endif
decl = base;
}
@@ -9307,6 +9421,16 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
**outlist = grp;
*outlist = &grp->next;
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "+++ EMIT NODE/GROUP:\n");
+ {
+ tree tmp = OMP_CLAUSE_CHAIN (grp->grp_end);
+ OMP_CLAUSE_CHAIN (grp->grp_end) = NULL_TREE;
+ debug_generic_expr (*grp->grp_start);
+ OMP_CLAUSE_CHAIN (grp->grp_end) = tmp;
+ }
+#endif
+
return true;
}
@@ -9327,6 +9451,10 @@ omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
cursor = &outlist;
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "TOPOLOGICALLY SORT MAPPING GROUPS\n");
+#endif
+
FOR_EACH_VEC_ELT (*groups, i, grp)
{
if (grp->mark != PERMANENT)
@@ -9348,6 +9476,10 @@ omp_segregate_mapping_groups (omp_mapping_group *inlist)
omp_mapping_group *ard_groups = NULL, *tf_groups = NULL;
omp_mapping_group **ard_tail = &ard_groups, **tf_tail = &tf_groups;
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "SEGREGATE MAPPING GROUPS\n");
+#endif
+
for (omp_mapping_group *w = inlist; w;)
{
tree c = *w->grp_start;
@@ -9452,7 +9584,7 @@ omp_reorder_mapping_groups (vec<omp_mapping_group> *groups,
-->
a {l m n} e {o p q} h i j (chain last group to old successor)
^new_grp_tail
- */
+ */
*new_grp_tail = old_succs[i - 1];
}
else
@@ -10285,9 +10417,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
grpmap = omp_index_mapping_groups (groups);
outlist = omp_tsort_mapping_groups (groups, grpmap);
+#ifdef NOISY_TOPOSORT
+ if (!outlist)
+ goto failure;
+ fprintf (stderr, "TOPO-SORTED ORDER:\n");
+ for (omp_mapping_group *w = outlist; w; w = w->next)
+ debug_mapping_group (w);
+#endif
outlist = omp_segregate_mapping_groups (outlist);
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "FINAL ORDER:\n");
+ for (omp_mapping_group *w = outlist; w; w = w->next)
+ debug_mapping_group (w);
+#endif
list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
-
+#ifdef NOISY_TOPOSORT
+ fprintf (stderr, "REORDERED CLAUSES:\n");
+ for (tree *w = list_p; *w; w = &OMP_CLAUSE_CHAIN (*w))
+ {
+ tree tmp = OMP_CLAUSE_CHAIN (*w);
+ OMP_CLAUSE_CHAIN (*w) = NULL_TREE;
+ debug_generic_expr (*w);
+ OMP_CLAUSE_CHAIN (*w) = tmp;
+ }
+#endif
failure:
delete grpmap;
delete groups;
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 09/11] Not for committing: noisy sibling-list handling output
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
` (7 preceding siblings ...)
2021-10-01 17:09 ` [PATCH 08/11] Not for committing: noisy topological sorting output Julian Brown
@ 2021-10-01 17:09 ` Julian Brown
2021-10-01 17:10 ` [PATCH 10/11] Not for committing: noisy mapping-group taxonomy Julian Brown
2021-10-01 17:10 ` [PATCH 11/11] OpenMP/OpenACC: [WIP] Add gcc_unreachable to apparently-dead path in build_struct_comp_nodes Julian Brown
10 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:09 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
As a possible aid to review, this is my "printf-style" debugging cruft for
the sibling list handling hoist/rework. It's not meant for committing.
---
gcc/gimplify.c | 131 +++++++++++++++++++++++++++++++++++++++++++++++++
1 file changed, 131 insertions(+)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 2ec83bf273b..ffb6eda5490 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -71,6 +71,7 @@ along with GCC; see the file COPYING3. If not see
#include "tree-nested.h"
//#define NOISY_TOPOSORT
+//#define NOISY_SIBLING_LISTS
/* Hash set of poisoned variables in a bind expr. */
static hash_set<tree> *asan_poisoned_variables = NULL;
@@ -9895,6 +9896,11 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
bool openmp = !(region_type & ORT_ACC);
tree *continue_at = NULL;
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "DECL starts out as:\n");
+ debug_generic_expr (ocd);
+#endif
+
while (TREE_CODE (ocd) == ARRAY_REF)
ocd = TREE_OPERAND (ocd, 0);
@@ -9903,6 +9909,11 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
tree base = extract_base_bit_offset (ocd, &cbitpos, &coffset);
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "BASE after extraction is (%p):\n", (void *) base);
+ debug_generic_expr (base);
+#endif
+
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)
@@ -9917,6 +9928,25 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
if (openmp && attach_detach)
return NULL;
+#ifdef NOISY_SIBLING_LISTS
+ if (struct_map_to_clause)
+ {
+ fprintf (stderr, "s_m_t_c->get (base) = ");
+ debug_generic_expr (base);
+ tree *r = struct_map_to_clause->get (base);
+ fprintf (stderr, "returns: ");
+ if (r)
+ {
+ tree tmp = OMP_CLAUSE_CHAIN (*r);
+ OMP_CLAUSE_CHAIN (*r) = NULL_TREE;
+ debug_generic_expr (*r);
+ OMP_CLAUSE_CHAIN (*r) = tmp;
+ }
+ else
+ fprintf (stderr, "(nothing)\n");
+ }
+#endif
+
if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
{
tree l = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP);
@@ -10026,6 +10056,11 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
{
tree *osc = struct_map_to_clause->get (base);
tree *sc = NULL, *scp = NULL;
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "looked up osc %p for decl (%p)\n", (void *) osc,
+ (void *) base);
+ debug_generic_expr (base);
+#endif
sc = &OMP_CLAUSE_CHAIN (*osc);
/* The struct mapping might be immediately followed by a
FIRSTPRIVATE_POINTER and/or FIRSTPRIVATE_REFERENCE -- if it's an
@@ -10098,6 +10133,17 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
return NULL;
}
}
+#ifdef NOISY_SIBLING_LISTS
+ if (known_eq (coffset, offset) && known_eq (cbitpos, bitpos))
+ {
+ fprintf (stderr, "duplicate offset!\n");
+ tree o1 = OMP_CLAUSE_DECL (*sc);
+ tree o2 = OMP_CLAUSE_DECL (grp_end);
+ debug_generic_expr (o1);
+ debug_generic_expr (o2);
+ }
+ else
+#endif
if (maybe_lt (coffset, offset)
|| (known_eq (coffset, offset)
&& maybe_lt (cbitpos, bitpos)))
@@ -10174,6 +10220,13 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
= cl ? move_concat_nodes_after (cl, tail_chain, grp_start_p,
grp_end, sc)
: move_nodes_after (grp_start_p, grp_end, sc);
+#ifdef NOISY_SIBLING_LISTS
+ if (continue_at)
+ {
+ fprintf (stderr, "continue at (1):\n");
+ debug_generic_expr (*continue_at);
+ }
+#endif
}
else if (*sc != grp_end)
{
@@ -10187,6 +10240,10 @@ build_struct_group (enum omp_region_type region_type, enum tree_code code,
the correct position in the struct component list, which in this
case is just SC. */
move_node_after (grp_end, grp_start_p, sc);
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "continue at (2):\n");
+ debug_generic_expr (*continue_at);
+#endif
}
}
return continue_at;
@@ -10218,6 +10275,16 @@ omp_build_struct_sibling_lists (enum tree_code code,
new_next = NULL;
+#ifdef NOISY_SIBLING_LISTS
+ {
+ tree *tmp = grp->grp_start;
+ grp->grp_start = grp_start_p;
+ fprintf (stderr, "processing group %u:\n", i);
+ debug_mapping_group (grp);
+ grp->grp_start = tmp;
+ }
+#endif
+
if (DECL_P (decl))
continue;
@@ -10252,6 +10319,11 @@ omp_build_struct_sibling_lists (enum tree_code code,
!= GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION))
decl = TREE_OPERAND (decl, 0);
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "using base pointer/decl:\n");
+ debug_generic_expr (decl);
+#endif
+
STRIP_NOPS (decl);
if (TREE_CODE (decl) != COMPONENT_REF)
@@ -10260,6 +10332,15 @@ omp_build_struct_sibling_lists (enum tree_code code,
omp_mapping_group **wholestruct = NULL;
tree wsdecl = omp_containing_struct (OMP_CLAUSE_DECL (c));
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "whole-struct decl is (%s original)\n",
+ wsdecl == OMP_CLAUSE_DECL (c) ? "same as" : "different from");
+
+ debug_generic_expr (wsdecl);
+ fprintf (stderr, "orig was:\n");
+ debug_generic_expr (OMP_CLAUSE_DECL (c));
+#endif
+
if (!(region_type & ORT_ACC) && wsdecl != OMP_CLAUSE_DECL (c))
{
wholestruct = grpmap->get (wsdecl);
@@ -10275,16 +10356,27 @@ omp_build_struct_sibling_lists (enum tree_code code,
if (wholestruct)
{
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "it looks like the whole struct is mapped by:\n");
+ debug_mapping_group (*wholestruct);
+#endif
if (*grp_start_p == grp_end)
{
/* Remove the whole of this mapping -- redundant. */
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "removing this group\n");
+#endif
new_next = grp_start_p;
*grp_start_p = OMP_CLAUSE_CHAIN (grp_end);
}
continue;
}
+#ifdef NOISY_SIBLING_LISTS
+ else
+ fprintf (stderr, "whole struct is not mapped\n");
+#endif
if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
@@ -10360,6 +10452,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
int handled_depend_iterators = -1;
int nowait = -1;
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "GIMPLIFY_SCAN_OMP_CLAUSES, list:\n");
+ debug_generic_expr (*list_p);
+#endif
+
ctx = new_omp_context (region_type);
ctx->code = code;
outer_ctx = ctx->outer_context;
@@ -10399,7 +10496,14 @@ 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);
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "BUILDING STRUCT SIBLING LISTS\n");
+#endif
omp_build_struct_sibling_lists (code, region_type, groups, grpmap);
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "result:\n");
+ debug_generic_expr (*list_p);
+#endif
omp_mapping_group *outlist = NULL;
@@ -10455,6 +10559,9 @@ 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);
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "BUILDING STRUCT SIBLING LISTS\n");
+#endif
omp_build_struct_sibling_lists (code, region_type, groups, grpmap);
delete groups;
@@ -10775,6 +10882,22 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_MAP:
decl = OMP_CLAUSE_DECL (c);
+
+#ifdef NOISY_SIBLING_LISTS
+ {
+ fprintf (stderr, "gimplify_scan_omp_clauses processing: ");
+ debug_generic_expr (c);
+ if (DECL_P (decl))
+ {
+ splay_tree_node n
+ = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+ fprintf (stderr, "decl has node: %p\n", (void *) n);
+ if (n)
+ fprintf (stderr, "flags are: %x\n", n->value);
+ }
+ }
+#endif
+
if (error_operand_p (decl))
remove = true;
switch (code)
@@ -10870,6 +10993,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
{
tree base = strip_components_and_deref (decl);
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "struct, base=");
+ debug_generic_expr (base);
+#endif
if (DECL_P (base))
{
decl = base;
@@ -10951,6 +11078,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (TREE_CODE (cref) == COMPONENT_REF)
{
+#ifdef NOISY_SIBLING_LISTS
+ fprintf (stderr, "we see a component_ref for:\n");
+ debug_generic_expr (c);
+#endif
tree base = cref;
while (base && !DECL_P (base))
{
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 10/11] Not for committing: noisy mapping-group taxonomy
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
` (8 preceding siblings ...)
2021-10-01 17:09 ` [PATCH 09/11] Not for committing: noisy sibling-list handling output Julian Brown
@ 2021-10-01 17:10 ` Julian Brown
2021-10-01 17:10 ` [PATCH 11/11] OpenMP/OpenACC: [WIP] Add gcc_unreachable to apparently-dead path in build_struct_comp_nodes Julian Brown
10 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:10 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
As a possible aid to review, this is code that can be used to enumerate
all the mapping group forms currently in use across the GCC/libgomp
testsuites for OpenMP/OpenACC. These groups have been added somewhat
organically, so there might be a couple of surprises: see e.g. the patch
following this one.
It's not meant for committing.
---
gcc/gimplify.c | 327 +++++++++++++++++++++++++++++++++++++++++++++++++
1 file changed, 327 insertions(+)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index ffb6eda5490..d9fda21413d 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -72,6 +72,7 @@ along with GCC; see the file COPYING3. If not see
//#define NOISY_TOPOSORT
//#define NOISY_SIBLING_LISTS
+//#define NOISY_TAXONOMY
/* Hash set of poisoned variables in a bind expr. */
static hash_set<tree> *asan_poisoned_variables = NULL;
@@ -9010,6 +9011,326 @@ omp_gather_mapping_groups (tree *list_p)
}
}
+#ifdef NOISY_TAXONOMY
+
+static void
+omp_mapping_group_taxonomy (vec<omp_mapping_group> *groups)
+{
+ int num = 0;
+
+ for (auto &it : *groups)
+ {
+ tree node, grp_start = *it.grp_start, grp_end = it.grp_end;
+ gomp_map_kind kind0 = OMP_CLAUSE_MAP_KIND (grp_start), kind1, kind2,
+ kind3;
+ int count = 1;
+ node = grp_start;
+ if (node != grp_end)
+ {
+ node = OMP_CLAUSE_CHAIN (node);
+ kind1 = OMP_CLAUSE_MAP_KIND (node);
+ count++;
+ if (node != grp_end)
+ {
+ node = OMP_CLAUSE_CHAIN (node);
+ kind2 = OMP_CLAUSE_MAP_KIND (node);
+ count++;
+ if (node != grp_end)
+ {
+ node = OMP_CLAUSE_CHAIN (node);
+ kind3 = OMP_CLAUSE_MAP_KIND (node);
+ count++;
+ gcc_assert (node == grp_end);
+ }
+ }
+ }
+
+ fprintf (stderr, "group %d: ", num);
+
+ switch (count)
+ {
+ case 1:
+ if (kind0 == GOMP_MAP_TO
+ || kind0 == GOMP_MAP_FROM
+ || kind0 == GOMP_MAP_TOFROM)
+ fprintf (stderr, "scalar to/from\n");
+ else if (kind0 == GOMP_MAP_ALLOC)
+ fprintf (stderr, "alloc\n");
+ else if (kind0 == GOMP_MAP_POINTER)
+ fprintf (stderr, "pointer (by itself)\n");
+ else if (kind0 == GOMP_MAP_TO_PSET)
+ fprintf (stderr, "map-to-pset (by itself)\n");
+ else if (kind0 == GOMP_MAP_FORCE_PRESENT)
+ fprintf (stderr, "force present\n");
+ else if (kind0 == GOMP_MAP_DELETE)
+ fprintf (stderr, "delete\n");
+ else if (kind0 == GOMP_MAP_FORCE_DEVICEPTR)
+ fprintf (stderr, "force deviceptr\n");
+ else if (kind0 == GOMP_MAP_DEVICE_RESIDENT)
+ fprintf (stderr, "device resident\n");
+ else if (kind0 == GOMP_MAP_LINK)
+ fprintf (stderr, "link\n");
+ else if (kind0 == GOMP_MAP_IF_PRESENT)
+ fprintf (stderr, "if present\n");
+ else if (kind0 == GOMP_MAP_FIRSTPRIVATE)
+ fprintf (stderr, "firstprivate (by itself)\n");
+ else if (kind0 == GOMP_MAP_FIRSTPRIVATE_INT)
+ fprintf (stderr, "firstprivate_int (by itself)\n");
+ else if (kind0 == GOMP_MAP_USE_DEVICE_PTR)
+ fprintf (stderr, "use device ptr\n");
+ else if (kind0 == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
+ fprintf (stderr, "zero-length array section (by itself)\n");
+ else if (kind0 == GOMP_MAP_FORCE_ALLOC)
+ fprintf (stderr, "force alloc\n");
+ else if (kind0 == GOMP_MAP_FORCE_TO
+ || kind0 == GOMP_MAP_FORCE_FROM
+ || kind0 == GOMP_MAP_FORCE_TOFROM)
+ fprintf (stderr, "force to/from (scalar)\n");
+ else if (kind0 == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
+ fprintf (stderr, "use device ptr if present\n");
+ else if (kind0 == GOMP_MAP_ALWAYS_TO
+ || kind0 == GOMP_MAP_ALWAYS_FROM
+ || kind0 == GOMP_MAP_ALWAYS_TOFROM)
+ fprintf (stderr, "always to/from (scalar)\n");
+ else if (kind0 == GOMP_MAP_STRUCT)
+ fprintf (stderr, "struct\n");
+ else if (kind0 == GOMP_MAP_ALWAYS_POINTER)
+ fprintf (stderr, "always pointer (by itself)\n");
+ else if (kind0 == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
+ fprintf (stderr, "ptr to 0-length array section (by itself)\n");
+ else if (kind0 == GOMP_MAP_RELEASE)
+ fprintf (stderr, "release\n");
+ else if (kind0 == GOMP_MAP_ATTACH)
+ fprintf (stderr, "attach\n");
+ else if (kind0 == GOMP_MAP_DETACH)
+ fprintf (stderr, "detach\n");
+ else if (kind0 == GOMP_MAP_FORCE_DETACH)
+ fprintf (stderr, "force detach\n");
+ else if (kind0 == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
+ fprintf (stderr, "attach 0-length array section\n");
+ else if (kind0 == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ fprintf (stderr, "firstprivate ptr (by itself)\n");
+ else if (kind0 == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ fprintf (stderr, "firstprivate ref (by itself)\n");
+ else if (kind0 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "attach/detach (by itself)\n");
+ else
+ fprintf (stderr, "unknown code %d\n", (int) kind0);
+ break;
+
+ case 2:
+ if (kind0 == GOMP_MAP_TO
+ || kind0 == GOMP_MAP_FROM
+ || kind0 == GOMP_MAP_TOFROM)
+ {
+ if (kind1 == GOMP_MAP_POINTER)
+ fprintf (stderr, "to/from, pointer\n");
+ else if (kind1 == GOMP_MAP_ALWAYS_POINTER)
+ fprintf (stderr, "to/from, always pointer\n");
+ else if (kind1 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "to/from, attach/detach\n");
+ else if (kind1 == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ fprintf (stderr, "to/from, firstprivate pointer\n");
+ else if (kind1 == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ fprintf (stderr, "to/from, firstprivate reference\n");
+ else
+ fprintf (stderr, "to/from, unknown code %d\n", (int) kind1);
+ }
+ else if (kind0 == GOMP_MAP_ALWAYS_FROM
+ || kind0 == GOMP_MAP_ALWAYS_TO
+ || kind0 == GOMP_MAP_ALWAYS_TOFROM)
+ {
+ if (kind1 == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ fprintf (stderr, "always to/from, firstprivate pointer\n");
+ else if (kind1 == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ fprintf (stderr, "always to/from, firstprivate reference\n");
+ else if (kind1 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "always to/from, attach/detach\n");
+ else if (kind1 == GOMP_MAP_ALWAYS_POINTER)
+ fprintf (stderr, "always to/from, always pointer\n");
+ else
+ fprintf (stderr, "always to/from, unknown code %d\n",
+ (int) kind1);
+ }
+ else if (kind0 == GOMP_MAP_FORCE_TO
+ || kind0 == GOMP_MAP_FORCE_FROM
+ || kind0 == GOMP_MAP_FORCE_TOFROM)
+ {
+ if (kind1 == GOMP_MAP_POINTER)
+ fprintf (stderr, "force to/from, pointer\n");
+ else
+ fprintf (stderr, "force to/from, unknown code %d\n",
+ (int) kind1);
+ }
+ else if (kind0 == GOMP_MAP_FORCE_PRESENT)
+ {
+ if (kind1 == GOMP_MAP_POINTER)
+ fprintf (stderr, "force present, pointer\n");
+ else if (kind1 == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ fprintf (stderr, "force present, firstprivate pointer\n");
+ else
+ fprintf (stderr, "force present, unknown code %d\n",
+ (int) kind1);
+ }
+ else if (kind0 == GOMP_MAP_ALLOC)
+ {
+ if (kind1 == GOMP_MAP_POINTER)
+ fprintf (stderr, "alloc, pointer\n");
+ else if (kind1 == GOMP_MAP_ALWAYS_POINTER)
+ fprintf (stderr, "alloc, always pointer\n");
+ else if (kind1 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "alloc, attach/detach\n");
+ else if (kind1 == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ fprintf (stderr, "alloc, firstprivate pointer\n");
+ else if (kind1 == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ fprintf (stderr, "alloc, firstprivate reference\n");
+ else
+ fprintf (stderr, "alloc, unknown code %d\n", (int) kind1);
+ }
+ else if (kind0 == GOMP_MAP_RELEASE)
+ {
+ if (kind1 == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ fprintf (stderr, "release, firstprivate reference\n");
+ else if (kind1 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "release, attach/detach\n");
+ else if (kind1 == GOMP_MAP_ALWAYS_POINTER)
+ fprintf (stderr, "release, always pointer\n");
+ else if (kind1 == GOMP_MAP_POINTER)
+ fprintf (stderr, "release, pointer\n");
+ else
+ fprintf (stderr, "release, unknown code %d\n", (int) kind1);
+ }
+ else if (kind0 == GOMP_MAP_DELETE)
+ {
+ if (kind1 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "delete, attach/detach\n");
+ else
+ fprintf (stderr, "delete, unknown code %d\n", (int) kind1);
+ }
+ else if (kind0 == GOMP_MAP_TO_PSET)
+ {
+ if (kind1 == GOMP_MAP_ATTACH)
+ fprintf (stderr, "pset, attach\n");
+ else if (kind1 == GOMP_MAP_DETACH)
+ fprintf (stderr, "pset, detach\n");
+ else
+ fprintf (stderr, "pset, unknown code %d\n", (int) kind1);
+ }
+ else
+ fprintf (stderr, "unknown code %d, unknown code %d\n",
+ (int) kind0, (int) kind1);
+ break;
+
+ case 3:
+ if (kind0 == GOMP_MAP_TO
+ || kind0 == GOMP_MAP_FROM
+ || kind0 == GOMP_MAP_TOFROM)
+ {
+ if (kind1 == GOMP_MAP_POINTER
+ && kind2 == GOMP_MAP_POINTER)
+ fprintf (stderr, "to/from, pointer, pointer\n");
+ else if (kind1 == GOMP_MAP_ALWAYS_POINTER
+ && kind2 == GOMP_MAP_ALWAYS_POINTER)
+ fprintf (stderr, "to/from, always-pointer, always-pointer\n");
+ else if (kind1 == GOMP_MAP_TO_PSET
+ && kind2 == GOMP_MAP_POINTER)
+ fprintf (stderr, "to/from, pset, pointer\n");
+ else if (kind1 == GOMP_MAP_TO_PSET
+ && kind2 == GOMP_MAP_ALWAYS_POINTER)
+ fprintf (stderr, "to/from, pset, always-pointer\n");
+ else if (kind1 == GOMP_MAP_TO_PSET
+ && kind2 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "to/from, pset, attach/detach\n");
+ else if (kind1 == GOMP_MAP_POINTER
+ && kind2 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "to/from, pointer, attach/detach\n");
+ else if (kind1 == GOMP_MAP_ALWAYS_POINTER
+ && kind2 == GOMP_MAP_ATTACH_DETACH)
+ fprintf (stderr, "to/from, always-pointer, attach/detach\n");
+ else
+ fprintf (stderr, "to/from, unknown code %d, unknown code %d\n",
+ (int) kind1, (int) kind2);
+ }
+ else if (kind0 == GOMP_MAP_FORCE_TO
+ || kind0 == GOMP_MAP_FORCE_FROM
+ || kind0 == GOMP_MAP_FORCE_TOFROM)
+ {
+ if (kind1 == GOMP_MAP_TO_PSET
+ && kind2 == GOMP_MAP_POINTER)
+ fprintf (stderr, "force to/from, pset, pointer\n");
+ else
+ fprintf (stderr, "force to/from, unknown code %d, "
+ "unknown code %d\n", (int) kind1, (int) kind2);
+ }
+ else if (kind0 == GOMP_MAP_ALLOC)
+ {
+ if (kind1 == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION
+ && kind2 == GOMP_MAP_ATTACH)
+ fprintf (stderr, "alloc, pointer to z-l-a-s, attach\n");
+ else if (kind1 == GOMP_MAP_TO_PSET
+ && kind2 == GOMP_MAP_POINTER)
+ fprintf (stderr, "alloc, pset, pointer\n");
+ else if (kind1 == GOMP_MAP_TO_PSET
+ && kind2 == GOMP_MAP_ALWAYS_POINTER)
+ fprintf (stderr, "alloc, pset, always-pointer\n");
+ else
+ fprintf (stderr, "alloc, unknown code %d, unknown code %d\n",
+ (int) kind1, (int) kind2);
+ }
+ else if (kind0 == GOMP_MAP_DELETE)
+ {
+ if (kind1 == GOMP_MAP_TO_PSET
+ || kind1 == GOMP_MAP_POINTER)
+ fprintf (stderr, "delete, pset, pointer\n");
+ else
+ fprintf (stderr, "delete, unknown code %d, unknown code %d\n",
+ (int) kind1, (int) kind2);
+ }
+ else
+ fprintf (stderr, "unknown code %d, unknown code %d, "
+ "unknown code %d\n", (int) kind0, (int) kind1,
+ (int) kind2);
+ break;
+
+ case 4:
+ if (kind0 == GOMP_MAP_TO
+ || kind0 == GOMP_MAP_FROM
+ || kind0 == GOMP_MAP_TOFROM)
+ {
+ if (kind1 == GOMP_MAP_TO_PSET
+ && kind2 == GOMP_MAP_POINTER
+ && kind3 == GOMP_MAP_POINTER)
+ fprintf (stderr, "to/from, pset, pointer, pointer\n");
+ else if (kind1 == GOMP_MAP_TO_PSET
+ && kind2 == GOMP_MAP_ALWAYS_POINTER
+ && kind3 == GOMP_MAP_POINTER)
+ fprintf (stderr, "to/from, pset, always pointer, pointer\n");
+ else
+ fprintf (stderr, "to/from, unknown code %d, unknown code %d, "
+ "unknown code %d\n", (int) kind1, (int) kind2,
+ (int) kind3);
+ }
+ else
+ fprintf (stderr, "unknown code %d, unknown code %d, "
+ "unknown code %d, unknown code %d\n", (int) kind0,
+ (int) kind1, (int) kind2, (int) kind3);
+ break;
+
+ default:
+ gcc_unreachable ();
+ }
+
+ tree tmp = OMP_CLAUSE_CHAIN (grp_end);
+ OMP_CLAUSE_CHAIN (grp_end) = NULL_TREE;
+ debug_generic_expr (grp_start);
+ OMP_CLAUSE_CHAIN (grp_end) = tmp;
+
+ num++;
+ }
+}
+
+#endif
+
/* A pointer mapping group GRP may define a block of memory starting at some
base address, and maybe also define a firstprivate pointer or firstprivate
reference that points to that block. The return value is a node containing
@@ -10493,6 +10814,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
groups = omp_gather_mapping_groups (list_p);
if (groups)
{
+#ifdef NOISY_TAXONOMY
+ omp_mapping_group_taxonomy (groups);
+#endif
hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
grpmap = omp_index_mapping_groups (groups);
@@ -10556,6 +10880,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
groups = omp_gather_mapping_groups (list_p);
if (groups)
{
+#ifdef NOISY_TAXONOMY
+ omp_mapping_group_taxonomy (groups);
+#endif
hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
grpmap = omp_index_mapping_groups (groups);
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* [PATCH 11/11] OpenMP/OpenACC: [WIP] Add gcc_unreachable to apparently-dead path in build_struct_comp_nodes
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
` (9 preceding siblings ...)
2021-10-01 17:10 ` [PATCH 10/11] Not for committing: noisy mapping-group taxonomy Julian Brown
@ 2021-10-01 17:10 ` Julian Brown
10 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:10 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Chung-Lin Tang, Thomas Schwinge
The previous "not for committing" taxonomy patch shows that the path
handling "extra nodes" in build_struct_comp_nodes is probably now dead,
at least across the current testsuite. This patch adds gcc_unreachable
on that path: this passes testing, which suggests that the extra node
handling can probably be removed completely. (Otherwise we need test
coverage for that path, ideally!)
This is mostly posted as an FYI: a real patch would probably just remove
the unused code path, if it really isn't needed any more.
Thanks,
Julian
2021-09-29 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (build_struct_comp_nodes): Add gcc_unreachable on code
path that appears to now be unused.
---
gcc/gimplify.c | 3 +++
1 file changed, 3 insertions(+)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index d9fda21413d..3d444d1836f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8625,6 +8625,9 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
OMP_CLAUSE_SIZE (c3) = TYPE_SIZE_UNIT (ptr_type_node);
OMP_CLAUSE_CHAIN (c3) = NULL_TREE;
+ /* Apparently? */
+ gcc_unreachable ();
+
*extra_node = c3;
}
else
--
2.29.2
^ permalink raw reply [flat|nested] 16+ messages in thread
* Re: [PATCH 05/11] OpenMP/OpenACC: Hoist struct sibling list handling in gimplification
2021-10-01 17:09 ` [PATCH 05/11] OpenMP/OpenACC: Hoist struct sibling list handling in gimplification Julian Brown
@ 2021-10-01 17:16 ` Julian Brown
0 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-01 17:16 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge, cltang
Oops, editing:
On Fri, 1 Oct 2021 10:09:03 -0700
Julian Brown <julian@codesourcery.com> wrote:
> 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
...in the main loop in gimplify_scan_omp_clauses.
^ permalink raw reply [flat|nested] 16+ messages in thread
* Re: [PATCH 07/11] OpenMP: Fix non-zero attach/detach bias for struct dereferences
2021-10-01 17:09 ` [PATCH 07/11] OpenMP: Fix non-zero attach/detach bias for struct dereferences Julian Brown
@ 2021-10-11 14:49 ` Julian Brown
0 siblings, 0 replies; 16+ messages in thread
From: Julian Brown @ 2021-10-11 14:49 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
On Fri, 1 Oct 2021 10:09:05 -0700
Julian Brown <julian@codesourcery.com> wrote:
> libgomp/
> * testsuite/libgomp.c++/baseptrs-3.C: Add test (XFAILed for
> now).
This XFAILed test is addressed in the followup patch:
https://gcc.gnu.org/pipermail/gcc-patches/2021-October/581342.html
Cheers,
Julian
^ permalink raw reply [flat|nested] 16+ messages in thread
* Re: [PATCH 01/11] libgomp: Release device lock on cbuf error path
2021-10-01 17:07 ` [PATCH 01/11] libgomp: Release device lock on cbuf error path Julian Brown
@ 2021-10-12 9:23 ` Jakub Jelinek
0 siblings, 0 replies; 16+ messages in thread
From: Jakub Jelinek @ 2021-10-12 9:23 UTC (permalink / raw)
To: Julian Brown; +Cc: gcc-patches, Thomas Schwinge
On Fri, Oct 01, 2021 at 10:07:48AM -0700, Julian Brown wrote:
> This patch releases the device lock on a sanity-checking error path in
> transfer combining (cbuf) handling in libgomp:target.c. This shouldn't
> happen when handling well-formed mapping clauses, but erroneous clauses
> can currently cause a hang if the condition triggers.
>
> Tested with offloading to NVPTX. OK?
>
> 2021-09-29 Julian Brown <julian@codesourcery.com>
>
> libgomp/
> * target.c (gomp_copy_host2dev): Release device lock on cbuf
> error path.
Ok, thanks. This doesn't seem to depend on anything else, so
can be committed separately right away.
> diff --git a/libgomp/target.c b/libgomp/target.c
> index 65bb40100e5..84c6fdf2c47 100644
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -385,7 +385,10 @@ gomp_copy_host2dev (struct gomp_device_descr *devicep,
> else if (cbuf->chunks[middle].start <= doff)
> {
> if (doff + sz > cbuf->chunks[middle].end)
> - gomp_fatal ("internal libgomp cbuf error");
> + {
> + gomp_mutex_unlock (&devicep->lock);
> + gomp_fatal ("internal libgomp cbuf error");
> + }
> memcpy ((char *) cbuf->buf + (doff - cbuf->chunks[0].start),
> h, sz);
> return;
> --
> 2.29.2
Jakub
^ permalink raw reply [flat|nested] 16+ messages in thread
* Re: [PATCH 02/11] Remove base_ind/base_ref handling from extract_base_bit_offset
2021-10-01 17:07 ` [PATCH 02/11] Remove base_ind/base_ref handling from extract_base_bit_offset Julian Brown
@ 2021-10-12 9:27 ` Jakub Jelinek
0 siblings, 0 replies; 16+ messages in thread
From: Jakub Jelinek @ 2021-10-12 9:27 UTC (permalink / raw)
To: Julian Brown; +Cc: gcc-patches, Chung-Lin Tang, Thomas Schwinge
On Fri, Oct 01, 2021 at 10:07:49AM -0700, Julian Brown wrote:
> In preparation for follow-up patches extending struct dereference
> handling for OpenMP, this patch removes base_ind/base_ref handling from
> gimplify.c:extract_base_bit_offset. This arguably simplifies some of the
> code around the callers of the function also, though subsequent patches
> modify those parts further.
>
> OK for mainline?
>
> Thanks,
>
> Julian
>
> 2021-09-29 Julian Brown <julian@codesourcery.com>
>
> gcc/
> * gimplify.c (extract_base_bit_offset): Remove BASE_IND, BASE_REF and
> OPENMP parameters.
> (strip_indirections): New function.
> (build_struct_group): Update calls to extract_base_bit_offset.
> Rearrange indirect/reference handling accordingly. Use extracted base
> instead of passed-in decl when grouping component accesses together.
This is ok for trunk once the whole series is approved.
Jakub
^ permalink raw reply [flat|nested] 16+ messages in thread
end of thread, other threads:[~2021-10-12 9:27 UTC | newest]
Thread overview: 16+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-10-01 17:07 [PATCH 00/11] OpenMP: Deep struct dereferences Julian Brown
2021-10-01 17:07 ` [PATCH 01/11] libgomp: Release device lock on cbuf error path Julian Brown
2021-10-12 9:23 ` Jakub Jelinek
2021-10-01 17:07 ` [PATCH 02/11] Remove base_ind/base_ref handling from extract_base_bit_offset Julian Brown
2021-10-12 9:27 ` Jakub Jelinek
2021-10-01 17:07 ` [PATCH 03/11] OpenMP 5.0: Clause ordering for OpenMP 5.0 (topological sorting by base pointer) Julian Brown
2021-10-01 17:07 ` [PATCH 04/11] Remove omp_target_reorder_clauses Julian Brown
2021-10-01 17:09 ` [PATCH 05/11] OpenMP/OpenACC: Hoist struct sibling list handling in gimplification Julian Brown
2021-10-01 17:16 ` Julian Brown
2021-10-01 17:09 ` [PATCH 06/11] OpenMP: Allow array ref components for C & C++ Julian Brown
2021-10-01 17:09 ` [PATCH 07/11] OpenMP: Fix non-zero attach/detach bias for struct dereferences Julian Brown
2021-10-11 14:49 ` Julian Brown
2021-10-01 17:09 ` [PATCH 08/11] Not for committing: noisy topological sorting output Julian Brown
2021-10-01 17:09 ` [PATCH 09/11] Not for committing: noisy sibling-list handling output Julian Brown
2021-10-01 17:10 ` [PATCH 10/11] Not for committing: noisy mapping-group taxonomy Julian Brown
2021-10-01 17:10 ` [PATCH 11/11] OpenMP/OpenACC: [WIP] Add gcc_unreachable to apparently-dead path in build_struct_comp_nodes 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).