public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Lambda capturing of pointers and references in target directives
@ 2021-05-13 16:20 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:20 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:9228d5a2ce3d0f5c19f2068b1ad42dd4ba4936c7
commit 9228d5a2ce3d0f5c19f2068b1ad42dd4ba4936c7
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date: Fri Mar 19 02:29:20 2021 +0800
Lambda capturing of pointers and references in target directives
This patch implements proper lambda capturing of pointer and reference variables
as specified in OpenMP 5.0. We map the entire closure object as a to-map,
attach pointers to zero-length array sections, and perform mapping of
references.
2021-03-18 Chung-Lin Tang <cltang@codesourcery.com>
gcc/cp/ChangeLog:
* cp-tree.h (set_omp_target_this_expr): Delete.
(finish_omp_target_clauses): New prototype.
* lambda.c (lambda_expr_this_capture): Remove call to
set_omp_target_this_expr.
* parser.c (cp_parser_omp_target): Likewise.
* pt.c (tsubst_expr): Add call to finish_omp_target_clauses for target
directives.
* semantics.c (omp_target_this_expr): Delete.
(omp_target_ptr_members_accessed): Delete.
(finish_non_static_data_member): Remove call to
set_omp_target_this_expr. Remove use of omp_target_ptr_members_accessed.
(finish_this_expr): Remove call to set_omp_target_this_expr.
(struct omp_target_walk_data): New struct for walking over
target-directive tree body.
(finish_omp_target_clauses_r): New function for tree walk.
(finish_omp_target_clauses): New function, with code factored out from
finish_omp_target. Add lambda object handling case.
(finish_omp_target): Factor code out and adjust to use
finish_omp_target_clauses.
(finish_omp_clauses): Revert prior "Adjustments to allow '*ptr' and
'ptr->member' cases in map clausess.", since not needed with new
organization of target-directive clause processing.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-lambda-1.C: New test.
libgomp/testsuite/ChangeLog:
* libgomp.c++/target-lambda-1.C: New test.
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, 518 insertions(+), 168 deletions(-)
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index d6c6e2aad46..0f1cf8d4576 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -7565,7 +7565,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 set_omp_target_this_expr (tree);
+extern void finish_omp_target_clauses (location_t, tree, 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 23c1682893c..16e2b4c18b4 100644
--- a/gcc/cp/lambda.c
+++ b/gcc/cp/lambda.c
@@ -845,9 +845,6 @@ 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 f9a2092e0e1..68af2efb443 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -41998,7 +41998,6 @@ 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:
@@ -42093,7 +42092,6 @@ 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 bd04e856181..edb26fdde1c 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -18978,6 +18978,11 @@ 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 46eb5807eb9..ba7d869cdb5 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -61,11 +61,6 @@ 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
---------------------------------
@@ -2080,7 +2075,6 @@ 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;
@@ -2119,14 +2113,6 @@ 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)
@@ -2187,13 +2173,6 @@ 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;
}
@@ -2935,9 +2914,6 @@ 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;
}
@@ -6600,7 +6576,6 @@ 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;
@@ -7785,14 +7760,6 @@ 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)
@@ -7828,12 +7795,6 @@ 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;
@@ -7897,7 +7858,6 @@ 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
@@ -7982,8 +7942,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else
{
- if (!indirect_ref_p && !indir_component_ref_p)
- bitmap_set_bit (&map_head, DECL_UID (t));
+ 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));
@@ -9008,26 +8967,126 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses)
return add_stmt (stmt);
}
-void
-set_omp_target_this_expr (tree this_val)
+/* Used to walk OpenMP target directive body. */
+
+struct omp_target_walk_data
{
- omp_target_this_expr = this_val;
+ tree current_object;
+ bool this_expr_accessed;
+
+ hash_map<tree, tree> ptr_members_accessed;
+ hash_set<tree> lambda_objects_accessed;
- if (omp_target_this_expr == NULL_TREE)
- omp_target_ptr_members_accessed.empty ();
+ tree current_closure;
+ hash_set<tree> closure_vars_accessed;
+};
+
+static tree
+finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
+{
+ 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;
+
+ 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;
}
-tree
-finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
+void
+finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
{
- tree last_inserted_clause = NULL_TREE;
+ omp_target_walk_data data;
+ data.this_expr_accessed = false;
- if (omp_target_this_expr)
+ tree ct = current_nonlambda_class_type ();
+ if (ct)
{
+ 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; *c; c = &OMP_CLAUSE_CHAIN (*c))
+ for (tree *c = clauses_ptr; *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),
@@ -9047,23 +9106,72 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
OMP_CLAUSE_DECL (c)
= build_indirect_ref (loc, closure, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c)
- = (processing_template_decl
- ? NULL_TREE
- : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure))));
+ = 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));
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;
- OMP_CLAUSE_CHAIN (c2) = clauses;
- OMP_CLAUSE_CHAIN (c) = c2;
- last_inserted_clause = c2;
- clauses = c;
+ new_clauses.safe_push (c2);
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. */
@@ -9078,12 +9186,13 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
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. */
- *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;
+ new_clauses.safe_push (this_map);
+ new_clauses.safe_push (nc);
}
else
{
@@ -9092,9 +9201,7 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
OMP_CLAUSE_DECL (c3)
= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c3)
- = (processing_template_decl
- ? NULL_TREE
- : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))));
+ = 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);
@@ -9102,10 +9209,8 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
OMP_CLAUSE_DECL (c4) = omp_target_this_expr;
OMP_CLAUSE_SIZE (c4) = size_zero_node;
- OMP_CLAUSE_CHAIN (c3) = c4;
- OMP_CLAUSE_CHAIN (c4) = OMP_CLAUSE_CHAIN (c2);
- OMP_CLAUSE_CHAIN (c2) = c3;
- last_inserted_clause = c4;
+ new_clauses.safe_push (c3);
+ new_clauses.safe_push (c4);
}
}
else
@@ -9119,112 +9224,177 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
OMP_CLAUSE_DECL (c)
= build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c)
- = (processing_template_decl
- ? NULL_TREE
- : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))));
+ = 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;
- OMP_CLAUSE_CHAIN (c2) = clauses;
- OMP_CLAUSE_CHAIN (c) = c2;
- clauses = c;
- last_inserted_clause = c2;
+
+ new_clauses.safe_push (c);
+ new_clauses.safe_push (c2);
}
}
- omp_target_this_expr = NULL_TREE;
- }
-
- 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;
- for (tree nc = OMP_CLAUSE_CHAIN (last_inserted_clause);
- nc != NULL_TREE; nc = OMP_CLAUSE_CHAIN (nc))
+ 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)
{
- /* 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 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;
- if (!cxx_mark_addressable (ptr_member))
- gcc_unreachable ();
+ 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 (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;
+ 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:
+ ;
}
- else
- gcc_unreachable ();
+ }
- next_ptr_member:
- ;
- }
+ 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));
+
+ 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 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);
+ }
+ }
+ }
+ }
+
+ 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;
+}
+
+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);
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
new file mode 100644
index 00000000000..7dceef80f47
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -0,0 +1,94 @@
+// 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
new file mode 100644
index 00000000000..06c6470b4ff
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-1.C
@@ -0,0 +1,86 @@
+#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-13 16:20 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:20 [gcc/devel/omp/gcc-11] Lambda capturing of pointers and references in target directives Kwok Yeung
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).