public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Chung-Lin Tang <cltang@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] Revert "Lambda capturing of pointers and references in target directives" Date: Mon, 31 May 2021 09:02:27 +0000 (GMT) [thread overview] Message-ID: <20210531090227.D714638515D4@sourceware.org> (raw) 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; -}
reply other threads:[~2021-05-31 9:02 UTC|newest] Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
Reply instructions: You may reply publicly to this message via plain-text email using any one of the following methods: * Save the following mbox file, import it into your mail client, and reply-to-all from there: mbox Avoid top-posting and favor interleaved quoting: https://en.wikipedia.org/wiki/Posting_style#Interleaved_style * Reply using the --to, --cc, and --in-reply-to switches of git-send-email(1): git send-email \ --in-reply-to=20210531090227.D714638515D4@sourceware.org \ --to=cltang@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /path/to/YOUR_REPLY https://kernel.org/pub/software/scm/git/docs/git-send-email.html * If your mail client supports setting the In-Reply-To header via mailto: links, try the mailto: linkBe sure your reply has a Subject: header at the top and a blank line before the message body.
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).