public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Revert "Lambda capturing of pointers and references in target directives"
@ 2021-05-31 9:02 Chung-Lin Tang
0 siblings, 0 replies; only message in thread
From: Chung-Lin Tang @ 2021-05-31 9:02 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:a9458daf1368df52a03cb97968e5d546b4b993af
commit a9458daf1368df52a03cb97968e5d546b4b993af
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date: Wed May 26 19:11:39 2021 +0800
Revert "Lambda capturing of pointers and references in target directives"
This reverts commit 9228d5a2ce3d0f5c19f2068b1ad42dd4ba4936c7.
Diff:
---
gcc/cp/cp-tree.h | 2 +-
gcc/cp/lambda.c | 3 +
gcc/cp/parser.c | 2 +
gcc/cp/pt.c | 5 -
gcc/cp/semantics.c | 494 ++++++++----------------
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 94 -----
libgomp/testsuite/libgomp.c++/target-lambda-1.C | 86 -----
7 files changed, 168 insertions(+), 518 deletions(-)
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 33d87e9ee36..4ee5e0fd324 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -7567,7 +7567,7 @@ extern void finish_lambda_function (tree body);
extern bool regenerated_lambda_fn_p (tree);
extern tree most_general_lambda (tree);
extern tree finish_omp_target (location_t, tree, tree, bool);
-extern void finish_omp_target_clauses (location_t, tree, tree *);
+extern void set_omp_target_this_expr (tree);
/* in tree.c */
extern int cp_tree_operand_length (const_tree);
diff --git a/gcc/cp/lambda.c b/gcc/cp/lambda.c
index 16e2b4c18b4..23c1682893c 100644
--- a/gcc/cp/lambda.c
+++ b/gcc/cp/lambda.c
@@ -845,6 +845,9 @@ lambda_expr_this_capture (tree lambda, int add_capture_p)
type cast (_expr.cast_ 5.4) to the type of 'this'. [ The cast
ensures that the transformed expression is an rvalue. ] */
result = rvalue (result);
+
+ /* Acknowledge to OpenMP target that 'this' was referenced. */
+ set_omp_target_this_expr (result);
}
return result;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 916a96482d7..0927f29bfe7 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -42144,6 +42144,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
keep_next_level (true);
tree sb = begin_omp_structured_block (), ret;
unsigned save = cp_parser_begin_omp_structured_block (parser);
+ set_omp_target_this_expr (NULL_TREE);
switch (ccode)
{
case OMP_TEAMS:
@@ -42238,6 +42239,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
"#pragma omp target", pragma_tok);
c_omp_adjust_map_clauses (clauses, true);
keep_next_level (true);
+ set_omp_target_this_expr (NULL_TREE);
tree body = cp_parser_omp_structured_block (parser, if_p);
finish_omp_target (pragma_tok->location, clauses, body, false);
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 59f5513fc79..56217f86831 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -18978,11 +18978,6 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
t = copy_node (t);
OMP_BODY (t) = stmt;
OMP_CLAUSES (t) = tmp;
-
- if (TREE_CODE (t) == OMP_TARGET)
- finish_omp_target_clauses (EXPR_LOCATION (t), OMP_BODY (t),
- &OMP_CLAUSES (t));
-
if (TREE_CODE (t) == OMP_TARGET && OMP_TARGET_COMBINED (t))
{
tree teams = cp_walk_tree (&stmt, tsubst_find_omp_teams, NULL, NULL);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 66b44ddb590..2eeab999841 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -61,6 +61,11 @@ static hash_map<tree, tree> *omp_private_member_map;
static vec<tree> omp_private_member_vec;
static bool omp_private_member_ignore_next;
+/* Used for OpenMP target region 'this' references. */
+static tree omp_target_this_expr = NULL_TREE;
+
+static hash_map<tree, tree> omp_target_ptr_members_accessed;
+
/* Deferred Access Checking Overview
---------------------------------
@@ -2075,6 +2080,7 @@ tree
finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
{
gcc_assert (TREE_CODE (decl) == FIELD_DECL);
+ tree orig_object = object;
bool try_omp_private = !object && omp_private_member_map;
tree ret;
@@ -2113,6 +2119,14 @@ finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
return error_mark_node;
}
+ if (orig_object == NULL_TREE)
+ {
+ tree this_expr = TREE_OPERAND (object, 0);
+
+ /* Acknowledge to OpenMP target that 'this' was referenced. */
+ set_omp_target_this_expr (this_expr);
+ }
+
if (current_class_ptr)
TREE_USED (current_class_ptr) = 1;
if (processing_template_decl)
@@ -2173,6 +2187,13 @@ finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
if (v)
ret = convert_from_reference (*v);
}
+ else if (omp_target_this_expr
+ && TREE_TYPE (ret)
+ && POINTER_TYPE_P (TREE_TYPE (ret)))
+ {
+ if (omp_target_ptr_members_accessed.get (decl) == NULL)
+ omp_target_ptr_members_accessed.put (decl, ret);
+ }
return ret;
}
@@ -2914,6 +2935,9 @@ finish_this_expr (void)
/* The keyword 'this' is a prvalue expression. */
result = rvalue (result);
+ /* Acknowledge to OpenMP target that 'this' was referenced. */
+ set_omp_target_this_expr (result);
+
return result;
}
@@ -6573,6 +6597,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bool order_seen = false;
bool schedule_seen = false;
bool oacc_async = false;
+ bool indirect_ref_p = false;
bool indir_component_ref_p = false;
tree last_iterators = NULL_TREE;
bool last_iterators_remove = false;
@@ -7821,6 +7846,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
indir_component_ref_p = true;
STRIP_NOPS (t);
}
+ indirect_ref_p = false;
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && INDIRECT_REF_P (t))
+ {
+ t = TREE_OPERAND (t, 0);
+ indirect_ref_p = true;
+ STRIP_NOPS (t);
+ }
if (TREE_CODE (t) == COMPONENT_REF
&& ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
|| ort == C_ORT_ACC)
@@ -7856,6 +7889,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
t = TREE_OPERAND (t, 0);
+ if (INDIRECT_REF_P (t))
+ {
+ t = TREE_OPERAND (t, 0);
+ indir_component_ref_p = true;
+ STRIP_NOPS (t);
+ }
}
if (remove)
break;
@@ -7919,6 +7958,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|| (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER))
&& !indir_component_ref_p
+ && !indirect_ref_p
&& !cxx_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -8013,7 +8053,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else
{
- bitmap_set_bit (&map_head, DECL_UID (t));
+ if (!indirect_ref_p && !indir_component_ref_p)
+ bitmap_set_bit (&map_head, DECL_UID (t));
if (t != OMP_CLAUSE_DECL (c)
&& TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
bitmap_set_bit (&map_field_head, DECL_UID (t));
@@ -9038,126 +9079,26 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses)
return add_stmt (stmt);
}
-/* Used to walk OpenMP target directive body. */
-
-struct omp_target_walk_data
-{
- tree current_object;
- bool this_expr_accessed;
-
- hash_map<tree, tree> ptr_members_accessed;
- hash_set<tree> lambda_objects_accessed;
-
- tree current_closure;
- hash_set<tree> closure_vars_accessed;
-};
-
-static tree
-finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
+void
+set_omp_target_this_expr (tree this_val)
{
- tree t = *tp;
- struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr;
- tree current_object = data->current_object;
- tree current_closure = data->current_closure;
+ omp_target_this_expr = this_val;
- if (current_object)
- {
- tree this_expr = TREE_OPERAND (current_object, 0);
-
- if (operand_equal_p (t, this_expr))
- {
- data->this_expr_accessed = true;
- *walk_subtrees = 0;
- return NULL_TREE;
- }
-
- if (TREE_CODE (t) == COMPONENT_REF
- && POINTER_TYPE_P (TREE_TYPE (t))
- && operand_equal_p (TREE_OPERAND (t, 0), current_object)
- && TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL)
- {
- data->this_expr_accessed = true;
- tree fld = TREE_OPERAND (t, 1);
- if (data->ptr_members_accessed.get (fld) == NULL)
- {
- if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
- t = convert_from_reference (t);
- data->ptr_members_accessed.put (fld, t);
- }
- *walk_subtrees = 0;
- return NULL_TREE;
- }
- }
-
- /* When the current_function_decl is a lambda function, the closure object
- argument's type seems to not yet have fields layed out, so a recording
- of DECL_VALUE_EXPRs during the target body walk seems the only way to
- find them. */
- if (current_closure
- && (TREE_CODE (t) == VAR_DECL
- || TREE_CODE (t) == PARM_DECL
- || TREE_CODE (t) == RESULT_DECL)
- && DECL_HAS_VALUE_EXPR_P (t)
- && TREE_CODE (DECL_VALUE_EXPR (t)) == COMPONENT_REF
- && operand_equal_p (current_closure,
- TREE_OPERAND (DECL_VALUE_EXPR (t), 0)))
- {
- if (!data->closure_vars_accessed.contains (t))
- data->closure_vars_accessed.add (t);
- *walk_subtrees = 0;
- return NULL_TREE;
- }
-
- if (TREE_TYPE(t) && LAMBDA_TYPE_P (TREE_TYPE (t)))
- {
- tree lt = TREE_TYPE (t);
- gcc_assert (CLASS_TYPE_P (lt));
-
- if (!data->lambda_objects_accessed.contains (t))
- data->lambda_objects_accessed.add (t);
- *walk_subtrees = 0;
- return NULL_TREE;
- }
-
- return NULL_TREE;
+ if (omp_target_this_expr == NULL_TREE)
+ omp_target_ptr_members_accessed.empty ();
}
-void
-finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
+tree
+finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
{
- omp_target_walk_data data;
- data.this_expr_accessed = false;
+ tree last_inserted_clause = NULL_TREE;
- tree ct = current_nonlambda_class_type ();
- if (ct)
+ if (omp_target_this_expr)
{
- tree object = maybe_dummy_object (ct, NULL);
- object = maybe_resolve_dummy (object, true);
- data.current_object = object;
- }
- else
- data.current_object = NULL_TREE;
-
- if (DECL_LAMBDA_FUNCTION_P (current_function_decl))
- {
- tree closure = DECL_ARGUMENTS (current_function_decl);
- data.current_closure = build_indirect_ref (loc, closure, RO_UNARY_STAR);
- }
- else
- data.current_closure = NULL_TREE;
-
- cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data);
-
- auto_vec<tree, 16> new_clauses;
-
- if (data.this_expr_accessed)
- {
- tree omp_target_this_expr = TREE_OPERAND (data.current_object, 0);
-
/* See if explicit user-specified map(this[:]) clause already exists.
If not, we create an implicit map(tofrom:this[:1]) clause. */
tree *explicit_this_deref_map = NULL;
- for (tree *c = clauses_ptr; *c; c = &OMP_CLAUSE_CHAIN (*c))
+ for (tree *c = &clauses; *c; c = &OMP_CLAUSE_CHAIN (*c))
if (OMP_CLAUSE_CODE (*c) == OMP_CLAUSE_MAP
&& TREE_CODE (OMP_CLAUSE_DECL (*c)) == INDIRECT_REF
&& operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*c), 0),
@@ -9177,72 +9118,23 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
OMP_CLAUSE_DECL (c)
= build_indirect_ref (loc, closure, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure)));
- new_clauses.safe_push (c);
-
- tree closure_obj = OMP_CLAUSE_DECL (c);
- tree closure_type = TREE_TYPE (closure_obj);
-
- gcc_assert (LAMBDA_TYPE_P (closure_type)
- && CLASS_TYPE_P (closure_type));
+ = (processing_template_decl
+ ? NULL_TREE
+ : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure))));
tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
OMP_CLAUSE_DECL (c2) = closure;
OMP_CLAUSE_SIZE (c2) = size_zero_node;
- new_clauses.safe_push (c2);
+ OMP_CLAUSE_CHAIN (c2) = clauses;
+ OMP_CLAUSE_CHAIN (c) = c2;
+ last_inserted_clause = c2;
+ clauses = c;
STRIP_NOPS (omp_target_this_expr);
gcc_assert (DECL_HAS_VALUE_EXPR_P (omp_target_this_expr));
omp_target_this_expr = DECL_VALUE_EXPR (omp_target_this_expr);
- for (hash_set<tree>::iterator i = data.closure_vars_accessed.begin ();
- i != data.closure_vars_accessed.end (); ++i)
- {
- tree orig_decl = *i;
- tree closure_expr = DECL_VALUE_EXPR (orig_decl);
-
- if (TREE_CODE (TREE_TYPE (orig_decl)) == POINTER_TYPE)
- {
- /* this-pointer is processed outside this loop. */
- if (operand_equal_p (closure_expr, omp_target_this_expr))
- continue;
-
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
- OMP_CLAUSE_DECL (c)
- = build_indirect_ref (loc, closure_expr, RO_UNARY_STAR);
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
- new_clauses.safe_push (c);
-
- c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND
- (c, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
- OMP_CLAUSE_DECL (c) = closure_expr;
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- new_clauses.safe_push (c);
- }
- else if (TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE)
- {
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
- OMP_CLAUSE_DECL (c)
- = build1 (INDIRECT_REF,
- TREE_TYPE (TREE_TYPE (closure_expr)),
- closure_expr);
- OMP_CLAUSE_SIZE (c)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure_expr)));
- new_clauses.safe_push (c);
-
- c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
- OMP_CLAUSE_DECL (c) = closure_expr;
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- new_clauses.safe_push (c);
- }
- }
-
if (explicit_this_deref_map)
{
/* Transform *this into *__closure->this in maps. */
@@ -9257,13 +9149,12 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
OMP_CLAUSE_DECL (nc) = omp_target_this_expr;
OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_ALWAYS_POINTER);
- /* Unlink this two-map sequence away from the chain. */
- *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc);
-
/* Move map(*__closure->this) map(always_pointer:__closure->this)
sequence to right after __closure map. */
- new_clauses.safe_push (this_map);
- new_clauses.safe_push (nc);
+ *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc);
+ OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c2);
+ OMP_CLAUSE_CHAIN (c2) = this_map;
+ last_inserted_clause = nc;
}
else
{
@@ -9272,7 +9163,9 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
OMP_CLAUSE_DECL (c3)
= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c3)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
+ = (processing_template_decl
+ ? NULL_TREE
+ : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))));
tree c4 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_ALWAYS_POINTER);
@@ -9280,8 +9173,10 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
OMP_CLAUSE_DECL (c4) = omp_target_this_expr;
OMP_CLAUSE_SIZE (c4) = size_zero_node;
- new_clauses.safe_push (c3);
- new_clauses.safe_push (c4);
+ OMP_CLAUSE_CHAIN (c3) = c4;
+ OMP_CLAUSE_CHAIN (c4) = OMP_CLAUSE_CHAIN (c2);
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ last_inserted_clause = c4;
}
}
else
@@ -9295,177 +9190,112 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
OMP_CLAUSE_DECL (c)
= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
+ = (processing_template_decl
+ ? NULL_TREE
+ : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))));
tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
STRIP_NOPS (omp_target_this_expr);
OMP_CLAUSE_DECL (c2) = omp_target_this_expr;
OMP_CLAUSE_SIZE (c2) = size_zero_node;
-
- new_clauses.safe_push (c);
- new_clauses.safe_push (c2);
+ OMP_CLAUSE_CHAIN (c2) = clauses;
+ OMP_CLAUSE_CHAIN (c) = c2;
+ clauses = c;
+ last_inserted_clause = c2;
}
}
-
- if (!data.ptr_members_accessed.is_empty ())
- for (hash_map<tree, tree>::iterator i
- = data.ptr_members_accessed.begin ();
- i != data.ptr_members_accessed.end (); ++i)
- {
- /* For each referenced member that is of pointer or
- reference-to-pointer type, create the equivalent of
- map(alloc:this->ptr[:0]). */
- tree field_decl = (*i).first;
- tree ptr_member = (*i).second;
-
- for (tree c = *clauses_ptr; c; c = OMP_CLAUSE_CHAIN (c))
- {
- /* If map(this->ptr[:N] already exists, avoid creating another
- such map. */
- tree decl = OMP_CLAUSE_DECL (c);
- if ((TREE_CODE (decl) == INDIRECT_REF
- || TREE_CODE (decl) == MEM_REF)
- && operand_equal_p (TREE_OPERAND (decl, 0),
- ptr_member))
- goto next_ptr_member;
- }
-
- if (!cxx_mark_addressable (ptr_member))
- gcc_unreachable ();
-
- if (TREE_CODE (TREE_TYPE (field_decl)) == REFERENCE_TYPE)
- {
- /* For reference to pointers, we need to map the referenced
- pointer first for things to be correct. */
- tree ptr_member_type = TREE_TYPE (ptr_member);
-
- /* Map pointer target as zero-length array section. */
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
- OMP_CLAUSE_DECL (c)
- = build1 (INDIRECT_REF, TREE_TYPE (ptr_member_type), ptr_member);
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
-
- /* Map pointer to zero-length array section. */
- tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND
- (c2, GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION);
- OMP_CLAUSE_DECL (c2) = ptr_member;
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
-
- /* Attach reference-to-pointer field to pointer. */
- tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH);
- OMP_CLAUSE_DECL (c3) = TREE_OPERAND (ptr_member, 0);
- OMP_CLAUSE_SIZE (c3) = size_zero_node;
-
- new_clauses.safe_push (c);
- new_clauses.safe_push (c2);
- new_clauses.safe_push (c3);
- }
- else if (TREE_CODE (TREE_TYPE (field_decl)) == POINTER_TYPE)
- {
- /* Map pointer target as zero-length array section. */
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
- OMP_CLAUSE_DECL (c)
- = build_indirect_ref (loc, ptr_member, RO_UNARY_STAR);
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
-
- /* Attach zero-length array section to pointer. */
- tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND
- (c2, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
- OMP_CLAUSE_DECL (c2) = ptr_member;
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
-
- new_clauses.safe_push (c);
- new_clauses.safe_push (c2);
- }
- else
- gcc_unreachable ();
-
- next_ptr_member:
- ;
- }
+ omp_target_this_expr = NULL_TREE;
}
- if (!data.lambda_objects_accessed.is_empty ())
- {
- for (hash_set<tree>::iterator i = data.lambda_objects_accessed.begin ();
- i != data.lambda_objects_accessed.end (); ++i)
- {
- tree lobj = *i;
- tree lt = TREE_TYPE (lobj);
- gcc_assert (LAMBDA_TYPE_P (lt) && CLASS_TYPE_P (lt));
+ if (last_inserted_clause && !omp_target_ptr_members_accessed.is_empty ())
+ for (hash_map<tree, tree>::iterator i
+ = omp_target_ptr_members_accessed.begin ();
+ i != omp_target_ptr_members_accessed.end (); ++i)
+ {
+ /* For each referenced member that is of pointer or reference-to-pointer
+ type, create the equivalent of map(alloc:this->ptr[:0]). */
+ tree field_decl = (*i).first;
+ tree ptr_member = (*i).second;
- tree lc = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (lc, GOMP_MAP_TO);
- OMP_CLAUSE_DECL (lc) = lobj;
- OMP_CLAUSE_SIZE (lc) = TYPE_SIZE_UNIT (lt);
- new_clauses.truncate (0);
- new_clauses.safe_push (lc);
+ for (tree nc = OMP_CLAUSE_CHAIN (last_inserted_clause);
+ nc != NULL_TREE; nc = OMP_CLAUSE_CHAIN (nc))
+ {
+ /* If map(this->ptr[:N] already exists, avoid creating another
+ such map. */
+ tree decl = OMP_CLAUSE_DECL (nc);
+ if ((TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == MEM_REF)
+ && operand_equal_p (TREE_OPERAND (decl, 0),
+ ptr_member))
+ goto next_ptr_member;
+ }
- for (tree fld = TYPE_FIELDS (lt); fld; fld = DECL_CHAIN (fld))
- {
- if (TREE_CODE (TREE_TYPE (fld)) == POINTER_TYPE)
- {
- tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld),
- lobj, fld, NULL_TREE);
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
- OMP_CLAUSE_DECL (c)
- = build_indirect_ref (loc, exp, RO_UNARY_STAR);
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
- new_clauses.safe_push (c);
-
- c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND
- (c, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
- OMP_CLAUSE_DECL (c) = exp;
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- new_clauses.safe_push (c);
- }
- else if (TREE_CODE (TREE_TYPE (fld)) == REFERENCE_TYPE)
- {
- tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld),
- lobj, fld, NULL_TREE);
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
- OMP_CLAUSE_DECL (c)
- = build1 (INDIRECT_REF, TREE_TYPE (TREE_TYPE (exp)), exp);
- OMP_CLAUSE_SIZE (c)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (exp)));
- new_clauses.safe_push (c);
-
- c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
- OMP_CLAUSE_DECL (c) = exp;
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- new_clauses.safe_push (c);
- }
- }
- }
- }
+ if (!cxx_mark_addressable (ptr_member))
+ gcc_unreachable ();
- tree c = *clauses_ptr;
- for (int i = new_clauses.length () - 1; i >= 0; i--)
- {
- OMP_CLAUSE_CHAIN (new_clauses[i]) = c;
- c = new_clauses[i];
- }
- *clauses_ptr = c;
-}
+ if (TREE_CODE (TREE_TYPE (field_decl)) == REFERENCE_TYPE)
+ {
+ /* For reference to pointers, we need to map the referenced pointer
+ first for things to be correct. */
+ tree ptr_member_type = TREE_TYPE (ptr_member);
+
+ /* Map pointer target as zero-length array section. */
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_DECL (c)
+ = build1 (INDIRECT_REF, TREE_TYPE (ptr_member_type), ptr_member);
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+
+ /* Map pointer to zero-length array section. */
+ tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND
+ (c2, GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION);
+ OMP_CLAUSE_DECL (c2) = ptr_member;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+
+ /* Attach reference-to-pointer field to pointer. */
+ tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH);
+ OMP_CLAUSE_DECL (c3) = TREE_OPERAND (ptr_member, 0);
+ OMP_CLAUSE_SIZE (c3) = size_zero_node;
+
+ OMP_CLAUSE_CHAIN (c) = c2;
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (last_inserted_clause);
+
+ OMP_CLAUSE_CHAIN (last_inserted_clause) = c;
+ last_inserted_clause = c3;
+ }
+ else if (TREE_CODE (TREE_TYPE (field_decl)) == POINTER_TYPE)
+ {
+ /* Map pointer target as zero-length array section. */
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_DECL (c)
+ = build_indirect_ref (loc, ptr_member, RO_UNARY_STAR);
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+
+ /* Attach zero-length array section to pointer. */
+ tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND
+ (c2, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
+ OMP_CLAUSE_DECL (c2) = ptr_member;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+
+ OMP_CLAUSE_CHAIN (c) = c2;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (last_inserted_clause);
+ OMP_CLAUSE_CHAIN (last_inserted_clause) = c;
+ last_inserted_clause = c2;
+ }
+ else
+ gcc_unreachable ();
-tree
-finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
-{
- if (!processing_template_decl)
- finish_omp_target_clauses (loc, body, &clauses);
+ next_ptr_member:
+ ;
+ }
tree stmt = make_node (OMP_TARGET);
TREE_TYPE (stmt) = void_type_node;
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
deleted file mode 100644
index 7dceef80f47..00000000000
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ /dev/null
@@ -1,94 +0,0 @@
-// We use 'auto' without a function return type, so specify dialect here
-// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
-#include <cstdlib>
-#include <cstring>
-
-template <typename L>
-void
-omp_target_loop (int begin, int end, L loop)
-{
- #pragma omp target teams distribute parallel for
- for (int i = begin; i < end; i++)
- loop (i);
-}
-
-struct S
-{
- int a, len;
- int *ptr;
-
- auto merge_data_func (int *iptr, int &b)
- {
- auto fn = [=](void) -> bool
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- mapped = (ptr != NULL && iptr != NULL);
- if (mapped)
- {
- for (int i = 0; i < len; i++)
- ptr[i] += a + b + iptr[i];
- }
- }
- return mapped;
- };
- return fn;
- }
-};
-
-int x = 1;
-
-int main (void)
-{
- const int N = 10;
- int *data1 = new int[N];
- int *data2 = new int[N];
- memset (data1, 0xab, sizeof (int) * N);
- memset (data1, 0xcd, sizeof (int) * N);
-
- int val = 1;
- int &valref = val;
- #pragma omp target enter data map(alloc: data1[:N], data2[:N])
-
- omp_target_loop (0, N, [=](int i) { data1[i] = val; });
- omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; });
-
- #pragma omp target update from(data1[:N], data2[:N])
-
- for (int i = 0; i < N; i++)
- {
- if (data1[i] != 1) abort ();
- if (data2[i] != 2) abort ();
- }
-
- #pragma omp target exit data map(delete: data1[:N], data2[:N])
-
- int b = 8;
- S s = { 4, N, data1 };
- auto f = s.merge_data_func (data2, b);
-
- if (f ()) abort ();
-
- #pragma omp target enter data map(to: data1[:N])
- if (f ()) abort ();
-
- #pragma omp target enter data map(to: data2[:N])
- if (!f ()) abort ();
-
- #pragma omp target exit data map(from: data1[:N], data2[:N])
-
- for (int i = 0; i < N; i++)
- {
- if (data1[i] != 0xf) abort ();
- if (data2[i] != 2) abort ();
- }
-
- return 0;
-}
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-1.C b/libgomp/testsuite/libgomp.c++/target-lambda-1.C
deleted file mode 100644
index 06c6470b4ff..00000000000
--- a/libgomp/testsuite/libgomp.c++/target-lambda-1.C
+++ /dev/null
@@ -1,86 +0,0 @@
-#include <cstdlib>
-#include <cstring>
-
-template <typename L>
-void
-omp_target_loop (int begin, int end, L loop)
-{
- #pragma omp target teams distribute parallel for
- for (int i = begin; i < end; i++)
- loop (i);
-}
-
-struct S
-{
- int a, len;
- int *ptr;
-
- auto merge_data_func (int *iptr, int &b)
- {
- auto fn = [=](void) -> bool
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- mapped = (ptr != NULL && iptr != NULL);
- if (mapped)
- {
- for (int i = 0; i < len; i++)
- ptr[i] += a + b + iptr[i];
- }
- }
- return mapped;
- };
- return fn;
- }
-};
-
-int x = 1;
-
-int main (void)
-{
- const int N = 10;
- int *data1 = new int[N];
- int *data2 = new int[N];
- memset (data1, 0xab, sizeof (int) * N);
- memset (data1, 0xcd, sizeof (int) * N);
-
- int val = 1;
- int &valref = val;
- #pragma omp target enter data map(alloc: data1[:N], data2[:N])
-
- omp_target_loop (0, N, [=](int i) { data1[i] = val; });
- omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; });
-
- #pragma omp target update from(data1[:N], data2[:N])
-
- for (int i = 0; i < N; i++)
- {
- if (data1[i] != 1) abort ();
- if (data2[i] != 2) abort ();
- }
-
- #pragma omp target exit data map(delete: data1[:N], data2[:N])
-
- int b = 8;
- S s = { 4, N, data1 };
- auto f = s.merge_data_func (data2, b);
-
- if (f ()) abort ();
-
- #pragma omp target enter data map(to: data1[:N])
- if (f ()) abort ();
-
- #pragma omp target enter data map(to: data2[:N])
- if (!f ()) abort ();
-
- #pragma omp target exit data map(from: data1[:N], data2[:N])
-
- for (int i = 0; i < N; i++)
- {
- if (data1[i] != 0xf) abort ();
- if (data2[i] != 2) abort ();
- }
-
- return 0;
-}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-31 9:02 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-31 9:02 [gcc/devel/omp/gcc-11] Revert "Lambda capturing of pointers and references in target directives" Chung-Lin Tang
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).