* [PATCH 00/16] OpenMP: lvalues in "map" clauses and struct handling rework
@ 2021-11-25 14:07 Julian Brown
2021-11-25 14:07 ` [PATCH 01/16] Rewrite GOMP_MAP_ATTACH_DETACH mappings unconditionally Julian Brown
` (3 more replies)
0 siblings, 4 replies; 5+ messages in thread
From: Julian Brown @ 2021-11-25 14:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
Hi Jakub,
This is a rebased/slightly bug-fixed version of several previously posted
patch series, all in one place for ease of reference. The series should
be applied on top of Chung-Lin's two patches:
"Improve OpenMP target support for C++ [PR92120 v5]"
https://gcc.gnu.org/pipermail/gcc-patches/2021-November/584602.html
"Remove array section base-pointer mapping semantics, and other front-end
adjustments (mainline trunk)"
https://gcc.gnu.org/pipermail/gcc-patches/2021-November/584994.html
And supersedes the following three patch series:
"Topological sort for OpenMP 5.0 base pointers"
https://gcc.gnu.org/pipermail/gcc-patches/2021-August/577211.html
"OpenMP: Deep struct dereferences"
https://gcc.gnu.org/pipermail/gcc-patches/2021-October/580721.html
"Parsing of lvalues for "map" clauses for C and C++"
https://gcc.gnu.org/pipermail/gcc-patches/2021-November/584445.html
Tested with offloading to NVPTX and bootstrapped. Further commentary
on individual patches. OK?
Thanks,
Julian
Julian Brown (16):
Rewrite GOMP_MAP_ATTACH_DETACH mappings unconditionally
OpenMP/OpenACC: Move array_ref/indirect_ref handling code out of
extract_base_bit_offset
OpenACC/OpenMP: Refactor struct lowering in gimplify.c
OpenACC: Rework indirect struct handling in gimplify.c
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
OpenMP: Handle reference-typed struct members
OpenACC: Make deep-copy-arrayofstruct.c a libgomp/runtime test
Add debug_omp_expr
OpenMP: Add inspector class to unify mapped address analysis
OpenMP: lvalue parsing for map clauses (C++)
OpenMP: lvalue parsing for map clauses (C)
gcc/c-family/c-common.h | 45 +
gcc/c-family/c-omp.c | 210 ++
gcc/c/c-parser.c | 150 +-
gcc/c/c-tree.h | 1 +
gcc/c/c-typeck.c | 250 +-
gcc/cp/error.c | 9 +
gcc/cp/parser.c | 141 +-
gcc/cp/parser.h | 3 +
gcc/cp/semantics.c | 290 +-
gcc/fortran/trans-openmp.c | 20 +-
gcc/gimplify.c | 2458 +++++++++++------
gcc/omp-low.c | 23 +-
gcc/testsuite/c-c++-common/gomp/map-1.c | 3 +-
gcc/testsuite/c-c++-common/gomp/map-6.c | 6 +-
gcc/testsuite/g++.dg/goacc/member-array-acc.C | 13 +
gcc/testsuite/g++.dg/gomp/ind-base-3.C | 38 +
gcc/testsuite/g++.dg/gomp/map-assignment-1.C | 12 +
gcc/testsuite/g++.dg/gomp/map-inc-1.C | 10 +
gcc/testsuite/g++.dg/gomp/map-lvalue-ref-1.C | 19 +
gcc/testsuite/g++.dg/gomp/map-ptrmem-1.C | 36 +
gcc/testsuite/g++.dg/gomp/map-ptrmem-2.C | 39 +
.../g++.dg/gomp/map-static-cast-lvalue-1.C | 17 +
gcc/testsuite/g++.dg/gomp/map-ternary-1.C | 20 +
gcc/testsuite/g++.dg/gomp/member-array-2.C | 86 +
gcc/testsuite/g++.dg/gomp/member-array-omp.C | 13 +
gcc/testsuite/g++.dg/gomp/pr67522.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 +-
.../g++.dg/gomp/unmappable-component-1.C | 21 +
gcc/tree-pretty-print.c | 45 +
gcc/tree-pretty-print.h | 1 +
gcc/tree.def | 3 +
libgomp/testsuite/libgomp.c++/baseptrs-3.C | 275 ++
libgomp/testsuite/libgomp.c++/ind-base-1.C | 162 ++
libgomp/testsuite/libgomp.c++/ind-base-2.C | 49 +
libgomp/testsuite/libgomp.c++/map-comma-1.C | 15 +
.../testsuite/libgomp.c++/map-rvalue-ref-1.C | 22 +
.../testsuite/libgomp.c++/member-array-1.C | 89 +
libgomp/testsuite/libgomp.c++/struct-ref-1.C | 97 +
.../libgomp.c-c++-common/baseptrs-1.c | 50 +
.../libgomp.c-c++-common/baseptrs-2.c | 70 +
.../libgomp.c-c++-common/ind-base-4.c | 50 +
.../libgomp.c-c++-common/unary-ptr-1.c | 16 +
.../testsuite/libgomp.oacc-c++/deep-copy-17.C | 101 +
.../libgomp.oacc-c-c++-common/deep-copy-15.c | 68 +
.../libgomp.oacc-c-c++-common/deep-copy-16.c | 231 ++
.../deep-copy-arrayofstruct.c | 2 +-
50 files changed, 4114 insertions(+), 1187 deletions(-)
create mode 100644 gcc/testsuite/g++.dg/goacc/member-array-acc.C
create mode 100644 gcc/testsuite/g++.dg/gomp/ind-base-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/map-assignment-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/map-inc-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/map-lvalue-ref-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/map-ptrmem-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/map-ptrmem-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/map-static-cast-lvalue-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/map-ternary-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/member-array-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/member-array-omp.C
create mode 100644 gcc/testsuite/g++.dg/gomp/unmappable-component-1.C
create mode 100644 libgomp/testsuite/libgomp.c++/baseptrs-3.C
create mode 100644 libgomp/testsuite/libgomp.c++/ind-base-1.C
create mode 100644 libgomp/testsuite/libgomp.c++/ind-base-2.C
create mode 100644 libgomp/testsuite/libgomp.c++/map-comma-1.C
create mode 100644 libgomp/testsuite/libgomp.c++/map-rvalue-ref-1.C
create mode 100644 libgomp/testsuite/libgomp.c++/member-array-1.C
create mode 100644 libgomp/testsuite/libgomp.c++/struct-ref-1.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
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/ind-base-4.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/unary-ptr-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
rename {gcc/testsuite/c-c++-common/goacc => libgomp/testsuite/libgomp.oacc-c-c++-common}/deep-copy-arrayofstruct.c (98%)
--
2.29.2
^ permalink raw reply [flat|nested] 5+ messages in thread
* [PATCH 01/16] Rewrite GOMP_MAP_ATTACH_DETACH mappings unconditionally
2021-11-25 14:07 [PATCH 00/16] OpenMP: lvalues in "map" clauses and struct handling rework Julian Brown
@ 2021-11-25 14:07 ` Julian Brown
2021-11-25 14:07 ` [PATCH 02/16] OpenMP/OpenACC: Move array_ref/indirect_ref handling code out of extract_base_bit_offset Julian Brown
` (2 subsequent siblings)
3 siblings, 0 replies; 5+ messages in thread
From: Julian Brown @ 2021-11-25 14:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
It never makes sense for a GOMP_MAP_ATTACH_DETACH mapping to survive
beyond gimplify.c, so this patch rewrites such mappings to GOMP_MAP_ATTACH
or GOMP_MAP_DETACH unconditionally (rather than checking for a list
of types of OpenACC or OpenMP constructs), in cases where it hasn't
otherwise been done already in the preceding code.
Previously posted here:
https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570399.html
https://gcc.gnu.org/pipermail/gcc-patches/2021-June/571711.html (og11)
OK?
Thanks,
Julian
2021-06-02 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Simplify condition
for changing GOMP_MAP_ATTACH_DETACH to GOMP_MAP_ATTACH or
GOMP_MAP_DETACH.
---
gcc/gimplify.c | 10 +---------
1 file changed, 1 insertion(+), 9 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 4cd62270a10..8d8735ae4c1 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9965,15 +9965,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
skip_map_struct:
;
}
- else if ((code == OACC_ENTER_DATA
- || code == OACC_EXIT_DATA
- || code == OACC_DATA
- || code == OACC_PARALLEL
- || code == OACC_KERNELS
- || code == OACC_SERIAL
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA)
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
{
gomp_map_kind k = ((code == OACC_EXIT_DATA
|| code == OMP_TARGET_EXIT_DATA)
--
2.29.2
^ permalink raw reply [flat|nested] 5+ messages in thread
* [PATCH 02/16] OpenMP/OpenACC: Move array_ref/indirect_ref handling code out of extract_base_bit_offset
2021-11-25 14:07 [PATCH 00/16] OpenMP: lvalues in "map" clauses and struct handling rework Julian Brown
2021-11-25 14:07 ` [PATCH 01/16] Rewrite GOMP_MAP_ATTACH_DETACH mappings unconditionally Julian Brown
@ 2021-11-25 14:07 ` Julian Brown
2021-11-25 14:07 ` [PATCH 03/16] OpenACC/OpenMP: Refactor struct lowering in gimplify.c Julian Brown
2021-11-25 14:07 ` [PATCH 04/16] OpenACC: Rework indirect struct handling " Julian Brown
3 siblings, 0 replies; 5+ messages in thread
From: Julian Brown @ 2021-11-25 14:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
This patch slightly cleans up the semantics of extract_base_bit_offset,
in that the stripping of ARRAY_REFS/INDIRECT_REFS out of
extract_base_bit_offset is moved back into the (two) call sites of the
function. This is done in preparation for follow-on patches that extend
the function.
Previously posted for the og11 branch here (patch & reversion/rework):
https://gcc.gnu.org/pipermail/gcc-patches/2021-June/571712.html
https://gcc.gnu.org/pipermail/gcc-patches/2021-June/571884.html
OK?
Thanks,
Julian
2021-06-03 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (extract_base_bit_offset): Don't look through ARRAY_REFs or
INDIRECT_REFs here.
(build_struct_group): Reinstate previous behaviour for handling
ARRAY_REFs/INDIRECT_REFs.
---
gcc/gimplify.c | 59 +++++++++++++++++++++++++-------------------------
1 file changed, 29 insertions(+), 30 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 8d8735ae4c1..1baea68920b 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8675,31 +8675,7 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
poly_offset_int poffset;
if (base_ref)
- {
- *base_ref = NULL_TREE;
-
- while (TREE_CODE (base) == ARRAY_REF)
- base = TREE_OPERAND (base, 0);
-
- if (TREE_CODE (base) == INDIRECT_REF)
- base = TREE_OPERAND (base, 0);
- }
- else
- {
- if (TREE_CODE (base) == ARRAY_REF)
- {
- while (TREE_CODE (base) == ARRAY_REF)
- base = TREE_OPERAND (base, 0);
- if (TREE_CODE (base) != COMPONENT_REF
- || TREE_CODE (TREE_TYPE (base)) != ARRAY_TYPE)
- return NULL_TREE;
- }
- else if (TREE_CODE (base) == INDIRECT_REF
- && TREE_CODE (TREE_OPERAND (base, 0)) == COMPONENT_REF
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0)))
- == REFERENCE_TYPE))
- base = TREE_OPERAND (base, 0);
- }
+ *base_ref = NULL_TREE;
base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode,
&unsignedp, &reversep, &volatilep);
@@ -9673,12 +9649,17 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
poly_offset_int offset1;
poly_int64 bitpos1;
tree tree_offset1;
- tree base_ref;
+ tree base_ref, ocd = OMP_CLAUSE_DECL (c);
- tree base
- = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
- &bitpos1, &offset1,
- &tree_offset1);
+ while (TREE_CODE (ocd) == ARRAY_REF)
+ ocd = TREE_OPERAND (ocd, 0);
+
+ if (TREE_CODE (ocd) == INDIRECT_REF)
+ ocd = TREE_OPERAND (ocd, 0);
+
+ tree base = extract_base_bit_offset (ocd, &base_ref,
+ &bitpos1, &offset1,
+ &tree_offset1);
bool do_map_struct = (base == decl && !tree_offset1);
@@ -9871,6 +9852,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
poly_offset_int offsetn;
poly_int64 bitposn;
tree tree_offsetn;
+
+ if (TREE_CODE (sc_decl) == ARRAY_REF)
+ {
+ while (TREE_CODE (sc_decl) == ARRAY_REF)
+ sc_decl = TREE_OPERAND (sc_decl, 0);
+ if (TREE_CODE (sc_decl) != COMPONENT_REF
+ || (TREE_CODE (TREE_TYPE (sc_decl))
+ != ARRAY_TYPE))
+ break;
+ }
+ else if (TREE_CODE (sc_decl) == INDIRECT_REF
+ && (TREE_CODE (TREE_OPERAND (sc_decl, 0))
+ == COMPONENT_REF)
+ && (TREE_CODE (TREE_TYPE
+ (TREE_OPERAND (sc_decl, 0)))
+ == REFERENCE_TYPE))
+ sc_decl = TREE_OPERAND (sc_decl, 0);
+
tree base
= extract_base_bit_offset (sc_decl, NULL,
&bitposn, &offsetn,
--
2.29.2
^ permalink raw reply [flat|nested] 5+ messages in thread
* [PATCH 03/16] OpenACC/OpenMP: Refactor struct lowering in gimplify.c
2021-11-25 14:07 [PATCH 00/16] OpenMP: lvalues in "map" clauses and struct handling rework Julian Brown
2021-11-25 14:07 ` [PATCH 01/16] Rewrite GOMP_MAP_ATTACH_DETACH mappings unconditionally Julian Brown
2021-11-25 14:07 ` [PATCH 02/16] OpenMP/OpenACC: Move array_ref/indirect_ref handling code out of extract_base_bit_offset Julian Brown
@ 2021-11-25 14:07 ` Julian Brown
2021-11-25 14:07 ` [PATCH 04/16] OpenACC: Rework indirect struct handling " Julian Brown
3 siblings, 0 replies; 5+ messages in thread
From: Julian Brown @ 2021-11-25 14:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
(Previously submitted here:
https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570398.html)
This patch is a second attempt at refactoring struct component mapping
handling for OpenACC/OpenMP during gimplification, after the patch I
posted here:
https://gcc.gnu.org/pipermail/gcc-patches/2018-November/510503.html
And improved here, post-review:
https://gcc.gnu.org/pipermail/gcc-patches/2019-November/533394.html
This patch goes further, in that the struct-handling code is outlined
into its own function (to create the "GOMP_MAP_STRUCT" node and the
sorted list of nodes immediately following it, from a set of mappings
of components of a given struct or derived type). I've also gone through
the list-handling code and attempted to add comments documenting how it
works to the best of my understanding, and broken out a couple of helper
functions in order to (hopefully) have the code self-document better also.
OK?
Thanks,
Julian
2021-06-02 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (insert_struct_comp_map): Refactor function into...
(build_struct_comp_nodes): This new function. Remove list handling
and improve self-documentation.
(insert_node_after, move_node_after, move_nodes_after,
move_concat_nodes_after): New helper functions.
(build_struct_group): New function to build up GOMP_MAP_STRUCT node
groups to map struct components. Outlined from...
(gimplify_scan_omp_clauses): Here. Call above function.
---
gcc/gimplify.c | 976 +++++++++++++++++++++++++++++++------------------
1 file changed, 611 insertions(+), 365 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 1baea68920b..c5e058d6d1f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8588,73 +8588,66 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
return 1;
}
-/* Insert a GOMP_MAP_ALLOC or GOMP_MAP_RELEASE node following a
- GOMP_MAP_STRUCT mapping. C is an always_pointer mapping. STRUCT_NODE is
- the struct node to insert the new mapping after (when the struct node is
- initially created). PREV_NODE is the first of two or three mappings for a
- pointer, and is either:
- - the node before C, when a pair of mappings is used, e.g. for a C/C++
- array section.
- - not the node before C. This is true when we have a reference-to-pointer
- type (with a mapping for the reference and for the pointer), or for
- Fortran derived-type mappings with a GOMP_MAP_TO_PSET.
- If SCP is non-null, the new node is inserted before *SCP.
- if SCP is null, the new node is inserted before PREV_NODE.
- The return type is:
- - PREV_NODE, if SCP is non-null.
- - The newly-created ALLOC or RELEASE node, if SCP is null.
- - The second newly-created ALLOC or RELEASE node, if we are mapping a
- reference to a pointer. */
+/* For a set of mappings describing an array section pointed to by a struct
+ (or derived type, etc.) component, create an "alloc" or "release" node to
+ insert into a list following a GOMP_MAP_STRUCT node. For some types of
+ mapping (e.g. Fortran arrays with descriptors), an additional mapping may
+ be created that is inserted into the list of mapping nodes attached to the
+ directive being processed -- not part of the sorted list of nodes after
+ GOMP_MAP_STRUCT.
+
+ CODE is the code of the directive being processed. GRP_START and GRP_END
+ are the first and last of two or three nodes representing this array section
+ mapping (e.g. a data movement node like GOMP_MAP_{TO,FROM}, optionally a
+ GOMP_MAP_TO_PSET, and finally a GOMP_MAP_ALWAYS_POINTER). EXTRA_NODE is
+ filled with the additional node described above, if needed.
+
+ This function does not add the new nodes to any lists itself. It is the
+ responsibility of the caller to do that. */
static tree
-insert_struct_comp_map (enum tree_code code, tree c, tree struct_node,
- tree prev_node, tree *scp)
+build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
+ tree *extra_node)
{
enum gomp_map_kind mkind
= (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)
? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
- tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- tree cl = scp ? prev_node : c2;
+ gcc_assert (grp_start != grp_end);
+
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
- OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c));
- OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node;
- if (OMP_CLAUSE_CHAIN (prev_node) != c
- && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
- == GOMP_MAP_TO_PSET))
- OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node));
+ OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (grp_end));
+ OMP_CLAUSE_CHAIN (c2) = NULL_TREE;
+ tree grp_mid = NULL_TREE;
+ if (OMP_CLAUSE_CHAIN (grp_start) != grp_end)
+ grp_mid = OMP_CLAUSE_CHAIN (grp_start);
+
+ if (grp_mid
+ && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET)
+ OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (grp_mid);
else
OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
- if (struct_node)
- OMP_CLAUSE_CHAIN (struct_node) = c2;
- /* We might need to create an additional mapping if we have a reference to a
- pointer (in C++). Don't do this if we have something other than a
- GOMP_MAP_ALWAYS_POINTER though, i.e. a GOMP_MAP_TO_PSET. */
- if (OMP_CLAUSE_CHAIN (prev_node) != c
- && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP
- && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
- == GOMP_MAP_ALWAYS_POINTER)
- || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
- == GOMP_MAP_ATTACH_DETACH)))
+ if (grp_mid
+ && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_ATTACH_DETACH))
{
- tree c4 = OMP_CLAUSE_CHAIN (prev_node);
- tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ tree c3
+ = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c3, mkind);
- OMP_CLAUSE_DECL (c3) = unshare_expr (OMP_CLAUSE_DECL (c4));
+ OMP_CLAUSE_DECL (c3) = unshare_expr (OMP_CLAUSE_DECL (grp_mid));
OMP_CLAUSE_SIZE (c3) = TYPE_SIZE_UNIT (ptr_type_node);
- OMP_CLAUSE_CHAIN (c3) = prev_node;
- if (!scp)
- OMP_CLAUSE_CHAIN (c2) = c3;
- else
- cl = c3;
+ OMP_CLAUSE_CHAIN (c3) = NULL_TREE;
+
+ *extra_node = c3;
}
+ else
+ *extra_node = NULL_TREE;
- if (scp)
- *scp = c2;
-
- return cl;
+ return c2;
}
/* Strip ARRAY_REFS or an indirect ref off BASE, find the containing object,
@@ -8992,6 +8985,558 @@ omp_lastprivate_for_combined_outer_constructs (struct gimplify_omp_ctx *octx,
omp_notice_variable (octx, decl, true);
}
+/* Link node NEWNODE so it is pointed to by chain INSERT_AT. NEWNODE's chain
+ is linked to the previous node pointed to by INSERT_AT. */
+
+static tree *
+insert_node_after (tree newnode, tree *insert_at)
+{
+ OMP_CLAUSE_CHAIN (newnode) = *insert_at;
+ *insert_at = newnode;
+ return &OMP_CLAUSE_CHAIN (newnode);
+}
+
+/* Move NODE (which is currently pointed to by the chain OLD_POS) so it is
+ pointed to by chain MOVE_AFTER instead. */
+
+static void
+move_node_after (tree node, tree *old_pos, tree *move_after)
+{
+ gcc_assert (node == *old_pos);
+ *old_pos = OMP_CLAUSE_CHAIN (node);
+ OMP_CLAUSE_CHAIN (node) = *move_after;
+ *move_after = node;
+}
+
+/* Move nodes from FIRST_PTR (pointed to by previous node's chain) to
+ LAST_NODE to after MOVE_AFTER chain. Similar to below function, but no
+ new nodes are prepended to the list before splicing into the new position.
+ Return the position we should continue scanning the list at, or NULL to
+ stay where we were. */
+
+static tree *
+move_nodes_after (tree *first_ptr, tree last_node, tree *move_after)
+{
+ if (first_ptr == move_after)
+ return NULL;
+
+ tree tmp = *first_ptr;
+ *first_ptr = OMP_CLAUSE_CHAIN (last_node);
+ OMP_CLAUSE_CHAIN (last_node) = *move_after;
+ *move_after = tmp;
+
+ return first_ptr;
+}
+
+/* Concatenate two lists described by [FIRST_NEW, LAST_NEW_TAIL] and
+ [FIRST_PTR, LAST_NODE], and insert them in the OMP clause list after chain
+ pointer MOVE_AFTER.
+
+ The latter list was previously part of the OMP clause list, and the former
+ (prepended) part is comprised of new nodes.
+
+ We start with a list of nodes starting with a struct mapping node. We
+ rearrange the list so that new nodes starting from FIRST_NEW and whose last
+ node's chain is LAST_NEW_TAIL comes directly after MOVE_AFTER, followed by
+ the group of mapping nodes we are currently processing (from the chain
+ FIRST_PTR to LAST_NODE). The return value is the pointer to the next chain
+ we should continue processing from, or NULL to stay where we were.
+
+ The transformation (in the case where MOVE_AFTER and FIRST_PTR are
+ different) is worked through below. Here we are processing LAST_NODE, and
+ FIRST_PTR points at the preceding mapping clause:
+
+ #. mapping node chain
+ ---------------------------------------------------
+ A. struct_node [->B]
+ B. comp_1 [->C]
+ C. comp_2 [->D (move_after)]
+ D. map_to_3 [->E]
+ E. attach_3 [->F (first_ptr)]
+ F. map_to_4 [->G (continue_at)]
+ G. attach_4 (last_node) [->H]
+ H. ...
+
+ *last_new_tail = *first_ptr;
+
+ I. new_node (first_new) [->F (last_new_tail)]
+
+ *first_ptr = OMP_CLAUSE_CHAIN (last_node)
+
+ #. mapping node chain
+ ----------------------------------------------------
+ A. struct_node [->B]
+ B. comp_1 [->C]
+ C. comp_2 [->D (move_after)]
+ D. map_to_3 [->E]
+ E. attach_3 [->H (first_ptr)]
+ F. map_to_4 [->G (continue_at)]
+ G. attach_4 (last_node) [->H]
+ H. ...
+
+ I. new_node (first_new) [->F (last_new_tail)]
+
+ OMP_CLAUSE_CHAIN (last_node) = *move_after;
+
+ #. mapping node chain
+ ---------------------------------------------------
+ A. struct_node [->B]
+ B. comp_1 [->C]
+ C. comp_2 [->D (move_after)]
+ D. map_to_3 [->E]
+ E. attach_3 [->H (continue_at)]
+ F. map_to_4 [->G]
+ G. attach_4 (last_node) [->D]
+ H. ...
+
+ I. new_node (first_new) [->F (last_new_tail)]
+
+ *move_after = first_new;
+
+ #. mapping node chain
+ ---------------------------------------------------
+ A. struct_node [->B]
+ B. comp_1 [->C]
+ C. comp_2 [->I (move_after)]
+ D. map_to_3 [->E]
+ E. attach_3 [->H (continue_at)]
+ F. map_to_4 [->G]
+ G. attach_4 (last_node) [->D]
+ H. ...
+ I. new_node (first_new) [->F (last_new_tail)]
+
+ or, in order:
+
+ #. mapping node chain
+ ---------------------------------------------------
+ A. struct_node [->B]
+ B. comp_1 [->C]
+ C. comp_2 [->I (move_after)]
+ I. new_node (first_new) [->F (last_new_tail)]
+ F. map_to_4 [->G]
+ G. attach_4 (last_node) [->D]
+ D. map_to_3 [->E]
+ E. attach_3 [->H (continue_at)]
+ H. ...
+*/
+
+static tree *
+move_concat_nodes_after (tree first_new, tree *last_new_tail, tree *first_ptr,
+ tree last_node, tree *move_after)
+{
+ tree *continue_at = NULL;
+ *last_new_tail = *first_ptr;
+ if (first_ptr == move_after)
+ *move_after = first_new;
+ else
+ {
+ *first_ptr = OMP_CLAUSE_CHAIN (last_node);
+ continue_at = first_ptr;
+ OMP_CLAUSE_CHAIN (last_node) = *move_after;
+ *move_after = first_new;
+ }
+ return continue_at;
+}
+
+/* Mapping struct members causes an additional set of nodes to be created,
+ starting with GOMP_MAP_STRUCT followed by a number of mappings equal to the
+ number of members being mapped, in order of ascending position (address or
+ bitwise).
+
+ We scan through the list of mapping clauses, calling this function for each
+ struct member mapping we find, and build up the list of mappings after the
+ initial GOMP_MAP_STRUCT node. For pointer members, these will be
+ newly-created ALLOC nodes. For non-pointer members, the existing mapping is
+ moved into place in the sorted list.
+
+ struct {
+ int *a;
+ int *b;
+ int c;
+ int *d;
+ };
+
+ #pragma (acc|omp directive) copy(struct.a[0:n], struct.b[0:n], struct.c,
+ struct.d[0:n])
+
+ GOMP_MAP_STRUCT (4)
+ [GOMP_MAP_FIRSTPRIVATE_REFERENCE -- for refs to structs]
+ GOMP_MAP_ALLOC (struct.a)
+ GOMP_MAP_ALLOC (struct.b)
+ GOMP_MAP_TO (struct.c)
+ GOMP_MAP_ALLOC (struct.d)
+ ...
+
+ In the case where we are mapping references to pointers, or in Fortran if
+ we are mapping an array with a descriptor, additional nodes may be created
+ after the struct node list also.
+
+ The return code is:
+ - DECL, if we just created the initial GOMP_MAP_STRUCT node.
+ - NULL_TREE, if we inserted the new struct member successfully.
+ - error_mark_node if an error occurred.
+
+ *CONT is set to TRUE if we should skip further processing and move to the
+ 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,
+ 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)
+{
+ poly_offset_int coffset;
+ poly_int64 cbitpos;
+ tree base_ref, tree_coffset;
+ tree ocd = OMP_CLAUSE_DECL (c);
+
+ while (TREE_CODE (ocd) == ARRAY_REF)
+ ocd = TREE_OPERAND (ocd, 0);
+
+ if (TREE_CODE (ocd) == INDIRECT_REF)
+ ocd = TREE_OPERAND (ocd, 0);
+
+ tree base = extract_base_bit_offset (ocd, &base_ref,
+ &cbitpos, &coffset, &tree_coffset);
+
+ bool do_map_struct = (base == 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. */
+ if ((code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_UPDATE
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ && 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_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)))
+ {
+ tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT;
+
+ OMP_CLAUSE_SET_MAP_KIND (l, k);
+
+ if (base_ref)
+ OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
+ else
+ {
+ OMP_CLAUSE_DECL (l) = unshare_expr (decl);
+ if (!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))
+ ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
+ : 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);
+
+ if (ptr || attach_detach)
+ {
+ tree extra_node;
+ tree alloc_node
+ = build_struct_comp_nodes (code, *prev_list_p, c, &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;
+
+ if (extra_node)
+ {
+ OMP_CLAUSE_CHAIN (extra_node) = *insert_node_pos;
+ OMP_CLAUSE_CHAIN (alloc_node) = extra_node;
+ }
+ else
+ 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);
+
+ /* 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;
+ }
+ *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 (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;
+ }
+
+ return decl;
+ }
+ else if (struct_map_to_clause)
+ {
+ tree *osc = struct_map_to_clause->get (decl);
+ 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_REFERENCE if it is a reference. (This added node is
+ removed in omp-low.c after it has been processed there.) */
+ if (*sc != c
+ && 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)
+ break;
+ else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
+ && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != INDIRECT_REF
+ && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != ARRAY_REF)
+ break;
+ else
+ {
+ tree sc_decl = OMP_CLAUSE_DECL (*sc);
+ poly_offset_int offset;
+ poly_int64 bitpos;
+ tree tree_offset;
+
+ if (TREE_CODE (sc_decl) == ARRAY_REF)
+ {
+ while (TREE_CODE (sc_decl) == ARRAY_REF)
+ sc_decl = TREE_OPERAND (sc_decl, 0);
+ if (TREE_CODE (sc_decl) != COMPONENT_REF
+ || TREE_CODE (TREE_TYPE (sc_decl)) != ARRAY_TYPE)
+ break;
+ }
+ else if (TREE_CODE (sc_decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (sc_decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (sc_decl, 0)))
+ == REFERENCE_TYPE))
+ sc_decl = TREE_OPERAND (sc_decl, 0);
+
+ tree base = extract_base_bit_offset (sc_decl, NULL, &bitpos,
+ &offset, &tree_offset);
+ if (base != decl)
+ break;
+ if (scp)
+ continue;
+ if ((region_type & ORT_ACC) != 0)
+ {
+ /* This duplicate checking code is currently only enabled for
+ OpenACC. */
+ tree d1 = OMP_CLAUSE_DECL (*sc);
+ tree d2 = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (d1) == ARRAY_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ while (TREE_CODE (d2) == ARRAY_REF)
+ d2 = TREE_OPERAND (d2, 0);
+ if (TREE_CODE (d1) == INDIRECT_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ if (TREE_CODE (d2) == INDIRECT_REF)
+ d2 = TREE_OPERAND (d2, 0);
+ while (TREE_CODE (d1) == COMPONENT_REF)
+ if (TREE_CODE (d2) == COMPONENT_REF
+ && TREE_OPERAND (d1, 1) == TREE_OPERAND (d2, 1))
+ {
+ d1 = TREE_OPERAND (d1, 0);
+ d2 = TREE_OPERAND (d2, 0);
+ }
+ else
+ break;
+ if (d1 == d2)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once in map clauses",
+ OMP_CLAUSE_DECL (c));
+ return error_mark_node;
+ }
+ }
+ if (maybe_lt (coffset, offset)
+ || (known_eq (coffset, offset)
+ && maybe_lt (cbitpos, bitpos)))
+ {
+ if (ptr || attach_detach)
+ scp = sc;
+ else
+ break;
+ }
+ }
+
+ if (!attach)
+ OMP_CLAUSE_SIZE (*osc)
+ = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), size_one_node);
+ 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 *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.
+ 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
+ should insert in the middle of the struct component list
+ (else NULL to insert at end).
+ alloc_node : the "alloc" node for the structure (pointer-type)
+ component. We insert at SCP (if present), else SC
+ (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.
+ tail_chain : pointer to chain of last prepended node.
+
+ The general idea is we move the nodes for this struct mapping
+ together: the alloc node goes into the sorted list directly after
+ the struct mapping, and any extra nodes (together with the nodes
+ mapping arrays pointed to by struct components) get moved after
+ that list. When SCP is NULL, we insert the nodes at SC, i.e. at
+ the end of the struct component mapping list. It's important that
+ the alloc_node comes first in that case because it's part of the
+ sorted component mapping list (but subsequent nodes are not!). */
+
+ if (scp)
+ insert_node_after (alloc_node, scp);
+
+ /* Make [cl,tail_chain] a list of the alloc node (if we haven't
+ already inserted it) and the extra_node (if it is present). The
+ list can be empty if we added alloc_node above and there is no
+ extra node. */
+ if (scp && extra_node)
+ {
+ cl = extra_node;
+ tail_chain = &OMP_CLAUSE_CHAIN (extra_node);
+ }
+ else if (extra_node)
+ {
+ OMP_CLAUSE_CHAIN (alloc_node) = extra_node;
+ cl = alloc_node;
+ tail_chain = &OMP_CLAUSE_CHAIN (extra_node);
+ }
+ else if (!scp)
+ {
+ cl = alloc_node;
+ 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;
+ }
+ }
+ else if (*sc != c)
+ {
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
+ == GS_ERROR)
+ return error_mark_node;
+ /* 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;
+ }
+ }
+ return NULL_TREE;
+}
+
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
and previous omp contexts. */
@@ -9645,324 +10190,25 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
}
}
-
- poly_offset_int offset1;
- poly_int64 bitpos1;
- tree tree_offset1;
- tree base_ref, ocd = OMP_CLAUSE_DECL (c);
-
- while (TREE_CODE (ocd) == ARRAY_REF)
- ocd = TREE_OPERAND (ocd, 0);
-
- if (TREE_CODE (ocd) == INDIRECT_REF)
- ocd = TREE_OPERAND (ocd, 0);
-
- tree base = extract_base_bit_offset (ocd, &base_ref,
- &bitpos1, &offset1,
- &tree_offset1);
-
- bool do_map_struct = (base == decl && !tree_offset1);
-
- 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))
-
+ 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)
{
- /* 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;
+ remove = true;
+ break;
}
-
- /* We currently don't handle non-constant offset accesses wrt to
- GOMP_MAP_STRUCT elements. */
- if (!do_map_struct)
- goto skip_map_struct;
-
- /* Nor for attach_detach for OpenMP. */
- if ((code == OMP_TARGET
- || code == OMP_TARGET_DATA
- || code == OMP_TARGET_UPDATE
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA)
- && attach_detach)
+ else if (add_decl && DECL_P (add_decl))
{
- 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);
- }
-
- goto skip_map_struct;
+ decl = add_decl;
+ goto do_add_decl;
}
-
- 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)))
- {
- tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT
- : GOMP_MAP_STRUCT;
-
- OMP_CLAUSE_SET_MAP_KIND (l, k);
- if (base_ref)
- OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
- else
- {
- OMP_CLAUSE_DECL (l) = unshare_expr (decl);
- if (!DECL_P (OMP_CLAUSE_DECL (l))
- && (gimplify_expr (&OMP_CLAUSE_DECL (l),
- pre_p, NULL, is_gimple_lvalue,
- fb_lvalue)
- == GS_ERROR))
- {
- remove = true;
- break;
- }
- }
- OMP_CLAUSE_SIZE (l)
- = (!attach
- ? size_int (1)
- : DECL_P (OMP_CLAUSE_DECL (l))
- ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
- : 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);
- if (ptr || attach_detach)
- {
- tree **sc = (struct_seen_clause
- ? struct_seen_clause->get (decl)
- : NULL);
- tree *insert_node_pos = sc ? *sc : prev_list_p;
-
- insert_struct_comp_map (code, c, l, *insert_node_pos,
- NULL);
- *insert_node_pos = l;
- prev_list_p = NULL;
- }
- else
- {
- OMP_CLAUSE_CHAIN (l) = c;
- *list_p = l;
- list_p = &OMP_CLAUSE_CHAIN (l);
- }
- 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;
- }
- 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 (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;
- }
-
- if (DECL_P (decl))
- goto do_add_decl;
- }
- else if (struct_map_to_clause)
- {
- tree *osc = struct_map_to_clause->get (decl);
- tree *sc = NULL, *scp = NULL;
- if (n != NULL
- && (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
- || ptr
- || attach_detach))
- n->value |= GOVD_SEEN;
- sc = &OMP_CLAUSE_CHAIN (*osc);
- if (*sc != c
- && (OMP_CLAUSE_MAP_KIND (*sc)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
- sc = &OMP_CLAUSE_CHAIN (*sc);
- /* Here "prev_list_p" is the end of the inserted
- alloc/release nodes after the struct node, OSC. */
- for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
- if ((ptr || attach_detach) && sc == prev_list_p)
- break;
- else if (TREE_CODE (OMP_CLAUSE_DECL (*sc))
- != COMPONENT_REF
- && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
- != INDIRECT_REF)
- && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
- != ARRAY_REF))
- break;
- else
- {
- tree sc_decl = OMP_CLAUSE_DECL (*sc);
- poly_offset_int offsetn;
- poly_int64 bitposn;
- tree tree_offsetn;
-
- if (TREE_CODE (sc_decl) == ARRAY_REF)
- {
- while (TREE_CODE (sc_decl) == ARRAY_REF)
- sc_decl = TREE_OPERAND (sc_decl, 0);
- if (TREE_CODE (sc_decl) != COMPONENT_REF
- || (TREE_CODE (TREE_TYPE (sc_decl))
- != ARRAY_TYPE))
- break;
- }
- else if (TREE_CODE (sc_decl) == INDIRECT_REF
- && (TREE_CODE (TREE_OPERAND (sc_decl, 0))
- == COMPONENT_REF)
- && (TREE_CODE (TREE_TYPE
- (TREE_OPERAND (sc_decl, 0)))
- == REFERENCE_TYPE))
- sc_decl = TREE_OPERAND (sc_decl, 0);
-
- tree base
- = extract_base_bit_offset (sc_decl, NULL,
- &bitposn, &offsetn,
- &tree_offsetn);
- if (base != decl)
- break;
- if (scp)
- continue;
- if ((region_type & ORT_ACC) != 0)
- {
- /* This duplicate checking code is currently only
- enabled for OpenACC. */
- tree d1 = OMP_CLAUSE_DECL (*sc);
- tree d2 = OMP_CLAUSE_DECL (c);
- while (TREE_CODE (d1) == ARRAY_REF)
- d1 = TREE_OPERAND (d1, 0);
- while (TREE_CODE (d2) == ARRAY_REF)
- d2 = TREE_OPERAND (d2, 0);
- if (TREE_CODE (d1) == INDIRECT_REF)
- d1 = TREE_OPERAND (d1, 0);
- if (TREE_CODE (d2) == INDIRECT_REF)
- d2 = TREE_OPERAND (d2, 0);
- while (TREE_CODE (d1) == COMPONENT_REF)
- if (TREE_CODE (d2) == COMPONENT_REF
- && TREE_OPERAND (d1, 1)
- == TREE_OPERAND (d2, 1))
- {
- d1 = TREE_OPERAND (d1, 0);
- d2 = TREE_OPERAND (d2, 0);
- }
- else
- break;
- if (d1 == d2)
- {
- error_at (OMP_CLAUSE_LOCATION (c),
- "%qE appears more than once in map "
- "clauses", OMP_CLAUSE_DECL (c));
- remove = true;
- break;
- }
- }
- if (maybe_lt (offset1, offsetn)
- || (known_eq (offset1, offsetn)
- && maybe_lt (bitpos1, bitposn)))
- {
- if (ptr || attach_detach)
- scp = sc;
- else
- break;
- }
- }
- if (remove)
- break;
- if (!attach)
- OMP_CLAUSE_SIZE (*osc)
- = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
- size_one_node);
- if (ptr || attach_detach)
- {
- tree cl = insert_struct_comp_map (code, c, NULL,
- *prev_list_p, scp);
- if (sc == prev_list_p)
- {
- *sc = cl;
- prev_list_p = NULL;
- }
- else
- {
- *prev_list_p = OMP_CLAUSE_CHAIN (c);
- list_p = prev_list_p;
- prev_list_p = NULL;
- OMP_CLAUSE_CHAIN (c) = *sc;
- *sc = cl;
- continue;
- }
- }
- else if (*sc != c)
- {
- if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
- fb_lvalue)
- == GS_ERROR)
- {
- remove = true;
- break;
- }
- *list_p = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = *sc;
- *sc = c;
- continue;
- }
- }
- skip_map_struct:
- ;
+ if (cont)
+ continue;
}
else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
{
--
2.29.2
^ permalink raw reply [flat|nested] 5+ messages in thread
* [PATCH 04/16] OpenACC: Rework indirect struct handling in gimplify.c
2021-11-25 14:07 [PATCH 00/16] OpenMP: lvalues in "map" clauses and struct handling rework Julian Brown
` (2 preceding siblings ...)
2021-11-25 14:07 ` [PATCH 03/16] OpenACC/OpenMP: Refactor struct lowering in gimplify.c Julian Brown
@ 2021-11-25 14:07 ` Julian Brown
3 siblings, 0 replies; 5+ messages in thread
From: Julian Brown @ 2021-11-25 14:07 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, Thomas Schwinge
(Previously posted here:
https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570400.html)
This patch reworks indirect struct handling in gimplify.c (i.e. for
struct components mapped with "mystruct->a[0:n]", "mystruct->b", etc.),
for OpenACC. The key observation leading to these changes was that
component mappings of references-to-structures is already implemented
and working, and indirect struct component handling via a pointer can
work quite similarly. That lets us remove some earlier, special-case
handling for mapping indirect struct component accesses for OpenACC,
which required the pointed-to struct to be manually mapped before the
indirect component mapping.
With this patch, you can map struct components directly (e.g. an array
slice "mystruct->a[0:n]") just like you can map a non-indirect struct
component slice ("mystruct.a[0:n]"). Both references-to-pointers (with
the former syntax) and references to structs (with the latter syntax)
work now.
For Fortran class pointers, we no longer re-use GOMP_MAP_TO_PSET for the
class metadata (the structure that points to the class data and vptr)
-- it is instead treated as any other struct.
For C++, the struct handling also works for class members ("this->foo"),
without having to explicitly map "this[:1]" first.
For OpenACC, we permit chained indirect component references
("mystruct->a->b[0:n]"), though only the last part of such mappings will
trigger an attach/detach operation. To properly use such a construct
on the target, you must still manually map "mystruct->a[:1]" first --
but there's no need to map "mystruct[:1]" explicitly before that.
This version of the patch avoids altering code paths for OpenMP,
where possible. (Those are dealt with by later patches in this series.)
OK?
Thanks,
Julian
2021-06-02 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Don't create GOMP_MAP_TO_PSET
mappings for class metadata, nor GOMP_MAP_POINTER mappings for
POINTER_TYPE_P decls.
gcc/
* gimplify.c (extract_base_bit_offset): Add BASE_IND and OPENMP
parameters. Handle pointer-typed indirect references for OpenACC
alongside reference-typed ones.
(strip_components_and_deref, aggregate_base_p): New functions.
(build_struct_group): Add pointer type indirect ref handling,
including chained references, for OpenACC. Also handle references to
structs for OpenACC. Conditionalise bits for OpenMP only where
appropriate.
(gimplify_scan_omp_clauses): Rework pointer-type indirect structure
access handling to work more like the reference-typed handling for
OpenACC only.
* omp-low.c (scan_sharing_clauses): Handle pointer-type indirect struct
references, and references to pointers to structs also.
gcc/testsuite/
* g++.dg/goacc/member-array-acc.C: New test.
* g++.dg/gomp/member-array-omp.C: New test.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c: New test.
* testsuite/libgomp.oacc-c++/deep-copy-17.C: New test.
---
gcc/fortran/trans-openmp.c | 20 +-
gcc/gimplify.c | 214 +++++++++++++---
gcc/omp-low.c | 16 +-
gcc/testsuite/g++.dg/goacc/member-array-acc.C | 13 +
gcc/testsuite/g++.dg/gomp/member-array-omp.C | 13 +
.../testsuite/libgomp.oacc-c++/deep-copy-17.C | 101 ++++++++
.../libgomp.oacc-c-c++-common/deep-copy-15.c | 68 ++++++
.../libgomp.oacc-c-c++-common/deep-copy-16.c | 231 ++++++++++++++++++
8 files changed, 618 insertions(+), 58 deletions(-)
create mode 100644 gcc/testsuite/g++.dg/goacc/member-array-acc.C
create mode 100644 gcc/testsuite/g++.dg/gomp/member-array-omp.C
create mode 100644 libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 7d761e90dd7..508e02306e9 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -3034,30 +3034,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
tree present = gfc_omp_check_optional_argument (decl, true);
if (openacc && n->sym->ts.type == BT_CLASS)
{
- tree type = TREE_TYPE (decl);
if (n->sym->attr.optional)
sorry ("optional class parameter");
- if (POINTER_TYPE_P (type))
- {
- node4 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
- OMP_CLAUSE_DECL (node4) = decl;
- OMP_CLAUSE_SIZE (node4) = size_int (0);
- decl = build_fold_indirect_ref (decl);
- }
tree ptr = gfc_class_data_get (decl);
ptr = build_fold_indirect_ref (ptr);
OMP_CLAUSE_DECL (node) = ptr;
OMP_CLAUSE_SIZE (node) = gfc_class_vtab_size_get (decl);
node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
- OMP_CLAUSE_DECL (node2) = decl;
- OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
- node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH_DETACH);
- OMP_CLAUSE_DECL (node3) = gfc_class_data_get (decl);
- OMP_CLAUSE_SIZE (node3) = size_int (0);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_ATTACH_DETACH);
+ OMP_CLAUSE_DECL (node2) = gfc_class_data_get (decl);
+ OMP_CLAUSE_SIZE (node2) = size_int (0);
goto finalize_map_clause;
}
else if (POINTER_TYPE_P (TREE_TYPE (decl))
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index c5e058d6d1f..fcc278d07cf 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8658,8 +8658,9 @@ 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_ref, poly_int64 *bitposp,
- poly_offset_int *poffsetp, tree *offsetp)
+extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref,
+ poly_int64 *bitposp, poly_offset_int *poffsetp,
+ tree *offsetp, bool openmp)
{
tree offset;
poly_int64 bitsize, bitpos;
@@ -8667,20 +8668,38 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
int unsignedp, reversep, volatilep = 0;
poly_offset_int poffset;
+ if (base_ind)
+ *base_ind = NULL_TREE;
+
if (base_ref)
*base_ref = NULL_TREE;
base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode,
&unsignedp, &reversep, &volatilep);
- tree orig_base = base;
-
+ 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)
- base = TREE_OPERAND (base, 0);
+ {
+ if (base_ref)
+ *base_ref = base;
+ base = TREE_OPERAND (base, 0);
+ }
+
+ if (!openmp)
+ STRIP_NOPS (base);
if (offset && poly_int_tree_p (offset))
{
@@ -8697,10 +8716,6 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
*poffsetp = poffset;
*offsetp = offset;
- /* Set *BASE_REF if BASE was a dereferenced reference variable. */
- if (base_ref && orig_base != base)
- *base_ref = orig_base;
-
return base;
}
@@ -8727,6 +8742,48 @@ is_or_contains_p (tree expr, tree base_ptr)
return operand_equal_p (expr, base_ptr);
}
+/* Remove COMPONENT_REFS and indirections from EXPR. */
+
+static tree
+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);
+
+ 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;
+}
+
/* 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. */
@@ -9192,8 +9249,9 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
{
poly_offset_int coffset;
poly_int64 cbitpos;
- tree base_ref, tree_coffset;
+ tree base_ind, base_ref, tree_coffset;
tree ocd = OMP_CLAUSE_DECL (c);
+ bool openmp = !(region_type & ORT_ACC);
while (TREE_CODE (ocd) == ARRAY_REF)
ocd = TREE_OPERAND (ocd, 0);
@@ -9201,8 +9259,8 @@ 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_ref,
- &cbitpos, &coffset, &tree_coffset);
+ tree base = extract_base_bit_offset (ocd, &base_ind, &base_ref, &cbitpos,
+ &coffset, &tree_coffset, openmp);
bool do_map_struct = (base == decl && !tree_coffset);
@@ -9241,12 +9299,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
return NULL_TREE;
/* Nor for attach_detach for OpenMP. */
- if ((code == OMP_TARGET
- || code == OMP_TARGET_DATA
- || code == OMP_TARGET_UPDATE
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA)
- && attach_detach)
+ if (openmp && attach_detach)
{
if (DECL_P (decl))
{
@@ -9269,12 +9322,15 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
OMP_CLAUSE_SET_MAP_KIND (l, k);
- if (base_ref)
+ 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 (!DECL_P (OMP_CLAUSE_DECL (l))
+ 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;
@@ -9330,6 +9386,48 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
OMP_CLAUSE_CHAIN (l) = c2;
}
+ else if (!openmp
+ && (base_ind || base_ref)
+ && (region_type & ORT_TARGET))
+ {
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ 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;
@@ -9338,7 +9436,8 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
/* If this is a *pointer-to-struct expression, make sure a
firstprivate map of the base-pointer exists. */
- if (component_ref_p
+ if (openmp
+ && component_ref_p
&& ((TREE_CODE (decl) == MEM_REF
&& integer_zerop (TREE_OPERAND (decl, 1)))
|| INDIRECT_REF_P (decl))
@@ -9368,10 +9467,12 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
n->value |= GOVD_SEEN;
sc = &OMP_CLAUSE_CHAIN (*osc);
/* The struct mapping might be immediately followed by a
- FIRSTPRIVATE_REFERENCE if it is a reference. (This added node is
- removed in omp-low.c after it has been processed there.) */
+ 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
- && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ && (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)
@@ -9401,9 +9502,10 @@ 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, &bitpos,
- &offset, &tree_offset);
- if (base != decl)
+ tree base = extract_base_bit_offset (sc_decl, NULL, NULL,
+ &bitpos, &offset,
+ &tree_offset, openmp);
+ if (!base || !operand_equal_p (base, decl, 0))
break;
if (scp)
continue;
@@ -9524,8 +9626,9 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
}
else if (*sc != c)
{
- if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
- == GS_ERROR)
+ if (openmp
+ && (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
+ == GS_ERROR))
return error_mark_node;
/* In the non-pointer case, the mapping clause itself is moved into
the correct position in the struct component list, which in this
@@ -10027,10 +10130,43 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree indir_base = NULL_TREE;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
- if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
- && TREE_CODE (*pd) == COMPONENT_REF
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
- && code != OACC_UPDATE)
+ 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)
{
@@ -10133,11 +10269,13 @@ 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 ((DECL_P (decl)
- || (component_ref_p
- && (INDIRECT_REF_P (decl)
- || TREE_CODE (decl) == MEM_REF
- || TREE_CODE (decl) == ARRAY_REF)))
+ 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
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 707cc4606c8..bb459e6273e 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1571,8 +1571,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (TREE_CODE (decl) == COMPONENT_REF
|| (TREE_CODE (decl) == INDIRECT_REF
&& TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == REFERENCE_TYPE)))
+ && (((TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE)
+ || (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == POINTER_TYPE)))))
break;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
@@ -13749,6 +13751,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
is_ref = false;
bool ref_to_array = false;
+ bool ref_to_ptr = false;
if (is_ref)
{
type = TREE_TYPE (type);
@@ -13767,6 +13770,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
new_var = decl2;
type = TREE_TYPE (new_var);
}
+ else if (TREE_CODE (type) == REFERENCE_TYPE
+ && TREE_CODE (TREE_TYPE (type)) == POINTER_TYPE)
+ {
+ type = TREE_TYPE (type);
+ ref_to_ptr = true;
+ }
x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx);
x = fold_convert_loc (clause_loc, type, x);
if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
@@ -13783,7 +13792,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (ref_to_array)
x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
- if (is_ref && !ref_to_array)
+ if ((is_ref && !ref_to_array)
+ || ref_to_ptr)
{
tree t = create_tmp_var_raw (type, get_name (var));
gimple_add_tmp_var (t);
diff --git a/gcc/testsuite/g++.dg/goacc/member-array-acc.C b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
new file mode 100644
index 00000000000..e0c11570f5d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
@@ -0,0 +1,13 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+struct Foo {
+ float *a;
+ void init(int N) {
+ a = new float[N];
+ #pragma acc enter data create(a[0:N])
+ }
+};
+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" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/member-array-omp.C b/gcc/testsuite/g++.dg/gomp/member-array-omp.C
new file mode 100644
index 00000000000..a53aa44592d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/member-array-omp.C
@@ -0,0 +1,13 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+struct Foo {
+ float *a;
+ void init(int N) {
+ a = new float[N];
+ #pragma omp target enter data map(alloc:a[0:N])
+ }
+};
+int main() { Foo x; x.init(1024); }
+
+/* { dg-final { scan-tree-dump {map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:this->a \[bias: 0\]\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
new file mode 100644
index 00000000000..dacbb520f3d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
@@ -0,0 +1,101 @@
+#include <cassert>
+
+/* Test attach/detach operation with pointers and references to structs. */
+
+typedef struct mystruct {
+ int *a;
+ int b;
+ int *c;
+ int d;
+ int *e;
+} mystruct;
+
+void str (void)
+{
+ int a[10], c[10], e[10];
+ mystruct m = { .a = a, .c = c, .e = e };
+ a[0] = 5;
+ c[0] = 7;
+ e[0] = 9;
+ #pragma acc parallel copy(m.a[0:10], m.b, m.c[0:10], m.d, m.e[0:10])
+ {
+ m.a[0] = m.c[0] + m.e[0];
+ }
+ assert (m.a[0] == 7 + 9);
+}
+
+void strp (void)
+{
+ int *a = new int[10];
+ int *c = new int[10];
+ int *e = new int[10];
+ mystruct *m = new mystruct;
+ m->a = a;
+ m->c = c;
+ m->e = e;
+ a[0] = 6;
+ c[0] = 8;
+ e[0] = 10;
+ #pragma acc parallel copy(m->a[0:10], m->b, m->c[0:10], m->d, m->e[0:10])
+ {
+ m->a[0] = m->c[0] + m->e[0];
+ }
+ assert (m->a[0] == 8 + 10);
+ delete m;
+ delete[] a;
+ delete[] c;
+ delete[] e;
+}
+
+void strr (void)
+{
+ int *a = new int[10];
+ int *c = new int[10];
+ int *e = new int[10];
+ mystruct m;
+ mystruct &n = m;
+ n.a = a;
+ n.c = c;
+ n.e = e;
+ a[0] = 7;
+ c[0] = 9;
+ e[0] = 11;
+ #pragma acc parallel copy(n.a[0:10], n.b, n.c[0:10], n.d, n.e[0:10])
+ {
+ n.a[0] = n.c[0] + n.e[0];
+ }
+ assert (n.a[0] == 9 + 11);
+ delete[] a;
+ delete[] c;
+ delete[] e;
+}
+
+void strrp (void)
+{
+ int a[10], c[10], e[10];
+ mystruct *m = new mystruct;
+ mystruct *&n = m;
+ n->a = a;
+ n->b = 3;
+ n->c = c;
+ n->d = 5;
+ n->e = e;
+ a[0] = 8;
+ c[0] = 10;
+ e[0] = 12;
+ #pragma acc parallel copy(n->a[0:10], n->c[0:10], n->e[0:10])
+ {
+ n->a[0] = n->c[0] + n->e[0];
+ }
+ assert (n->a[0] == 10 + 12);
+ delete m;
+}
+
+int main (int argc, char *argv[])
+{
+ str ();
+ strp ();
+ strr ();
+ strrp ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
new file mode 100644
index 00000000000..27fe1a9d07d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
@@ -0,0 +1,68 @@
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+ non-zero. */
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+ m->c = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->b[i] = 0;
+ m->c[i] = 0;
+ }
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+#pragma acc parallel loop copy(m->a[0:N])
+ for (j = 0; j < N; j++)
+ m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+ for (j = 0; j < N; j++)
+ {
+ m->b[j]++;
+ if (j > 5 && j < N - 5)
+ m->c[j]++;
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m->c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m->c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m->c);
+ free (m);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
new file mode 100644
index 00000000000..a7308e8c98b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
@@ -0,0 +1,231 @@
+#include <stdlib.h>
+
+/* Test mapping chained indirect struct accesses, mixed in different ways. */
+
+typedef struct {
+ int *a;
+ int b;
+ int *c;
+} str1;
+
+typedef struct {
+ int d;
+ int *e;
+ str1 *f;
+} str2;
+
+typedef struct {
+ int g;
+ int h;
+ str2 *s2;
+} str3;
+
+typedef struct {
+ str3 m;
+ str3 n;
+} str4;
+
+void
+zero_arrays (str4 *s, int N)
+{
+ for (int i = 0; i < N; i++)
+ {
+ s->m.s2->e[i] = 0;
+ s->m.s2->f->a[i] = 0;
+ s->m.s2->f->c[i] = 0;
+ s->n.s2->e[i] = 0;
+ s->n.s2->f->a[i] = 0;
+ s->n.s2->f->c[i] = 0;
+ }
+}
+
+void
+alloc_s2 (str2 **s, int N)
+{
+ (*s) = (str2 *) malloc (sizeof (str2));
+ (*s)->f = (str1 *) malloc (sizeof (str1));
+ (*s)->e = (int *) malloc (sizeof (int) * N);
+ (*s)->f->a = (int *) malloc (sizeof (int) * N);
+ (*s)->f->c = (int *) malloc (sizeof (int) * N);
+}
+
+int main (int argc, char* argv[])
+{
+ const int N = 1024;
+ str4 p, *q;
+ int i;
+
+ alloc_s2 (&p.m.s2, N);
+ alloc_s2 (&p.n.s2, N);
+ q = (str4 *) malloc (sizeof (str4));
+ alloc_s2 (&q->m.s2, N);
+ alloc_s2 (&q->n.s2, N);
+
+ zero_arrays (&p, N);
+
+ for (int i = 0; i < 99; i++)
+ {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc parallel loop copy(p.m.s2->e[:N])
+ for (int j = 0; j < N; j++)
+ p.m.s2->e[j]++;
+#pragma acc exit data delete(p.m.s2[:1])
+ }
+
+ for (i = 0; i < N; i++)
+ if (p.m.s2->e[i] != 99)
+ abort ();
+
+ zero_arrays (&p, N);
+
+ for (int i = 0; i < 99; i++)
+ {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N])
+ for (int j = 0; j < N; j++)
+ {
+ p.m.s2->f->a[j]++;
+ p.m.s2->f->c[j]++;
+ }
+#pragma acc exit data delete(p.m.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1])
+ }
+
+ for (i = 0; i < N; i++)
+ if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99)
+ abort ();
+
+ zero_arrays (&p, N);
+
+ for (int i = 0; i < 99; i++)
+ {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1]) copyin(p.n.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N]) \
+ copy(p.n.s2->f->a[:N]) copy(p.n.s2->f->c[:N])
+ for (int j = 0; j < N; j++)
+ {
+ p.m.s2->f->a[j]++;
+ p.m.s2->f->c[j]++;
+ p.n.s2->f->a[j]++;
+ p.n.s2->f->c[j]++;
+ }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+ }
+
+ for (i = 0; i < N; i++)
+ if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99
+ || p.n.s2->f->a[i] != 99 || p.n.s2->f->c[i] != 99)
+ abort ();
+
+ zero_arrays (&p, N);
+
+ for (int i = 0; i < 99; i++)
+ {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.n.s2->e[:N]) copyin(p.n.s2->f[:1]) \
+ copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.n.s2->f->a[:N])
+ for (int j = 0; j < N; j++)
+ {
+ p.m.s2->f->a[j]++;
+ p.n.s2->f->a[j]++;
+ p.n.s2->e[j]++;
+ }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1]) \
+ copyout(p.n.s2->e[:N])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+ }
+
+ for (i = 0; i < N; i++)
+ if (p.m.s2->f->a[i] != 99 || p.n.s2->f->a[i] != 99
+ || p.n.s2->e[i] != 99)
+ abort ();
+
+ zero_arrays (q, N);
+
+ for (int i = 0; i < 99; i++)
+ {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc parallel loop copy(q->m.s2->e[:N])
+ for (int j = 0; j < N; j++)
+ q->m.s2->e[j]++;
+#pragma acc exit data delete(q->m.s2[:1])
+ }
+
+ for (i = 0; i < N; i++)
+ if (q->m.s2->e[i] != 99)
+ abort ();
+
+ zero_arrays (q, N);
+
+ for (int i = 0; i < 99; i++)
+ {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N])
+ for (int j = 0; j < N; j++)
+ {
+ q->m.s2->f->a[j]++;
+ q->m.s2->f->c[j]++;
+ }
+#pragma acc exit data delete(q->m.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1])
+ }
+
+ for (i = 0; i < N; i++)
+ if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99)
+ abort ();
+
+ zero_arrays (q, N);
+
+ for (int i = 0; i < 99; i++)
+ {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1]) copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N]) \
+ copy(q->n.s2->f->a[:N]) copy(q->n.s2->f->c[:N])
+ for (int j = 0; j < N; j++)
+ {
+ q->m.s2->f->a[j]++;
+ q->m.s2->f->c[j]++;
+ q->n.s2->f->a[j]++;
+ q->n.s2->f->c[j]++;
+ }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+ }
+
+ for (i = 0; i < N; i++)
+ if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99
+ || q->n.s2->f->a[i] != 99 || q->n.s2->f->c[i] != 99)
+ abort ();
+
+ zero_arrays (q, N);
+
+ for (int i = 0; i < 99; i++)
+ {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->n.s2->e[:N]) copyin(q->m.s2->f[:1]) \
+ copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->n.s2->f->a[:N])
+ for (int j = 0; j < N; j++)
+ {
+ q->m.s2->f->a[j]++;
+ q->n.s2->f->a[j]++;
+ q->n.s2->e[j]++;
+ }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1]) \
+ copyout(q->n.s2->e[:N])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+ }
+
+ for (i = 0; i < N; i++)
+ if (q->m.s2->f->a[i] != 99 || q->n.s2->f->a[i] != 99
+ || q->n.s2->e[i] != 99)
+ abort ();
+
+ return 0;
+}
--
2.29.2
^ permalink raw reply [flat|nested] 5+ messages in thread
end of thread, other threads:[~2021-11-25 14:07 UTC | newest]
Thread overview: 5+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-11-25 14:07 [PATCH 00/16] OpenMP: lvalues in "map" clauses and struct handling rework Julian Brown
2021-11-25 14:07 ` [PATCH 01/16] Rewrite GOMP_MAP_ATTACH_DETACH mappings unconditionally Julian Brown
2021-11-25 14:07 ` [PATCH 02/16] OpenMP/OpenACC: Move array_ref/indirect_ref handling code out of extract_base_bit_offset Julian Brown
2021-11-25 14:07 ` [PATCH 03/16] OpenACC/OpenMP: Refactor struct lowering in gimplify.c Julian Brown
2021-11-25 14:07 ` [PATCH 04/16] OpenACC: Rework indirect struct handling " 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).