[PATCH 4/5] Rework indirect struct handling for OpenACC/OpenMP in gimplify.c
Julian Brown
julian@codesourcery.com
Fri May 14 21:26:26 GMT 2021
This patch reworks indirect struct handling in gimplify.c (i.e. for struct
components mapped with "mystruct->a[0:n]", "mystruct->b", etc.), for
both OpenACC and OpenMP. 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 patch incorporates parts of Chung-Lin's patch "Recommit "Enable
gimplify GOMP_MAP_STRUCT handling of (COMPONENT_REF (INDIRECT_REF
...)) map clauses"." from the og10 branch.
OK for trunk?
Thanks,
Julian
2021-05-14 Julian Brown <julian@codesourcery.com>
Chung-Lin Tang <cltang@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 (tree-hash-traits.h): Include.
(extract_base_bit_offset): Add BASE_IND parameter. Handle
pointer-typed indirect references alongside reference-typed ones.
(strip_components_and_deref, aggregate_base_p): New functions.
(build_struct_group): Update struct_map_to_clause type. Add pointer
type indirect ref handling, including chained references. Handle
pointers and references to structs in OpenACC regions as well as
OpenMP ones.
(gimplify_scan_omp_clauses): Remove struct_deref_set handling. Rework
pointer-type indirect structure access handling to work more like
the reference-typed handling.
* 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 (XFAILed for now).
* g++.dg/gomp/member-array-omp.C: New test (XFAILed for now).
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 | 285 ++++++++++--------
gcc/omp-low.c | 16 +-
gcc/testsuite/g++.dg/goacc/member-array-acc.C | 14 +
gcc/testsuite/g++.dg/gomp/member-array-omp.C | 14 +
.../testsuite/libgomp.oacc-c++/deep-copy-17.C | 101 +++++++
.../libgomp.oacc-c-c++-common/deep-copy-15.c | 71 +++++
.../libgomp.oacc-c-c++-common/deep-copy-16.c | 231 ++++++++++++++
8 files changed, 612 insertions(+), 140 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 5666cd68c7e..ff614ffe744 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2721,30 +2721,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 69ab637367c..c0f068a725d 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -53,6 +53,7 @@ along with GCC; see the file COPYING3. If not see
#include "langhooks.h"
#include "tree-cfg.h"
#include "tree-ssa.h"
+#include "tree-hash-traits.h"
#include "omp-general.h"
#include "omp-low.h"
#include "gimple-low.h"
@@ -8351,8 +8352,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_ref, poly_int64 *bitposp,
- poly_offset_int *poffsetp)
+extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref,
+ poly_int64 *bitposp, poly_offset_int *poffsetp)
{
tree offset;
poly_int64 bitsize, bitpos;
@@ -8360,6 +8361,9 @@ 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;
@@ -8380,14 +8384,27 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode,
&unsignedp, &reversep, &volatilep);
- tree orig_base = base;
-
+ if ((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);
+ }
+
+ STRIP_NOPS (base);
gcc_assert (offset == NULL_TREE || poly_int_tree_p (offset));
@@ -8402,10 +8419,6 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
*bitposp = bitpos;
*poffsetp = poffset;
- /* Set *BASE_REF if BASE was a dereferenced reference variable. */
- if (base_ref && orig_base != base)
- *base_ref = orig_base;
-
return base;
}
@@ -8422,6 +8435,48 @@ is_or_contains_p (tree expr, tree base_ptr)
return 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. */
@@ -8715,19 +8770,26 @@ static tree
build_struct_group (struct gimplify_omp_ctx *ctx,
enum omp_region_type region_type, enum tree_code code,
tree decl, unsigned int *flags, tree c,
- hash_map<tree, tree> *&struct_map_to_clause,
+ hash_map<tree_operand_hash, tree> *&struct_map_to_clause,
tree *&prev_list_p, tree *&list_p, bool *cont)
{
poly_offset_int coffset;
poly_int64 cbitpos;
- tree base_ref;
+ tree base_ind, base_ref;
+ tree *list_in_p = list_p, *prev_list_in_p = prev_list_p;
- tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
- &cbitpos, &coffset);
+ tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ind,
+ &base_ref, &cbitpos, &coffset);
gcc_assert (base == decl);
- splay_tree_node n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+ /* 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
@@ -8748,13 +8810,18 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
OMP_CLAUSE_SET_MAP_KIND (c, k);
has_attachments = true;
}
- if (n == NULL || (n->value & GOVD_MAP) == 0)
+ 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)
+ if (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) = decl;
@@ -8764,8 +8831,8 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
? 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, tree>;
- struct_map_to_clause->put (decl, l);
+ struct_map_to_clause = new hash_map<tree_operand_hash, tree>;
+ struct_map_to_clause->put (unshare_expr (decl), l);
if (ptr || attach_detach)
{
tree extra_node;
@@ -8784,15 +8851,50 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
}
else
list_p = insert_node_after (l, list_p);
- if (base_ref && code == OMP_TARGET)
+
+ /* 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_ind || base_ref) && (region_type & ORT_TARGET))
{
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- enum gomp_map_kind mkind = GOMP_MAP_FIRSTPRIVATE_REFERENCE;
+ enum gomp_map_kind mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE
+ : 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 (l);
- OMP_CLAUSE_CHAIN (l) = c2;
+ 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)
@@ -8805,11 +8907,18 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
{
tree *osc = struct_map_to_clause->get (decl);
tree *sc = NULL, *scp = NULL;
- if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach)
+ 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
- && 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)
@@ -8823,9 +8932,9 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
tree sc_decl = OMP_CLAUSE_DECL (*sc);
poly_offset_int offset;
poly_int64 bitpos;
- tree base = extract_base_bit_offset (sc_decl, NULL, &bitpos,
- &offset);
- if (base != decl)
+ tree base = extract_base_bit_offset (sc_decl, NULL, NULL,
+ &bitpos, &offset);
+ if (!base || !operand_equal_p (base, decl, 0))
break;
if (scp)
continue;
@@ -8966,8 +9075,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
{
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
- hash_map<tree, tree> *struct_map_to_clause = NULL;
- hash_set<tree> *struct_deref_set = NULL;
+ hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
tree *prev_list_p = NULL, *orig_list_p = list_p;
int handled_depend_iterators = -1;
int nowait = -1;
@@ -9484,99 +9592,38 @@ 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;
- 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)
- {
- while (TREE_CODE (decl) == COMPONENT_REF)
- {
- decl = TREE_OPERAND (decl, 0);
- 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;
- 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_ref = decl;
- decl = TREE_OPERAND (decl, 0);
- }
- }
- }
- else if (TREE_CODE (decl) == COMPONENT_REF)
+ if (TREE_CODE (decl) == COMPONENT_REF)
{
- while (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);
- }
- 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)))
- {
- 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;
+
+ STRIP_NOPS (decl);
}
/* An "attach/detach" operation on an update directive should
behave as a GOMP_MAP_ALWAYS_POINTER. Beware that
@@ -9585,7 +9632,7 @@ 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)
+ if (aggregate_base_p (decl)
&& 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
@@ -9640,7 +9687,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
}
- else if (add_decl)
+ else if (add_decl && DECL_P (add_decl))
{
decl = add_decl;
goto do_add_decl;
@@ -10284,8 +10331,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
gimplify_omp_ctxp = ctx;
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
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index cadca7e201f..15f1d63a836 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1504,8 +1504,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)
@@ -13046,6 +13048,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);
@@ -13064,6 +13067,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)))
@@ -13080,7 +13089,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..5662397507e
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/member-array-acc.C
@@ -0,0 +1,14 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+/* { dg-xfail-if "parse failure" { "*-*-*" } } */
+
+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..6671713d97c
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/member-array-omp.C
@@ -0,0 +1,14 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+/* { dg-xfail-if "parse failure" { "*-*-*" } } */
+
+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 {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/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..dda5b413136
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
@@ -0,0 +1,71 @@
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+ non-zero.
+
+ This version of the test omits the explicit mapping of the indirected
+ struct. */
+
+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
More information about the Gcc-patches
mailing list