* [PATCH, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v4]
@ 2021-06-18 14:25 Chung-Lin Tang
2021-06-24 13:15 ` Jakub Jelinek
0 siblings, 1 reply; 6+ messages in thread
From: Chung-Lin Tang @ 2021-06-18 14:25 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek, Catherine Moore
[-- Attachment #1: Type: text/plain, Size: 4586 bytes --]
Hi Jakub,
this patch is the "v4" version of my PR92120 patch, v3 was here:
https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570886.html
(there I listed the various patches from devel/omp/gcc-10 branch that was combined,
which I won't repeat here).
Basically this v4 adds fixes for lambda capture, which was already pushed to
devel/omp/gcc-11 yesterday:
https://gcc.gnu.org/pipermail/gcc-patches/2021-June/572988.html
I have attached both the combined v4 version, and the v3-to-v4 diff.
Tested on x86_64-linux with nvptx offloading, seeking for approval to trunk.
Thanks,
Chung-Lin
gcc/cp/
* cp-tree.h (finish_omp_target): New declaration.
(finish_omp_target_clauses): Likewise.
* parser.c (cp_parser_omp_clause_map): Adjust call to
cp_parser_omp_var_list_no_open to set 'allow_deref' argument to true.
(cp_parser_omp_target): Factor out code, adjust into calls to new
function finish_omp_target.
* pt.c (tsubst_expr): Add call to finish_omp_target_clauses for
OMP_TARGET case.
* semantics.c (handle_omp_array_sections_1): Add handling to create
'this->member' from 'member' FIELD_DECL.
(handle_omp_array_sections): Likewise.
(finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP
map clauses. Handle 'A->member' case in map clauses.
(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.
(finish_omp_target): New function.
gcc/c/
* c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in
call to c_parser_omp_variable_list to 'true'.
* c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in
array base handling.
(c_finish_omp_clauses): Handle 'A->member' case in map clauses.
gcc/
* gimplify.c ("tree-hash-traits.h"): Add include.
(gimplify_scan_omp_clauses): Change struct_map_to_clause to type
hash_map<tree_operand, tree> *. Adjust struct map handling to handle
cases of *A and A->B expressions. Under !DECL_P case of
GOMP_CLAUSE_MAP handling, add STRIP_NOPS for indir_p case, add to
struct_deref_set for map(*ptr_to_struct) cases. Add MEM_REF case when
handling component_ref_p case. Add unshare_expr and gimplification
when created GOMP_MAP_STRUCT is not a DECL. Add code to add
firstprivate pointer for *pointer-to-struct case.
(gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for
exit data directives code to earlier position.
* omp-low.c (lower_omp_target):
Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds.
* tree-pretty-print.c (dump_omp_clause): Likewise.
gcc/testsuite/
* gcc.dg/gomp/target-3.c: New testcase.
* g++.dg/gomp/target-3.C: New testcase.
* g++.dg/gomp/target-lambda-1.C: New testcase.
* g++.dg/gomp/target-lambda-2.C: New testcase.
* g++.dg/gomp/target-this-1.C: New testcase.
* g++.dg/gomp/target-this-2.C: New testcase.
* g++.dg/gomp/target-this-3.C: New testcase.
* g++.dg/gomp/target-this-4.C: New testcase.
* g++.dg/gomp/target-this-5.C: New testcase.
* g++.dg/gomp/this-2.C: Adjust testcase.
include/
* gomp-constants.h (enum gomp_map_kind):
Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds.
(GOMP_MAP_POINTER_P):
Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION.
libgomp/
* libgomp.h (gomp_attach_pointer): Add bool parameter.
* oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer.
(goacc_enter_data_internal): Likewise.
* target.c (gomp_map_vars_existing): Update assert condition to
include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION.
(gomp_map_pointer): Add 'bool allow_zero_length_array_sections'
parameter, add support for mapping a pointer with NULL target.
(gomp_attach_pointer): Add 'bool allow_zero_length_array_sections'
parameter, add support for attaching a pointer with NULL target.
(gomp_map_vars_internal): Update calls to gomp_map_pointer and
gomp_attach_pointer, add handling for
GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases.
* testsuite/libgomp.c++/target-23.C: New testcase.
* testsuite/libgomp.c++/target-lambda-1.C: New testcase.
* testsuite/libgomp.c++/target-lambda-2.C: New testcase.
* testsuite/libgomp.c++/target-this-1.C: New testcase.
* testsuite/libgomp.c++/target-this-2.C: New testcase.
* testsuite/libgomp.c++/target-this-3.C: New testcase.
* testsuite/libgomp.c++/target-this-4.C: New testcase.
* testsuite/libgomp.c++/target-this-5.C: New testcase.
[-- Attachment #2: pr92120-v4.patch --]
[-- Type: text/plain, Size: 76420 bytes --]
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b90710cba2f..f1f7c43ad30 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15851,7 +15851,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
c_parser_consume_token (parser);
}
- nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list);
+ nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
+ true);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 845d50f3310..f64958eb7ba 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13104,6 +13104,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && TREE_CODE (t) == MEM_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ }
if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
@@ -13952,6 +13958,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
tree ordered_clause = NULL_TREE;
tree schedule_clause = NULL_TREE;
bool oacc_async = false;
+ bool indir_component_ref_p = false;
tree last_iterators = NULL_TREE;
bool last_iterators_remove = false;
tree *nogroup_seen = NULL;
@@ -14736,6 +14743,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
while (TREE_CODE (t) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == MEM_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ }
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -14802,6 +14814,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bias) to zero here, so it is not set erroneously to the pointer
size later on in gimplify.c. */
OMP_CLAUSE_SIZE (c) = size_zero_node;
+ indir_component_ref_p = false;
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF)
+ {
+ t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
+ indir_component_ref_p = true;
+ STRIP_NOPS (t);
+ }
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
@@ -14874,6 +14895,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER))
+ && !indir_component_ref_p
&& !c_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -14932,8 +14954,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
- && (ort != C_ORT_OMP
- || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
+ && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 36f99ccf189..6fa7fc26c16 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -7597,6 +7597,8 @@ extern tree start_lambda_function (tree fn, tree lambda_expr);
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 *);
/* in tree.c */
extern int cp_tree_operand_length (const_tree);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d57ddc4560d..9076bb2fb7e 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -38201,7 +38201,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
}
nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
- NULL);
+ NULL, true);
for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -42225,8 +42225,6 @@ static bool
cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
enum pragma_context context, bool *if_p)
{
- tree *pc = NULL, stmt;
-
if (flag_openmp)
omp_requires_mask
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
@@ -42330,16 +42328,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
}
}
- tree stmt = make_node (OMP_TARGET);
- TREE_TYPE (stmt) = void_type_node;
- OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
- c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
- OMP_TARGET_BODY (stmt) = body;
- OMP_TARGET_COMBINED (stmt) = 1;
- SET_EXPR_LOCATION (stmt, pragma_tok->location);
- add_stmt (stmt);
- pc = &OMP_TARGET_CLAUSES (stmt);
- goto check_clauses;
+ c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true);
+ finish_omp_target (pragma_tok->location,
+ cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true);
+ return true;
}
else if (!flag_openmp) /* flag_openmp_simd */
{
@@ -42376,49 +42368,13 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
return false;
}
- stmt = make_node (OMP_TARGET);
- TREE_TYPE (stmt) = void_type_node;
-
- OMP_TARGET_CLAUSES (stmt)
- = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
- "#pragma omp target", pragma_tok);
- c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
-
- pc = &OMP_TARGET_CLAUSES (stmt);
+ tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+ "#pragma omp target", pragma_tok);
+ c_omp_adjust_map_clauses (clauses, true);
keep_next_level (true);
- OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p);
-
- SET_EXPR_LOCATION (stmt, pragma_tok->location);
- add_stmt (stmt);
+ tree body = cp_parser_omp_structured_block (parser, if_p);
-check_clauses:
- while (*pc)
- {
- if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
- switch (OMP_CLAUSE_MAP_KIND (*pc))
- {
- case GOMP_MAP_TO:
- case GOMP_MAP_ALWAYS_TO:
- case GOMP_MAP_FROM:
- case GOMP_MAP_ALWAYS_FROM:
- case GOMP_MAP_TOFROM:
- case GOMP_MAP_ALWAYS_TOFROM:
- case GOMP_MAP_ALLOC:
- case GOMP_MAP_FIRSTPRIVATE_POINTER:
- case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
- case GOMP_MAP_ALWAYS_POINTER:
- case GOMP_MAP_ATTACH_DETACH:
- break;
- default:
- error_at (OMP_CLAUSE_LOCATION (*pc),
- "%<#pragma omp target%> with map-type other "
- "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
- "on %<map%> clause");
- *pc = OMP_CLAUSE_CHAIN (*pc);
- continue;
- }
- pc = &OMP_CLAUSE_CHAIN (*pc);
- }
+ finish_omp_target (pragma_tok->location, clauses, body, false);
return true;
}
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 15947b2c812..65621fd6b6f 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -18892,6 +18892,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 384c54bd025..455620c0663 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4990,6 +4990,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
if (error_operand_p (t))
return error_mark_node;
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && TREE_CODE (t) == FIELD_DECL)
+ t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
@@ -5018,8 +5021,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
- if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF)
- t = TREE_OPERAND (t, 0);
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && TREE_CODE (t) == INDIRECT_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ }
}
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
@@ -5043,6 +5050,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
else if (ort == C_ORT_OMP
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM
&& TREE_CODE (t) == PARM_DECL
&& DECL_ARTIFICIAL (t)
&& DECL_NAME (t) == this_identifier
@@ -5567,6 +5577,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
}
OMP_CLAUSE_DECL (c) = first;
OMP_CLAUSE_SIZE (c) = size;
+ if (TREE_CODE (t) == FIELD_DECL)
+ t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
@@ -6583,6 +6595,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 indir_component_ref_p = false;
tree last_iterators = NULL_TREE;
bool last_iterators_remove = false;
/* 1 if normal/task reduction has been seen, -1 if inscan reduction
@@ -7723,6 +7736,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = TREE_OPERAND (t, 0);
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == INDIRECT_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ STRIP_NOPS (t);
+ }
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (t))
@@ -7795,10 +7813,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_DECL (c) = t;
}
+ indir_component_ref_p = false;
if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
&& TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
- t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
+ {
+ t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
+ indir_component_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)
@@ -7847,6 +7870,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
goto handle_map_references;
}
}
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && !processing_template_decl
+ && TREE_CODE (t) == FIELD_DECL)
+ {
+ OMP_CLAUSE_DECL (c) = finish_non_static_data_member (t, NULL_TREE,
+ NULL_TREE);
+ break;
+ }
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl && TREE_CODE (t) != OVERLOAD)
@@ -7873,7 +7904,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if (ort != C_ORT_ACC && t == current_class_ptr)
+ else if (ort != C_ORT_ACC
+ && ort != C_ORT_OMP
+ && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -7886,6 +7919,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
&& (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER))
+ && !indir_component_ref_p
&& !cxx_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -9003,6 +9037,493 @@ 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;
+
+ hash_set<tree> local_decls;
+};
+
+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_CODE (t) == BIND_EXPR)
+ {
+ tree block = BIND_EXPR_BLOCK (t);
+ for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var))
+ if (!data->local_decls.contains (var))
+ data->local_decls.add (var);
+ 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)
+ /* Do not prepare to create target maps for locally declared
+ lambdas or anonymous ones. */
+ && !data->local_decls.contains (t)
+ && TREE_CODE (t) != TARGET_EXPR)
+ data->lambda_objects_accessed.add (t);
+ *walk_subtrees = 0;
+ return NULL_TREE;
+ }
+
+ return NULL_TREE;
+}
+
+void
+finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
+{
+ omp_target_walk_data data;
+ data.this_expr_accessed = false;
+
+ 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_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),
+ omp_target_this_expr))
+ {
+ explicit_this_deref_map = c;
+ break;
+ }
+
+ if (DECL_LAMBDA_FUNCTION_P (current_function_decl))
+ {
+ /* For lambda functions, we need to first create a copy of the
+ __closure object. */
+ tree closure = DECL_ARGUMENTS (current_function_decl);
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+ 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));
+
+ 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);
+
+ 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. */
+ tree this_map = *explicit_this_deref_map;
+ OMP_CLAUSE_DECL (this_map)
+ = build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
+
+ tree nc = OMP_CLAUSE_CHAIN (this_map);
+ gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER));
+ 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);
+ }
+ else
+ {
+ tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_TOFROM);
+ 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)));
+
+ tree c4 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_ALWAYS_POINTER);
+
+ 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);
+ }
+ }
+ else
+ {
+ /* For the non-lambda case, we only need to create map(this[:1]) when
+ it's not present, no transforming needed. */
+ if (!explicit_this_deref_map)
+ {
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+ 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)));
+
+ 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);
+ }
+ }
+
+ 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:
+ ;
+ }
+ }
+
+ 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;
+ if (TREE_CODE (lobj) == TARGET_EXPR)
+ lobj = TREE_OPERAND (lobj, 0);
+
+ 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_TOFROM);
+ 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;
+ OMP_TARGET_CLAUSES (stmt) = clauses;
+ OMP_TARGET_BODY (stmt) = body;
+ OMP_TARGET_COMBINED (stmt) = combined_p;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ tree c = clauses;
+ while (c)
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ break;
+ default:
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<#pragma omp target%> with map-type other "
+ "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+ "on %<map%> clause");
+ break;
+ }
+ c = OMP_CLAUSE_CHAIN (c);
+ }
+ return add_stmt (stmt);
+}
+
tree
finish_omp_parallel (tree clauses, tree body)
{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 2606998152d..86d328c99fc 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"
@@ -8738,7 +8739,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_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
hash_set<tree> *struct_deref_set = NULL;
tree *prev_list_p = NULL, *orig_list_p = list_p;
int handled_depend_iterators = -1;
@@ -9176,7 +9177,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
GOVD_FIRSTPRIVATE | GOVD_SEEN);
}
- if (!DECL_P (decl))
+ if (TREE_CODE (decl) == TARGET_EXPR)
+ {
+ if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
+ is_gimple_lvalue, fb_lvalue)
+ == GS_ERROR)
+ remove = true;
+ }
+ else if (!DECL_P (decl))
{
tree d = decl, *pd;
if (TREE_CODE (d) == ARRAY_REF)
@@ -9192,12 +9200,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
&& TREE_CODE (decl) == INDIRECT_REF
&& TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == REFERENCE_TYPE))
+ == REFERENCE_TYPE)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION))
{
pd = &TREE_OPERAND (decl, 0);
decl = TREE_OPERAND (decl, 0);
}
bool indir_p = false;
+ bool component_ref_p = false;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -9208,6 +9219,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
while (TREE_CODE (decl) == COMPONENT_REF)
{
decl = TREE_OPERAND (decl, 0);
+ component_ref_p = true;
if (((TREE_CODE (decl) == MEM_REF
&& integer_zerop (TREE_OPERAND (decl, 1)))
|| INDIRECT_REF_P (decl))
@@ -9216,6 +9228,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
{
indir_p = true;
decl = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (decl);
}
if (TREE_CODE (decl) == INDIRECT_REF
&& DECL_P (TREE_OPERAND (decl, 0))
@@ -9227,8 +9240,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
}
}
- else if (TREE_CODE (decl) == COMPONENT_REF)
+ else if (TREE_CODE (decl) == COMPONENT_REF
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
{
+ component_ref_p = true;
while (TREE_CODE (decl) == COMPONENT_REF)
decl = TREE_OPERAND (decl, 0);
if (TREE_CODE (decl) == INDIRECT_REF
@@ -9298,7 +9314,10 @@ 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 ((DECL_P (decl)
+ || (component_ref_p
+ && (INDIRECT_REF_P (decl)
+ || TREE_CODE (decl) == MEM_REF)))
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
@@ -9355,7 +9374,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
gcc_assert (base == decl);
splay_tree_node n
- = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+ = (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)
@@ -9381,7 +9403,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
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);
@@ -9392,7 +9418,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (base_ref)
OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
else
- OMP_CLAUSE_DECL (l) = decl;
+ {
+ OMP_CLAUSE_DECL (l) = unshare_expr (decl);
+ if (!DECL_P (OMP_CLAUSE_DECL (l))
+ && (gimplify_expr (&OMP_CLAUSE_DECL (l),
+ pre_p, NULL, is_gimple_lvalue,
+ fb_lvalue)
+ == GS_ERROR))
+ {
+ remove = true;
+ break;
+ }
+ }
OMP_CLAUSE_SIZE (l)
= (!attach
? size_int (1)
@@ -9400,7 +9437,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
? 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
+ = new hash_map<tree_operand_hash, tree>;
struct_map_to_clause->put (decl, l);
if (ptr || attach_detach)
{
@@ -9434,7 +9472,32 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
flags |= GOVD_SEEN;
if (has_attachments)
flags |= GOVD_MAP_HAS_ATTACHMENTS;
- goto do_add_decl;
+
+ /* If this is a *pointer-to-struct expression, make sure a
+ firstprivate map of the base-pointer exists. */
+ if (component_ref_p
+ && ((TREE_CODE (decl) == MEM_REF
+ && integer_zerop (TREE_OPERAND (decl, 1)))
+ || INDIRECT_REF_P (decl))
+ && DECL_P (TREE_OPERAND (decl, 0))
+ && !splay_tree_lookup (ctx->variables,
+ ((splay_tree_key)
+ TREE_OPERAND (decl, 0))))
+ {
+ decl = TREE_OPERAND (decl, 0);
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ enum gomp_map_kind mkind
+ = GOMP_MAP_FIRSTPRIVATE_POINTER;
+ OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+ OMP_CLAUSE_DECL (c2) = decl;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = c2;
+ }
+
+ if (DECL_P (decl))
+ goto do_add_decl;
}
else if (struct_map_to_clause)
{
@@ -9543,6 +9606,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
else if (*sc != c)
{
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+ fb_lvalue)
+ == GS_ERROR)
+ {
+ remove = true;
+ break;
+ }
*list_p = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = *sc;
*sc = c;
@@ -9570,6 +9640,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
+ /* If this was of the form map(*pointer_to_struct), then the
+ 'pointer_to_struct' DECL should be considered deref'ed. */
+ if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c))
+ || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c)))
+ && INDIRECT_REF_P (orig_decl)
+ && DECL_P (TREE_OPERAND (orig_decl, 0))
+ && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE)
+ {
+ tree ptr = TREE_OPERAND (orig_decl, 0);
+ if (!struct_deref_set || !struct_deref_set->contains (ptr))
+ {
+ if (!struct_deref_set)
+ struct_deref_set = new hash_set<tree> ();
+ struct_deref_set->add (ptr);
+ }
+ }
+
if (!remove
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
@@ -10846,6 +10934,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
}
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+ && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA))
+ {
+ remove = true;
+ break;
+ }
if (!DECL_P (decl))
{
if ((ctx->region_type & ORT_TARGET) != 0
@@ -10892,10 +10986,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
= OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c));
}
}
- else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- && (code == OMP_TARGET_EXIT_DATA
- || code == OACC_EXIT_DATA))
- remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 2cc2a186080..0307e97512b 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12268,6 +12268,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH:
case GOMP_MAP_DETACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C
new file mode 100644
index 00000000000..f4d40ec8e4b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-3.C
@@ -0,0 +1,36 @@
+// { dg-do compile }
+// { dg-options "-fopenmp -fdump-tree-gimple" }
+
+struct S
+{
+ int a, b;
+ void bar (int);
+};
+
+void
+S::bar (int x)
+{
+ #pragma omp target map (alloc: a, b)
+ ;
+ #pragma omp target enter data map (alloc: a, b)
+}
+
+template <int N>
+struct T
+{
+ int a, b;
+ void bar (int);
+};
+
+template <int N>
+void
+T<N>::bar (int x)
+{
+ #pragma omp target map (alloc: a, b)
+ ;
+ #pragma omp target enter data map (alloc: a, b)
+}
+
+template struct T<0>;
+
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
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/gcc/testsuite/g++.dg/gomp/target-lambda-2.C b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
new file mode 100644
index 00000000000..bdf2564cd04
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
@@ -0,0 +1,35 @@
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+ int X, Y;
+ #pragma omp target map(from: X, Y)
+ {
+ int x = 0, y = 0;
+
+ for (int i = 0; i < N; i++)
+ [&] (int v) { x += v; } (i);
+
+ auto yinc = [&y] { y++; };
+ for (int i = 0; i < N; i++)
+ yinc ();
+
+ X = x;
+ Y = y;
+ }
+
+ int Xs = 0;
+ for (int i = 0; i < N; i++)
+ Xs += i;
+ if (X != Xs)
+ abort ();
+
+ if (Y != N)
+ abort ();
+}
+
+/* Make sure lambda objects do NOT appear in target maps. */
+/* { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(from:Y \[len: [0-9]+\]\) map\(from:X \[len: [0-9]+\]\)$} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-1.C b/gcc/testsuite/g++.dg/gomp/target-this-1.C
new file mode 100644
index 00000000000..de93a3e5e57
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-1.C
@@ -0,0 +1,33 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+extern "C" void abort ();
+
+struct S
+{
+ int a, b, c, d;
+
+ int sum (void)
+ {
+ int val = 0;
+ val += a + b + this->c + this->d;
+ return val;
+ }
+
+ int sum_offload (void)
+ {
+ int val = 0;
+ #pragma omp target map(val)
+ val += a + b + this->c + this->d;
+ return val;
+ }
+};
+
+int main (void)
+{
+ S s = { 1, 2, 3, 4 };
+ if (s.sum () != s.sum_offload ())
+ abort ();
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C
new file mode 100644
index 00000000000..679c85a54dd
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -0,0 +1,49 @@
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-do compile }
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+
+extern "C" void abort ();
+
+struct T
+{
+ int x, y;
+
+ auto sum_func (int n)
+ {
+ auto fn = [=](int m) -> int
+ {
+ int v;
+ v = (x + y) * n + m;
+ return v;
+ };
+ return fn;
+ }
+
+ auto sum_func_offload (int n)
+ {
+ auto fn = [=](int m) -> int
+ {
+ int v;
+ #pragma omp target map(from:v)
+ v = (x + y) * n + m;
+ return v;
+ };
+ return fn;
+ }
+
+};
+
+int main (void)
+{
+ T a = { 1, 2 };
+
+ auto s1 = a.sum_func (3);
+ auto s2 = a.sum_func_offload (3);
+
+ if (s1 (1) != s2 (1))
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, 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\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
new file mode 100644
index 00000000000..08568f9284c
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -0,0 +1,105 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+#include <cstdlib>
+#include <cstring>
+extern "C" void abort ();
+
+struct S
+{
+ int * ptr;
+ int ptr_len;
+
+ int *&refptr;
+ int refptr_len;
+
+ bool set_ptr (int n)
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (ptr != NULL)
+ for (int i = 0; i < ptr_len; i++)
+ ptr[i] = n;
+ mapped = (ptr != NULL);
+ }
+ return mapped;
+ }
+
+ bool set_refptr (int n)
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (refptr != NULL)
+ for (int i = 0; i < refptr_len; i++)
+ refptr[i] = n;
+ mapped = (refptr != NULL);
+ }
+ return mapped;
+ }
+};
+
+int main (void)
+{
+ #define N 10
+ int *ptr1 = new int[N];
+ int *ptr2 = new int[N];
+
+ memset (ptr1, 0, sizeof (int) * N);
+ memset (ptr2, 0, sizeof (int) * N);
+
+ S s = { ptr1, N, ptr2, N };
+
+ bool mapped;
+ int val = 123;
+
+ mapped = s.set_ptr (val);
+ if (mapped)
+ abort ();
+ if (s.ptr != ptr1)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 0)
+ abort ();
+
+ mapped = s.set_refptr (val);
+ if (mapped)
+ abort ();
+ if (s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 0)
+ abort ();
+
+ #pragma omp target data map(ptr1[:N])
+ mapped = s.set_ptr (val);
+
+ if (!mapped)
+ abort ();
+ if (s.set_refptr (0))
+ abort ();
+ if (s.ptr != ptr1 || s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != val)
+ abort ();
+
+ #pragma omp target data map(ptr2[:N])
+ mapped = s.set_refptr (val);
+
+ if (!mapped)
+ abort ();
+ if (s.set_ptr (0))
+ abort ();
+ if (s.ptr != ptr1 || s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != val)
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
new file mode 100644
index 00000000000..3b2d5811350
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -0,0 +1,107 @@
+// 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>
+
+struct T
+{
+ int *ptr;
+ int ptr_len;
+
+ int *&refptr;
+ int refptr_len;
+
+ auto set_ptr_func (int n)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (ptr)
+ for (int i = 0; i < ptr_len; i++)
+ ptr[i] = n;
+ mapped = (ptr != NULL);
+ }
+ return mapped;
+ };
+ return fn;
+ }
+
+ auto set_refptr_func (int n)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (refptr)
+ for (int i = 0; i < refptr_len; i++)
+ refptr[i] = n;
+ mapped = (refptr != NULL);
+ }
+ return mapped;
+ };
+ return fn;
+ }
+};
+
+int main (void)
+{
+ #define N 10
+ int *ptr1 = new int[N];
+ int *ptr2 = new int[N];
+
+ memset (ptr1, 0, sizeof (int) * N);
+ memset (ptr2, 0, sizeof (int) * N);
+
+ T a = { ptr1, N, ptr2, N };
+
+ auto p1 = a.set_ptr_func (1);
+ auto r2 = a.set_refptr_func (2);
+
+ if (p1 ())
+ abort ();
+ if (r2 ())
+ abort ();
+
+ if (a.ptr != ptr1)
+ abort ();
+ if (a.refptr != ptr2)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 0)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 0)
+ abort ();
+
+ #pragma omp target data map(ptr1[:N], ptr2[:N])
+ {
+ if (!p1 ())
+ abort ();
+ if (!r2 ())
+ abort ();
+ }
+
+ if (a.ptr != ptr1)
+ abort ();
+ if (a.refptr != ptr2)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 1)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr2[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\(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: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) 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:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, 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\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-5.C b/gcc/testsuite/g++.dg/gomp/target-this-5.C
new file mode 100644
index 00000000000..a9ac74bcf1f
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-5.C
@@ -0,0 +1,34 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+extern "C" void abort ();
+
+template<typename T>
+struct S
+{
+ T a, b, c, d;
+
+ T sum (void)
+ {
+ T val = 0;
+ val += a + b + this->c + this->d;
+ return val;
+ }
+
+ T sum_offload (void)
+ {
+ T val = 0;
+ #pragma omp target map(val)
+ val += a + b + this->c + this->d;
+ return val;
+ }
+};
+
+int main (void)
+{
+ S<int> s = { 1, 2, 3, 4 };
+ if (s.sum () != s.sum_offload ())
+ abort ();
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/this-2.C b/gcc/testsuite/g++.dg/gomp/this-2.C
index d03b8a0728e..b521a4faf5e 100644
--- a/gcc/testsuite/g++.dg/gomp/this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/this-2.C
@@ -9,14 +9,14 @@ struct S
void
S::bar (int x)
{
- #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+ #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" }
;
- #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+ #pragma omp target map (this[0], x)
;
- #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
- #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
- #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
- #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+ #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" }
+ #pragma omp target update to (this[0], x)
+ #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" }
+ #pragma omp target update from (this[1], x)
}
template <int N>
@@ -29,14 +29,14 @@ template <int N>
void
T<N>::bar (int x)
{
- #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+ #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" }
;
- #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+ #pragma omp target map (this[0], x)
;
- #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
- #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
- #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
- #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" }
+ #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" }
+ #pragma omp target update to (this[0], x)
+ #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" }
+ #pragma omp target update from (this[1], x)
}
template struct T<0>;
diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c
new file mode 100644
index 00000000000..3e7921270c9
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/target-3.c
@@ -0,0 +1,16 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+struct S
+{
+ int a, b;
+};
+
+void foo (struct S *s)
+{
+ #pragma omp target map (alloc: s->a, s->b)
+ ;
+ #pragma omp target enter data map (alloc: s->a, s->b)
+}
+
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index bcbe669653c..6a855baa26e 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -836,6 +836,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_POINTER:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, "alloc");
break;
case GOMP_MAP_IF_PRESENT:
@@ -914,6 +915,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_ATTACH_DETACH:
pp_string (pp, "attach_detach");
break;
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ pp_string (pp, "attach_zero_length_array_section");
+ break;
default:
gcc_unreachable ();
}
@@ -932,6 +936,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_ALWAYS_POINTER:
pp_string (pp, " [pointer assign, bias: ");
break;
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ pp_string (pp, " [pointer assign, zero-length array section, bias: ");
+ break;
case GOMP_MAP_TO_PSET:
pp_string (pp, " [pointer set, len: ");
break;
@@ -939,6 +946,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_DETACH:
case GOMP_MAP_FORCE_DETACH:
case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, " [bias: ");
break;
default:
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 6e163b02560..4a4a14393d2 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -132,6 +132,11 @@ enum gomp_map_kind
No refcount is bumped by this, and the store is done unconditionally. */
GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_FLAG_SPECIAL | 1),
+ /* Like GOMP_MAP_POINTER, but allow zero-length array section, i.e. set to
+ NULL if target is not mapped. */
+ GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION
+ = (GOMP_MAP_FLAG_SPECIAL_2
+ | GOMP_MAP_FLAG_SPECIAL | 2),
/* Forced deallocation of zero length array section. */
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
= (GOMP_MAP_FLAG_SPECIAL_2
@@ -152,6 +157,12 @@ enum gomp_map_kind
GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY
| GOMP_MAP_FLAG_FORCE | 1),
+ /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections
+ (i.e. set to NULL when array section is not mapped) Currently only used
+ by OpenMP. */
+ GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
+ = (GOMP_MAP_DEEP_COPY | 2),
+
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1),
@@ -175,7 +186,8 @@ enum gomp_map_kind
((X) == GOMP_MAP_ALWAYS_POINTER)
#define GOMP_MAP_POINTER_P(X) \
- ((X) == GOMP_MAP_POINTER)
+ ((X) == GOMP_MAP_POINTER \
+ || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
#define GOMP_MAP_ALWAYS_TO_P(X) \
(((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 8d25dc8e2a8..235e919f191 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1234,7 +1234,7 @@ extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t);
extern void gomp_attach_pointer (struct gomp_device_descr *,
struct goacc_asyncqueue *, splay_tree,
splay_tree_key, uintptr_t, size_t,
- struct gomp_coalesce_buf *);
+ struct gomp_coalesce_buf *, bool);
extern void gomp_detach_pointer (struct gomp_device_descr *,
struct goacc_asyncqueue *, splay_tree_key,
uintptr_t, bool, struct gomp_coalesce_buf *);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index c21508f3739..5ee8647cddf 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -937,7 +937,7 @@ acc_attach_async (void **hostaddr, int async)
}
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
- 0, NULL);
+ 0, NULL, false);
gomp_mutex_unlock (&acc_dev->lock);
}
@@ -1141,7 +1141,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH)
{
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n,
- (uintptr_t) h, s, NULL);
+ (uintptr_t) h, s, NULL, false);
/* OpenACC 'attach'/'detach' doesn't affect structured/dynamic
reference counts ('n->refcount', 'n->dynamic_refcount'). */
}
@@ -1159,7 +1159,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
splay_tree_key m
= lookup_host (acc_dev, hostaddrs[j], sizeof (void *));
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m,
- (uintptr_t) hostaddrs[j], sizes[j], NULL);
+ (uintptr_t) hostaddrs[j], sizes[j], NULL,
+ false);
}
bool processed = false;
diff --git a/libgomp/target.c b/libgomp/target.c
index bb09d501dd6..1ff6a503fe4 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -496,7 +496,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
struct gomp_coalesce_buf *cbuf,
htab_t *refcount_set)
{
- assert (kind != GOMP_MAP_ATTACH);
+ assert (kind != GOMP_MAP_ATTACH
+ || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
tgt_var->key = oldn;
tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
@@ -536,7 +537,8 @@ get_kind (bool short_mapkind, void *kinds, int idx)
static void
gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
uintptr_t host_ptr, uintptr_t target_offset, uintptr_t bias,
- struct gomp_coalesce_buf *cbuf)
+ struct gomp_coalesce_buf *cbuf,
+ bool allow_zero_length_array_sections)
{
struct gomp_device_descr *devicep = tgt->device_descr;
struct splay_tree_s *mem_map = &devicep->mem_map;
@@ -558,16 +560,24 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
if (n == NULL)
{
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("Pointer target of array section wasn't mapped");
- }
- cur_node.host_start -= n->host_start;
- cur_node.tgt_offset
- = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
- /* At this point tgt_offset is target address of the
- array section. Now subtract bias to get what we want
- to initialize the pointer with. */
- cur_node.tgt_offset -= bias;
+ if (allow_zero_length_array_sections)
+ cur_node.tgt_offset = 0;
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("Pointer target of array section wasn't mapped");
+ }
+ }
+ else
+ {
+ cur_node.host_start -= n->host_start;
+ cur_node.tgt_offset
+ = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
+ /* At this point tgt_offset is target address of the
+ array section. Now subtract bias to get what we want
+ to initialize the pointer with. */
+ cur_node.tgt_offset -= bias;
+ }
gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset),
(void *) &cur_node.tgt_offset, sizeof (void *), cbuf);
}
@@ -638,7 +648,8 @@ attribute_hidden void
gomp_attach_pointer (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, splay_tree mem_map,
splay_tree_key n, uintptr_t attach_to, size_t bias,
- struct gomp_coalesce_buf *cbufp)
+ struct gomp_coalesce_buf *cbufp,
+ bool allow_zero_length_array_sections)
{
struct splay_tree_key_s s;
size_t size, idx;
@@ -690,11 +701,21 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
if (!tn)
{
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("pointer target not mapped for attach");
+ if (allow_zero_length_array_sections)
+ {
+ /* When allowing attachment to zero-length array sections, we
+ allow attaching to NULL pointers when the target region is not
+ mapped. */
+ data = 0;
+ }
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("pointer target not mapped for attach");
+ }
}
-
- data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
+ else
+ data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
gomp_debug (1,
"%s: attaching host %p, target %p (struct base %p) to %p\n",
@@ -950,7 +971,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
has_firstprivate = true;
continue;
}
- else if ((kind & typemask) == GOMP_MAP_ATTACH)
+ else if ((kind & typemask) == GOMP_MAP_ATTACH
+ || ((kind & typemask)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
{
tgt->list[i].key = NULL;
has_firstprivate = true;
@@ -1197,7 +1220,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
(uintptr_t) *(void **) hostaddrs[j],
k->tgt_offset + ((uintptr_t) hostaddrs[j]
- k->host_start),
- sizes[j], cbufp);
+ sizes[j], cbufp, false);
}
}
i = j - 1;
@@ -1325,6 +1348,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
++i;
continue;
case GOMP_MAP_ATTACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
{
cur_node.host_start = (uintptr_t) hostaddrs[i];
cur_node.host_end = cur_node.host_start + sizeof (void *);
@@ -1341,9 +1365,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
structured/dynamic reference counts ('n->refcount',
'n->dynamic_refcount'). */
+ bool zlas
+ = ((kind & typemask)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
gomp_attach_pointer (devicep, aq, mem_map, n,
(uintptr_t) hostaddrs[i], sizes[i],
- cbufp);
+ cbufp, zlas);
}
else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
{
@@ -1453,9 +1480,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
k->host_end - k->host_start, cbufp);
break;
case GOMP_MAP_POINTER:
- gomp_map_pointer (tgt, aq,
- (uintptr_t) *(void **) k->host_start,
- k->tgt_offset, sizes[i], cbufp);
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ gomp_map_pointer
+ (tgt, aq, (uintptr_t) *(void **) k->host_start,
+ k->tgt_offset, sizes[i], cbufp,
+ ((kind & typemask)
+ == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION));
break;
case GOMP_MAP_TO_PSET:
gomp_copy_host2dev (devicep, aq,
@@ -1496,7 +1526,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
k->tgt_offset
+ ((uintptr_t) hostaddrs[j]
- k->host_start),
- sizes[j], cbufp);
+ sizes[j], cbufp, false);
}
}
i = j - 1;
diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C
new file mode 100644
index 00000000000..d4f9ff3e983
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-23.C
@@ -0,0 +1,34 @@
+extern "C" void abort ();
+
+struct S
+{
+ int *data;
+};
+
+int
+main (void)
+{
+ #define SZ 10
+ S *s = new S ();
+ s->data = new int[SZ];
+
+ for (int i = 0; i < SZ; i++)
+ s->data[i] = 0;
+
+ #pragma omp target enter data map(to: s)
+ #pragma omp target enter data map(to: s->data[:SZ])
+ #pragma omp target
+ {
+ for (int i = 0; i < SZ; i++)
+ s->data[i] = i;
+ }
+ #pragma omp target exit data map(from: s->data[:SZ])
+ #pragma omp target exit data map(from: s)
+
+ for (int i = 0; i < SZ; i++)
+ if (s->data[i] != i)
+ abort ();
+
+ return 0;
+}
+
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;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-2.C b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
new file mode 100644
index 00000000000..1d3561ffbd7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
@@ -0,0 +1,30 @@
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+ int X, Y;
+ #pragma omp target map(from: X, Y)
+ {
+ int x = 0, y = 0;
+
+ for (int i = 0; i < N; i++)
+ [&] (int v) { x += v; } (i);
+
+ auto yinc = [&y] { y++; };
+ for (int i = 0; i < N; i++)
+ yinc ();
+
+ X = x;
+ Y = y;
+ }
+
+ int Xs = 0;
+ for (int i = 0; i < N; i++)
+ Xs += i;
+ if (X != Xs)
+ abort ();
+
+ if (Y != N)
+ abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-1.C b/libgomp/testsuite/libgomp.c++/target-this-1.C
new file mode 100644
index 00000000000..a591ea4c564
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-1.C
@@ -0,0 +1,29 @@
+extern "C" void abort ();
+
+struct S
+{
+ int a, b, c, d;
+
+ int sum (void)
+ {
+ int val = 0;
+ val += a + b + this->c + this->d;
+ return val;
+ }
+
+ int sum_offload (void)
+ {
+ int val = 0;
+ #pragma omp target map(val)
+ val += a + b + this->c + this->d;
+ return val;
+ }
+};
+
+int main (void)
+{
+ S s = { 1, 2, 3, 4 };
+ if (s.sum () != s.sum_offload ())
+ abort ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-2.C b/libgomp/testsuite/libgomp.c++/target-this-2.C
new file mode 100644
index 00000000000..8119be8c2c5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-2.C
@@ -0,0 +1,47 @@
+
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14" }
+
+extern "C" void abort ();
+
+struct T
+{
+ int x, y;
+
+ auto sum_func (int n)
+ {
+ auto fn = [=](int m) -> int
+ {
+ int v;
+ v = (x + y) * n + m;
+ return v;
+ };
+ return fn;
+ }
+
+ auto sum_func_offload (int n)
+ {
+ auto fn = [=](int m) -> int
+ {
+ int v;
+ #pragma omp target map(from:v)
+ v = (x + y) * n + m;
+ return v;
+ };
+ return fn;
+ }
+
+};
+
+int main (void)
+{
+ T a = { 1, 2 };
+
+ auto s1 = a.sum_func (3);
+ auto s2 = a.sum_func_offload (3);
+
+ if (s1 (1) != s2 (1))
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-3.C b/libgomp/testsuite/libgomp.c++/target-this-3.C
new file mode 100644
index 00000000000..e15f69a1623
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-3.C
@@ -0,0 +1,99 @@
+#include <stdio.h>
+#include <string.h>
+extern "C" void abort ();
+
+struct S
+{
+ int * ptr;
+ int ptr_len;
+
+ int *&refptr;
+ int refptr_len;
+
+ bool set_ptr (int n)
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (ptr != NULL)
+ for (int i = 0; i < ptr_len; i++)
+ ptr[i] = n;
+ mapped = (ptr != NULL);
+ }
+ return mapped;
+ }
+
+ bool set_refptr (int n)
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (refptr != NULL)
+ for (int i = 0; i < refptr_len; i++)
+ refptr[i] = n;
+ mapped = (refptr != NULL);
+ }
+ return mapped;
+ }
+};
+
+int main (void)
+{
+ #define N 10
+ int *ptr1 = new int[N];
+ int *ptr2 = new int[N];
+
+ memset (ptr1, 0, sizeof (int) * N);
+ memset (ptr2, 0, sizeof (int) * N);
+
+ S s = { ptr1, N, ptr2, N };
+
+ bool mapped;
+ int val = 123;
+
+ mapped = s.set_ptr (val);
+ if (mapped)
+ abort ();
+ if (s.ptr != ptr1)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 0)
+ abort ();
+
+ mapped = s.set_refptr (val);
+ if (mapped)
+ abort ();
+ if (s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 0)
+ abort ();
+
+ #pragma omp target data map(ptr1[:N])
+ mapped = s.set_ptr (val);
+
+ if (!mapped)
+ abort ();
+ if (s.set_refptr (0))
+ abort ();
+ if (s.ptr != ptr1 || s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != val)
+ abort ();
+
+ #pragma omp target data map(ptr2[:N])
+ mapped = s.set_refptr (val);
+
+ if (!mapped)
+ abort ();
+ if (s.set_ptr (0))
+ abort ();
+ if (s.ptr != ptr1 || s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != val)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-4.C b/libgomp/testsuite/libgomp.c++/target-this-4.C
new file mode 100644
index 00000000000..9f53677a240
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-4.C
@@ -0,0 +1,104 @@
+
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14" }
+#include <cstdlib>
+#include <cstring>
+
+struct T
+{
+ int *ptr;
+ int ptr_len;
+
+ int *&refptr;
+ int refptr_len;
+
+ auto set_ptr_func (int n)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (ptr)
+ for (int i = 0; i < ptr_len; i++)
+ ptr[i] = n;
+ mapped = (ptr != NULL);
+ }
+ return mapped;
+ };
+ return fn;
+ }
+
+ auto set_refptr_func (int n)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (refptr)
+ for (int i = 0; i < refptr_len; i++)
+ refptr[i] = n;
+ mapped = (refptr != NULL);
+ }
+ return mapped;
+ };
+ return fn;
+ }
+};
+
+int main (void)
+{
+ #define N 10
+ int *ptr1 = new int[N];
+ int *ptr2 = new int[N];
+
+ memset (ptr1, 0, sizeof (int) * N);
+ memset (ptr2, 0, sizeof (int) * N);
+
+ T a = { ptr1, N, ptr2, N };
+
+ auto p1 = a.set_ptr_func (1);
+ auto r2 = a.set_refptr_func (2);
+
+ if (p1 ())
+ abort ();
+ if (r2 ())
+ abort ();
+
+ if (a.ptr != ptr1)
+ abort ();
+ if (a.refptr != ptr2)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 0)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 0)
+ abort ();
+
+ #pragma omp target data map(ptr1[:N], ptr2[:N])
+ {
+ if (!p1 ())
+ abort ();
+ if (!r2 ())
+ abort ();
+ }
+
+ if (a.ptr != ptr1)
+ abort ();
+ if (a.refptr != ptr2)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 1)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 2)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-5.C b/libgomp/testsuite/libgomp.c++/target-this-5.C
new file mode 100644
index 00000000000..e71c566687d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-5.C
@@ -0,0 +1,30 @@
+extern "C" void abort ();
+
+template<typename T>
+struct S
+{
+ T a, b, c, d;
+
+ T sum (void)
+ {
+ T val = 0;
+ val += a + b + this->c + this->d;
+ return val;
+ }
+
+ T sum_offload (void)
+ {
+ T val = 0;
+ #pragma omp target map(val)
+ val += a + b + this->c + this->d;
+ return val;
+ }
+};
+
+int main (void)
+{
+ S<int> s = { 1, 2, 3, 4 };
+ if (s.sum () != s.sum_offload ())
+ abort ();
+ return 0;
+}
[-- Attachment #3: v3-to-v4.patch --]
[-- Type: text/plain, Size: 3741 bytes --]
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 41ddea405ea..455620c0663 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -9049,6 +9049,8 @@ struct omp_target_walk_data
tree current_closure;
hash_set<tree> closure_vars_accessed;
+
+ hash_set<tree> local_decls;
};
static tree
@@ -9107,12 +9109,25 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
return NULL_TREE;
}
+ if (TREE_CODE (t) == BIND_EXPR)
+ {
+ tree block = BIND_EXPR_BLOCK (t);
+ for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var))
+ if (!data->local_decls.contains (var))
+ data->local_decls.add (var);
+ 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))
+ if (!data->lambda_objects_accessed.contains (t)
+ /* Do not prepare to create target maps for locally declared
+ lambdas or anonymous ones. */
+ && !data->local_decls.contains (t)
+ && TREE_CODE (t) != TARGET_EXPR)
data->lambda_objects_accessed.add (t);
*walk_subtrees = 0;
return NULL_TREE;
@@ -9398,6 +9413,9 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
i != data.lambda_objects_accessed.end (); ++i)
{
tree lobj = *i;
+ if (TREE_CODE (lobj) == TARGET_EXPR)
+ lobj = TREE_OPERAND (lobj, 0);
+
tree lt = TREE_TYPE (lobj);
gcc_assert (LAMBDA_TYPE_P (lt) && CLASS_TYPE_P (lt));
@@ -9434,7 +9452,7 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
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_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
OMP_CLAUSE_DECL (c)
= build1 (INDIRECT_REF, TREE_TYPE (TREE_TYPE (exp)), exp);
OMP_CLAUSE_SIZE (c)
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-2.C b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
new file mode 100644
index 00000000000..bdf2564cd04
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
@@ -0,0 +1,35 @@
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+ int X, Y;
+ #pragma omp target map(from: X, Y)
+ {
+ int x = 0, y = 0;
+
+ for (int i = 0; i < N; i++)
+ [&] (int v) { x += v; } (i);
+
+ auto yinc = [&y] { y++; };
+ for (int i = 0; i < N; i++)
+ yinc ();
+
+ X = x;
+ Y = y;
+ }
+
+ int Xs = 0;
+ for (int i = 0; i < N; i++)
+ Xs += i;
+ if (X != Xs)
+ abort ();
+
+ if (Y != N)
+ abort ();
+}
+
+/* Make sure lambda objects do NOT appear in target maps. */
+/* { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(from:Y \[len: [0-9]+\]\) map\(from:X \[len: [0-9]+\]\)$} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-2.C b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
new file mode 100644
index 00000000000..1d3561ffbd7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
@@ -0,0 +1,30 @@
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+ int X, Y;
+ #pragma omp target map(from: X, Y)
+ {
+ int x = 0, y = 0;
+
+ for (int i = 0; i < N; i++)
+ [&] (int v) { x += v; } (i);
+
+ auto yinc = [&y] { y++; };
+ for (int i = 0; i < N; i++)
+ yinc ();
+
+ X = x;
+ Y = y;
+ }
+
+ int Xs = 0;
+ for (int i = 0; i < N; i++)
+ Xs += i;
+ if (X != Xs)
+ abort ();
+
+ if (Y != N)
+ abort ();
+}
^ permalink raw reply [flat|nested] 6+ messages in thread* Re: [PATCH, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v4] 2021-06-18 14:25 [PATCH, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v4] Chung-Lin Tang @ 2021-06-24 13:15 ` Jakub Jelinek 2021-11-16 12:43 ` [PATCH, v5, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v5] Chung-Lin Tang 0 siblings, 1 reply; 6+ messages in thread From: Jakub Jelinek @ 2021-06-24 13:15 UTC (permalink / raw) To: Chung-Lin Tang; +Cc: gcc-patches, Catherine Moore On Fri, Jun 18, 2021 at 10:25:16PM +0800, Chung-Lin Tang wrote: Note, you'll need to rebase your patch, it clashes with r12-1768-g7619d33471c10fe3d149dcbb701d99ed3dd23528. Sorry for that. And sorry for patch review delay. > --- a/gcc/c/c-typeck.c > +++ b/gcc/c/c-typeck.c > @@ -13104,6 +13104,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, > return error_mark_node; > } > t = TREE_OPERAND (t, 0); > + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) Map clauses never appear on declare simd, so (ort == C_ORT_ACC || ort == C_ORT_OMP) previously meant always and since the in_reduction change is incorrect (as C_ORT_OMP_TARGET is used for target construct but not for e.g. target data* or target update). > + && TREE_CODE (t) == MEM_REF) So please just use if (TREE_CODE (t) == MEM_REF) or explain when it shouldn't trigger. > @@ -14736,6 +14743,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) > { > while (TREE_CODE (t) == COMPONENT_REF) > t = TREE_OPERAND (t, 0); > + if (TREE_CODE (t) == MEM_REF) > + { > + t = TREE_OPERAND (t, 0); > + STRIP_NOPS (t); > + } This doesn't look correct. At least the parsing (and the spec AFAIK) doesn't ensure that if there is ->, it must come before all the dots. So, if one uses map (s->x.y) the above would work, but if map (s->x.y->z) or map (s.a->b->c->d->e) is used, it wouldn't. I'd expect a single while loop that looks through COMPONENT_REFs and MEM_REFs as they appear. Maybe the handle_omp_array_sections_1 MEM_REF case too? Or do you want to have it done incrementally, start with supporting only a single -> first before all the dots and later on add support for the rest? I think the 5.0 and especially 5.1 wording basically says that map clause operand is arbitrary lvalue expression that includes array section support too, so eventually we should just have somewhere in parsing scope a bool whether OpenMP array sections are allowed or not, add OMP_ARRAY_REF or similar tree code for those and after parsing the expression, ensure array sections appear only where they can appear and for a subset of the lvalue expressions where we have decl plus series of -> field or . field or [ index ] or [ array section stuff ] handle those specially. That arbitrary lvalue can certainly be done incrementally. map (foo(123)->a.b[3]->c.d[:7]) and the like. > if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP > && OMP_CLAUSE_MAP_IMPLICIT (c) > && (bitmap_bit_p (&map_head, DECL_UID (t)) > @@ -14802,6 +14814,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) > bias) to zero here, so it is not set erroneously to the pointer > size later on in gimplify.c. */ > OMP_CLAUSE_SIZE (c) = size_zero_node; > + indir_component_ref_p = false; > + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) Same comment about ort tests. > + && TREE_CODE (t) == COMPONENT_REF > + && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF) > + { > + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); > + indir_component_ref_p = true; > + STRIP_NOPS (t); > + } Again, this can handle only a single -> > @@ -42330,16 +42328,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, > cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc; > } > } > - tree stmt = make_node (OMP_TARGET); > - TREE_TYPE (stmt) = void_type_node; > - OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; > - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); > - OMP_TARGET_BODY (stmt) = body; > - OMP_TARGET_COMBINED (stmt) = 1; > - SET_EXPR_LOCATION (stmt, pragma_tok->location); > - add_stmt (stmt); > - pc = &OMP_TARGET_CLAUSES (stmt); > - goto check_clauses; > + c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true); > + finish_omp_target (pragma_tok->location, > + cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true); What is the advantage of finish_omp_target. Perhaps the check_clauses label can be renamed and more things common to both paths moved after the label if needed, but as long as it isn't something also called during instantiation, I find it cleaner to do it in cp_parser_omp_target at one place. The reason for e.g. finish_omp_parallel is that it is called from both parsing and instantiation. > --- a/gcc/cp/semantics.c > +++ b/gcc/cp/semantics.c > @@ -4990,6 +4990,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, > { > if (error_operand_p (t)) > return error_mark_node; > + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) See above about ort. Declare simd only allows uniform, linear, aligned, simdlen, inbranch and notinbranch clauses and none of those support array sections. > + && TREE_CODE (t) == FIELD_DECL) > + t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); handle_omp_array_sections_1 already has recent: if (TREE_CODE (t) == FIELD_DECL && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND)) ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); so shouldn't that be extended to map/to/from clauses too? And I guess we should check reduction/in_reduction/task_reduction clauses too. > @@ -9003,6 +9037,493 @@ 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; > + > + hash_set<tree> local_decls; > +}; > + > +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; This is something that we'll eventually need to do e.g. for declare mapper in all the 3 FEs, gather what variables might need to be mapped and for all their types look up the mappers (recursively for nested types and for all types mentioned in those declare mappers etc.) and remember that somehow until gimplification. If it is only preliminary and covers might appear rather than appears, I think it should be fine. What this routine does is ultimate, if you see this somewhere, you say it is accessed, if you see a lambda, again, it has to be accessed etc. I'm afraid that is unsafe though. The IL at this point isn't folded yet, one could have sizeof (this) or other unevaluated context appear there, or something could appear in a private clause on some inner construct that doesn't imply an access on the outer target, etc. So, I think either this function would need to be more careful, especially for nested OpenMP constructs, or can't it be done through langhooks at gimplification time when we should know exactly what appears and what doesn't in the body? > + if (TREE_TYPE(t) && LAMBDA_TYPE_P (TREE_TYPE (t))) Formatting, missing space before (. > + 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); Is it guaranteed everything added here can't have an explicit map clause already? Jakub ^ permalink raw reply [flat|nested] 6+ messages in thread
* [PATCH, v5, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v5] 2021-06-24 13:15 ` Jakub Jelinek @ 2021-11-16 12:43 ` Chung-Lin Tang 2021-12-03 16:47 ` Jakub Jelinek 2022-09-07 7:04 ` [committed] openmp: Fix handling of target constructs in static member functions [PR106829] Jakub Jelinek 0 siblings, 2 replies; 6+ messages in thread From: Chung-Lin Tang @ 2021-11-16 12:43 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc-patches, Catherine Moore [-- Attachment #1: Type: text/plain, Size: 15924 bytes --] Hi Jakub, On 2021/6/24 9:15 PM, Jakub Jelinek wrote: > On Fri, Jun 18, 2021 at 10:25:16PM +0800, Chung-Lin Tang wrote: > > Note, you'll need to rebase your patch, it clashes with > r12-1768-g7619d33471c10fe3d149dcbb701d99ed3dd23528. > Sorry for that. And sorry for patch review delay. > >> --- a/gcc/c/c-typeck.c >> +++ b/gcc/c/c-typeck.c >> @@ -13104,6 +13104,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, >> return error_mark_node; >> } >> t = TREE_OPERAND (t, 0); >> + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) > > Map clauses never appear on declare simd, so > (ort == C_ORT_ACC || ort == C_ORT_OMP) > previously meant always and since the in_reduction change is incorrect > (as C_ORT_OMP_TARGET is used for target construct but not for > e.g. target data* or target update). > >> + && TREE_CODE (t) == MEM_REF) Upon reviewing, it appears that most of these C_ORT_* tests are no longer needed, removed in new patch. > So please just use if (TREE_CODE (t) == MEM_REF) > or explain when it shouldn't trigger. > >> @@ -14736,6 +14743,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) >> { >> while (TREE_CODE (t) == COMPONENT_REF) >> t = TREE_OPERAND (t, 0); >> + if (TREE_CODE (t) == MEM_REF) >> + { >> + t = TREE_OPERAND (t, 0); >> + STRIP_NOPS (t); >> + } > > This doesn't look correct. At least the parsing (and the spec AFAIK) > doesn't ensure that if there is ->, it must come before all the dots. > So, if one uses map (s->x.y) the above would work, but if map (s->x.y->z) or > map (s.a->b->c->d->e) is used, it wouldn't. I'd expect a single > while loop that looks through COMPONENT_REFs and MEM_REFs as they appear. > Maybe the handle_omp_array_sections_1 MEM_REF case too? > > Or do you want to have it done incrementally, start with supporting only > a single -> first before all the dots and later on add support for the rest? > > I think the 5.0 and especially 5.1 wording basically says that map clause > operand is arbitrary lvalue expression that includes array section support > too, so eventually we should just have somewhere in parsing scope a bool > whether OpenMP array sections are allowed or not, add OMP_ARRAY_REF or > similar tree code for those and after parsing the expression, ensure > array sections appear only where they can appear and for a subset of the > lvalue expressions where we have decl plus series of -> field or . field > or [ index ] or [ array section stuff ] handle those specially. > That arbitrary lvalue can certainly be done incrementally. > map (foo(123)->a.b[3]->c.d[:7]) and the like. Indeed this kind of modification is sort of "as encountered", so there are probably many cases that are not completely handled yet; it's not just the front-end, but also changes in gimplify_scan_omp_clauses(). However, I had another patch that should've plowed a bit further on this: https://gcc.gnu.org/pipermail/gcc-patches/2021-May/570075.html as well as those patch sets that Julian is working on. (our current plan is to have my sets go in first, and Julian's on top, to minimize clashing) >> if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP >> && OMP_CLAUSE_MAP_IMPLICIT (c) >> && (bitmap_bit_p (&map_head, DECL_UID (t)) >> @@ -14802,6 +14814,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) >> bias) to zero here, so it is not set erroneously to the pointer >> size later on in gimplify.c. */ >> OMP_CLAUSE_SIZE (c) = size_zero_node; >> + indir_component_ref_p = false; >> + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) > > Same comment about ort tests. > >> + && TREE_CODE (t) == COMPONENT_REF >> + && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF) >> + { >> + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); >> + indir_component_ref_p = true; >> + STRIP_NOPS (t); >> + } > > Again, this can handle only a single -> > >> @@ -42330,16 +42328,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, >> cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc; >> } >> } >> - tree stmt = make_node (OMP_TARGET); >> - TREE_TYPE (stmt) = void_type_node; >> - OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; >> - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); >> - OMP_TARGET_BODY (stmt) = body; >> - OMP_TARGET_COMBINED (stmt) = 1; >> - SET_EXPR_LOCATION (stmt, pragma_tok->location); >> - add_stmt (stmt); >> - pc = &OMP_TARGET_CLAUSES (stmt); >> - goto check_clauses; >> + c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true); >> + finish_omp_target (pragma_tok->location, >> + cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true); > > What is the advantage of finish_omp_target. Perhaps the check_clauses label > can be renamed and more things common to both paths moved after the label if > needed, but as long as it isn't something also called during instantiation, > I find it cleaner to do it in cp_parser_omp_target at one place. > The reason for e.g. finish_omp_parallel is that it is called from both > parsing and instantiation. Originally, finish_omp_target was also meant for calling from both parsing and instantiation, but later factoring turned that part into finish_omp_target_clauses, while finish_omp_target was still kept. But having this part factored as a finish_omp_target function seems quite consistent with usual C/C++ front-end conventions. >> --- a/gcc/cp/semantics.c >> +++ b/gcc/cp/semantics.c >> @@ -4990,6 +4990,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, >> { >> if (error_operand_p (t)) >> return error_mark_node; >> + if ((ort == C_ORT_ACC || ort == C_ORT_OMP) > > See above about ort. > Declare simd only allows uniform, linear, aligned, simdlen, inbranch and > notinbranch clauses and none of those support array sections. Removed such conditions. >> + && TREE_CODE (t) == FIELD_DECL) >> + t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); > > handle_omp_array_sections_1 already has recent: > if (TREE_CODE (t) == FIELD_DECL > && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY > || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND)) > ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); > so shouldn't that be extended to map/to/from clauses too? > And I guess we should check reduction/in_reduction/task_reduction clauses > too. The use of finish_non_static_data_member() here appears to not need to be divided between clause codes now, so the OMP_CLAUSE_AFFINITY/DEPEND test has been removed here. Also, the cases in handle_omp_array_sections_1 and finish_omp_clauses that are related to "'this' is allowed in OpenMP only in declare simd" are removed, since these restrictions are now lifted. >> @@ -9003,6 +9037,493 @@ 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; >> + >> + hash_set<tree> local_decls; >> +}; >> + >> +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; > > This is something that we'll eventually need to do e.g. for declare mapper > in all the 3 FEs, gather what variables might need to be mapped and > for all their types look up the mappers (recursively for nested types and > for all types mentioned in those declare mappers etc.) and remember that > somehow until gimplification. Probably, which will probably need more foundational plumbing to establish ways to pass that kind of information to the middle-end. > If it is only preliminary and covers might appear rather than appears, > I think it should be fine. What this routine does is ultimate, if you see > this somewhere, you say it is accessed, if you see a lambda, again, it has > to be accessed etc. I'm afraid that is unsafe though. > The IL at this point isn't folded yet, one could have sizeof (this) or other > unevaluated context appear there, or something could appear in a private clause > on some inner construct that doesn't imply an access on the outer target, etc. We'll just have to handle those cases as they come up. Right now, 'this' related constructs do not seem to be allowed on private() clause at all. > So, I think either this function would need to be more careful, especially > for nested OpenMP constructs, or can't it be done through langhooks at > gimplification time when we should know exactly what appears and what > doesn't in the body? I have added a case to ignore SIZEOF_EXPR and ALIGNOF_EXPR expressions, not sure if there are more, but usually if we map something unnecessary, it should only be inefficient but not wrong code. Further cases can be handled as we encounter them. >> + if (TREE_TYPE(t) && LAMBDA_TYPE_P (TREE_TYPE (t))) > > Formatting, missing space before (. Fixed. >> + 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); > > Is it guaranteed everything added here can't have an explicit map clause > already? I have re-checked the cases, and reorganized the code a bit. Right now all cases in this C++ front-end patch should be ensured to not create new map clauses when a user-explicit one exists. Re-tested without regressions on trunk on x86_64-linux with nvptx offloading, is this okay for trunk? Thanks, Chung-Lin 2021-11-16 Chung-Lin Tang <cltang@codesourcery.com> PR middle-end/92120 gcc/cp/ChangeLog: * cp-tree.h (finish_omp_target): New declaration. (finish_omp_target_clauses): Likewise. * parser.c (cp_parser_omp_clause_map): Adjust call to cp_parser_omp_var_list_no_open to set 'allow_deref' argument to true. (cp_parser_omp_target): Factor out code, adjust into calls to new function finish_omp_target. * pt.c (tsubst_expr): Add call to finish_omp_target_clauses for OMP_TARGET case. * semantics.c (handle_omp_array_sections_1): Add handling to create 'this->member' from 'member' FIELD_DECL. Remove case of rejecting 'this' when not in declare simd. (handle_omp_array_sections): Likewise. (finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP map clauses. Handle 'A->member' case in map clauses. Remove case of rejecting 'this' when not in declare simd. (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. (finish_omp_target): New function. gcc/c/ChangeLog: * c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in call to c_parser_omp_variable_list to 'true'. * c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in array base handling. (c_finish_omp_clauses): Handle 'A->member' case in map clauses. gcc/ChangeLog: * gimplify.c ("tree-hash-traits.h"): Add include. (gimplify_scan_omp_clauses): Change struct_map_to_clause to type hash_map<tree_operand, tree> *. Adjust struct map handling to handle cases of *A and A->B expressions. Under !DECL_P case of GOMP_CLAUSE_MAP handling, add STRIP_NOPS for indir_p case, add to struct_deref_set for map(*ptr_to_struct) cases. Add MEM_REF case when handling component_ref_p case. Add unshare_expr and gimplification when created GOMP_MAP_STRUCT is not a DECL. Add code to add firstprivate pointer for *pointer-to-struct case. (gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for exit data directives code to earlier position. * omp-low.c (lower_omp_target): Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. * tree-pretty-print.c (dump_omp_clause): Likewise. gcc/testsuite/ChangeLog: * gcc.dg/gomp/target-3.c: New testcase. * g++.dg/gomp/target-3.C: New testcase. * g++.dg/gomp/target-lambda-1.C: New testcase. * g++.dg/gomp/target-lambda-2.C: New testcase. * g++.dg/gomp/target-this-1.C: New testcase. * g++.dg/gomp/target-this-2.C: New testcase. * g++.dg/gomp/target-this-3.C: New testcase. * g++.dg/gomp/target-this-4.C: New testcase. * g++.dg/gomp/target-this-5.C: New testcase. * g++.dg/gomp/this-2.C: Adjust testcase. include/ChangeLog: * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. (GOMP_MAP_POINTER_P): Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION. libgomp/ChangeLog: * libgomp.h (gomp_attach_pointer): Add bool parameter. * oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_existing): Update assert condition to include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION. (gomp_map_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for mapping a pointer with NULL target. (gomp_attach_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for attaching a pointer with NULL target. (gomp_map_vars_internal): Update calls to gomp_map_pointer and gomp_attach_pointer, add handling for GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases. * testsuite/libgomp.c++/target-23.C: New testcase. * testsuite/libgomp.c++/target-lambda-1.C: New testcase. * testsuite/libgomp.c++/target-lambda-2.C: New testcase. * testsuite/libgomp.c++/target-this-1.C: New testcase. * testsuite/libgomp.c++/target-this-2.C: New testcase. * testsuite/libgomp.c++/target-this-3.C: New testcase. * testsuite/libgomp.c++/target-this-4.C: New testcase. * testsuite/libgomp.c++/target-this-5.C: New testcase. [-- Attachment #2: pr92120-v5.patch --] [-- Type: text/plain, Size: 78369 bytes --] diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 3c9f5877481..acbf20dcb58 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -16175,7 +16175,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list) c_parser_consume_token (parser); } - nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list); + nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list, + true); for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 782414f8c8c..c0ebb319aff 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13238,6 +13238,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, return error_mark_node; } t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == MEM_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF) { if (maybe_ne (mem_ref_offset (t), 0)) @@ -14083,6 +14088,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) tree ordered_clause = NULL_TREE; tree schedule_clause = NULL_TREE; bool oacc_async = false; + bool indir_component_ref_p = false; tree last_iterators = NULL_TREE; bool last_iterators_remove = false; tree *nogroup_seen = NULL; @@ -14884,6 +14890,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { while (TREE_CODE (t) == COMPONENT_REF) t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == MEM_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c) && (bitmap_bit_p (&map_head, DECL_UID (t)) @@ -14950,6 +14961,14 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bias) to zero here, so it is not set erroneously to the pointer size later on in gimplify.c. */ OMP_CLAUSE_SIZE (c) = size_zero_node; + indir_component_ref_p = false; + if (TREE_CODE (t) == COMPONENT_REF + && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF) + { + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + indir_component_ref_p = true; + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { @@ -15022,6 +15041,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)) + && !indir_component_ref_p && !c_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -15078,8 +15098,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && (ort == C_ORT_ACC - || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index f387b5036d2..b180fad25da 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -7653,6 +7653,8 @@ extern tree start_lambda_function (tree fn, tree lambda_expr); 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 *); /* in tree.c */ extern int cp_tree_operand_length (const_tree); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index adfd3c1378d..8c27ea12013 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -39232,7 +39232,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) } nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, - NULL); + NULL, true); for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); @@ -44021,8 +44021,6 @@ static bool cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, enum pragma_context context, bool *if_p) { - tree *pc = NULL, stmt; - if (flag_openmp) omp_requires_mask = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED); @@ -44127,16 +44125,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc; } - tree stmt = make_node (OMP_TARGET); - TREE_TYPE (stmt) = void_type_node; - OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); - OMP_TARGET_BODY (stmt) = body; - OMP_TARGET_COMBINED (stmt) = 1; - SET_EXPR_LOCATION (stmt, pragma_tok->location); - add_stmt (stmt); - pc = &OMP_TARGET_CLAUSES (stmt); - goto check_clauses; + c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true); + finish_omp_target (pragma_tok->location, + cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true); + return true; } else if (!flag_openmp) /* flag_openmp_simd */ { @@ -44171,13 +44163,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, return false; } - stmt = make_node (OMP_TARGET); - TREE_TYPE (stmt) = void_type_node; - - OMP_TARGET_CLAUSES (stmt) - = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, - "#pragma omp target", pragma_tok, false); - for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, + "#pragma omp target", pragma_tok, + false); + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION) { tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -44186,45 +44175,13 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = nc; } - OMP_TARGET_CLAUSES (stmt) - = finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET); - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); + clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET); - pc = &OMP_TARGET_CLAUSES (stmt); + c_omp_adjust_map_clauses (clauses, true); keep_next_level (true); - OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); - - SET_EXPR_LOCATION (stmt, pragma_tok->location); - add_stmt (stmt); + tree body = cp_parser_omp_structured_block (parser, if_p); -check_clauses: - while (*pc) - { - if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_TO: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_FROM: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_TOFROM: - case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_ALLOC: - case GOMP_MAP_FIRSTPRIVATE_POINTER: - case GOMP_MAP_FIRSTPRIVATE_REFERENCE: - case GOMP_MAP_ALWAYS_POINTER: - case GOMP_MAP_ATTACH_DETACH: - break; - default: - error_at (OMP_CLAUSE_LOCATION (*pc), - "%<#pragma omp target%> with map-type other " - "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> " - "on %<map%> clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } - pc = &OMP_CLAUSE_CHAIN (*pc); - } + finish_omp_target (pragma_tok->location, clauses, body, false); return true; } diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 82bf7dc26f6..c57666691fb 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -18950,6 +18950,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 60e0982cc48..c64b45c0cee 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5054,15 +5054,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, return error_mark_node; } t = TREE_OPERAND (t, 0); - if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF) - t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } } if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); } - if (TREE_CODE (t) == FIELD_DECL - && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND)) + if (TREE_CODE (t) == FIELD_DECL) ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { @@ -5078,18 +5079,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } - else if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP - && TREE_CODE (t) == PARM_DECL - && DECL_ARTIFICIAL (t) - && DECL_NAME (t) == this_identifier - && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY - && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND) - { - error_at (OMP_CLAUSE_LOCATION (c), - "%<this%> allowed in OpenMP only in %<declare simd%>" - " clauses"); - return error_mark_node; - } else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND && VAR_P (t) && CP_DECL_THREAD_LOCAL_P (t)) @@ -5603,6 +5592,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } OMP_CLAUSE_DECL (c) = first; OMP_CLAUSE_SIZE (c) = size; + if (TREE_CODE (t) == FIELD_DECL) + t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) @@ -6616,6 +6607,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 indir_component_ref_p = false; tree last_iterators = NULL_TREE; bool last_iterators_remove = false; /* 1 if normal/task reduction has been seen, -1 if inscan reduction @@ -7867,6 +7859,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c) && (bitmap_bit_p (&map_head, DECL_UID (t)) @@ -7939,9 +7936,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } + indir_component_ref_p = false; if (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) - t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + { + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + indir_component_ref_p = true; + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { @@ -7988,6 +7990,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) goto handle_map_references; } } + if (!processing_template_decl + && TREE_CODE (t) == FIELD_DECL) + { + OMP_CLAUSE_DECL (c) = finish_non_static_data_member (t, NULL_TREE, + NULL_TREE); + break; + } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { if (processing_template_decl && TREE_CODE (t) != OVERLOAD) @@ -8014,19 +8023,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (ort != C_ORT_ACC && t == current_class_ptr) - { - error_at (OMP_CLAUSE_LOCATION (c), - "%<this%> allowed in OpenMP only in %<declare simd%>" - " clauses"); - remove = true; - break; - } else if (!processing_template_decl && !TYPE_REF_P (TREE_TYPE (t)) && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)) + && !indir_component_ref_p && !cxx_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -9182,6 +9184,511 @@ 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; + + hash_set<tree> local_decls; +}; + +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; + + /* References inside of these expression codes shouldn't incur any + form of mapping, so return early. */ + if (TREE_CODE (t) == SIZEOF_EXPR + || TREE_CODE (t) == ALIGNOF_EXPR) + { + *walk_subtrees = 0; + return NULL_TREE; + } + + 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_CODE (t) == BIND_EXPR) + { + tree block = BIND_EXPR_BLOCK (t); + for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var)) + if (!data->local_decls.contains (var)) + data->local_decls.add (var); + 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) + /* Do not prepare to create target maps for locally declared + lambdas or anonymous ones. */ + && !data->local_decls.contains (t) + && TREE_CODE (t) != TARGET_EXPR) + data->lambda_objects_accessed.add (t); + *walk_subtrees = 0; + return NULL_TREE; + } + + return NULL_TREE; +} + +void +finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr) +{ + omp_target_walk_data data; + data.this_expr_accessed = false; + + 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; + + tree omp_target_this_expr = NULL_TREE; + tree *explicit_this_deref_map = NULL; + if (data.this_expr_accessed) + { + 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. */ + for (tree *cp = clauses_ptr; *cp; cp = &OMP_CLAUSE_CHAIN (*cp)) + if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP + && (TREE_CODE (OMP_CLAUSE_DECL (*cp)) == INDIRECT_REF + || TREE_CODE (OMP_CLAUSE_DECL (*cp)) == MEM_REF) + && operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*cp), 0), + omp_target_this_expr)) + { + explicit_this_deref_map = cp; + break; + } + } + + if (DECL_LAMBDA_FUNCTION_P (current_function_decl) + && (data.this_expr_accessed + || !data.closure_vars_accessed.is_empty ())) + { + /* For lambda functions, we need to first create a copy of the + __closure object. */ + tree closure = DECL_ARGUMENTS (current_function_decl); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO); + 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)); + + 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); + } + + if (data.this_expr_accessed) + { + /* If the this-expr was accessed, create a map(*this) clause. */ + enum gomp_map_kind kind = GOMP_MAP_TOFROM; + if (explicit_this_deref_map) + { + tree this_map = *explicit_this_deref_map; + tree nc = OMP_CLAUSE_CHAIN (this_map); + gcc_assert (nc != NULL_TREE + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) + == GOMP_MAP_FIRSTPRIVATE_POINTER)); + kind = OMP_CLAUSE_MAP_KIND (this_map); + /* Remove the original 'map(*this) map(firstprivate_ptr:this)' + two-map sequence away from the chain. */ + *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc); + } + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, kind); + 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))); + new_clauses.safe_push (c); + + /* If we're in a lambda function, the this-pointer will actually be + '__closure->this', a mapped member of __closure, hence always_pointer. + Otherwise it's a firstprivate pointer. */ + enum gomp_map_kind ptr_kind + = (DECL_LAMBDA_FUNCTION_P (current_function_decl) + ? GOMP_MAP_ALWAYS_POINTER + : GOMP_MAP_FIRSTPRIVATE_POINTER); + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, ptr_kind); + OMP_CLAUSE_DECL (c) = omp_target_this_expr; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + + if (DECL_LAMBDA_FUNCTION_P (current_function_decl)) + { + if (omp_target_this_expr) + { + 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 + || TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE) + { + /* this-pointer is processed above, outside this loop. */ + if (omp_target_this_expr + && operand_equal_p (closure_expr, omp_target_this_expr)) + continue; + + bool ptr_p = TREE_CODE (TREE_TYPE (orig_decl)) == POINTER_TYPE; + enum gomp_map_kind kind, ptr_kind, nc_kind; + tree size; + + if (ptr_p) + { + /* For pointers, default mapped as zero-length array + section. */ + kind = GOMP_MAP_ALLOC; + nc_kind = GOMP_MAP_FIRSTPRIVATE_POINTER; + ptr_kind = GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION; + size = size_zero_node; + } + else + { + /* For references, default mapped as appearing on map + clause. */ + kind = GOMP_MAP_TOFROM; + nc_kind = GOMP_MAP_FIRSTPRIVATE_REFERENCE; + ptr_kind = GOMP_MAP_ALWAYS_POINTER; + size = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure_expr))); + } + + for (tree *p = clauses_ptr; *p; p = &OMP_CLAUSE_CHAIN (*p)) + if (OMP_CLAUSE_CODE (*p) == OMP_CLAUSE_MAP + && (TREE_CODE (OMP_CLAUSE_DECL (*p)) == INDIRECT_REF + || TREE_CODE (OMP_CLAUSE_DECL (*p)) == MEM_REF) + && operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*p), 0), + orig_decl)) + { + /* If this was already specified by user as a map, + save the user specified map kind, delete the + "map(*ptr/ref), map(firstprivate ptr/ref)" sequence, + and insert our own sequence: + "map(*__closure->ptr/ref), map(<ptr_kind>:__closure->ref" + */ + tree nc = OMP_CLAUSE_CHAIN (*p); + gcc_assert (nc != NULL_TREE + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (nc) == nc_kind); + /* Update with user specified kind and size. */ + kind = OMP_CLAUSE_MAP_KIND (*p); + size = OMP_CLAUSE_SIZE (*p); + *p = OMP_CLAUSE_CHAIN (nc); + break; + } + + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, kind); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, closure_expr, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) = size; + if (ptr_p) + 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, ptr_kind); + OMP_CLAUSE_DECL (c) = closure_expr; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + } + } + + 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: + ; + } + + for (hash_set<tree>::iterator i = data.lambda_objects_accessed.begin (); + i != data.lambda_objects_accessed.end (); ++i) + { + tree lobj = *i; + if (TREE_CODE (lobj) == TARGET_EXPR) + lobj = TREE_OPERAND (lobj, 0); + + 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.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_TOFROM); + 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; + OMP_TARGET_CLAUSES (stmt) = clauses; + OMP_TARGET_BODY (stmt) = body; + OMP_TARGET_COMBINED (stmt) = combined_p; + SET_EXPR_LOCATION (stmt, loc); + + tree c = clauses; + while (c) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + break; + default: + error_at (OMP_CLAUSE_LOCATION (c), + "%<#pragma omp target%> with map-type other " + "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> " + "on %<map%> clause"); + break; + } + c = OMP_CLAUSE_CHAIN (c); + } + return add_stmt (stmt); +} + tree finish_omp_parallel (tree clauses, tree body) { diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 4e022d860a1..ed46fe3c461 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" @@ -8927,7 +8928,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_map<tree_operand_hash, tree> *struct_map_to_clause = NULL; hash_set<tree> *struct_deref_set = NULL; tree *prev_list_p = NULL, *orig_list_p = list_p; int handled_depend_iterators = -1; @@ -9371,7 +9372,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, GOVD_FIRSTPRIVATE | GOVD_SEEN); } - if (!DECL_P (decl)) + if (TREE_CODE (decl) == TARGET_EXPR) + { + if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + is_gimple_lvalue, fb_lvalue) + == GS_ERROR) + remove = true; + } + else if (!DECL_P (decl)) { tree d = decl, *pd; if (TREE_CODE (d) == ARRAY_REF) @@ -9387,12 +9395,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && TREE_CODE (decl) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE)) + == REFERENCE_TYPE) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)) { pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } bool indir_p = false; + bool component_ref_p = false; tree orig_decl = decl; tree decl_ref = NULL_TREE; if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0 @@ -9403,6 +9414,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, while (TREE_CODE (decl) == COMPONENT_REF) { decl = TREE_OPERAND (decl, 0); + component_ref_p = true; if (((TREE_CODE (decl) == MEM_REF && integer_zerop (TREE_OPERAND (decl, 1))) || INDIRECT_REF_P (decl)) @@ -9411,6 +9423,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { indir_p = true; decl = TREE_OPERAND (decl, 0); + STRIP_NOPS (decl); } if (TREE_CODE (decl) == INDIRECT_REF && DECL_P (TREE_OPERAND (decl, 0)) @@ -9422,8 +9435,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } - else if (TREE_CODE (decl) == COMPONENT_REF) + else if (TREE_CODE (decl) == COMPONENT_REF + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) { + component_ref_p = true; while (TREE_CODE (decl) == COMPONENT_REF) decl = TREE_OPERAND (decl, 0); if (TREE_CODE (decl) == INDIRECT_REF @@ -9493,7 +9509,10 @@ 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 ((DECL_P (decl) + || (component_ref_p + && (INDIRECT_REF_P (decl) + || TREE_CODE (decl) == MEM_REF))) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH @@ -9550,7 +9569,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gcc_assert (base == decl); splay_tree_node n - = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); + = (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) @@ -9576,7 +9598,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, 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); @@ -9587,7 +9613,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (base_ref) OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); else - OMP_CLAUSE_DECL (l) = decl; + { + OMP_CLAUSE_DECL (l) = unshare_expr (decl); + if (!DECL_P (OMP_CLAUSE_DECL (l)) + && (gimplify_expr (&OMP_CLAUSE_DECL (l), + pre_p, NULL, is_gimple_lvalue, + fb_lvalue) + == GS_ERROR)) + { + remove = true; + break; + } + } OMP_CLAUSE_SIZE (l) = (!attach ? size_int (1) @@ -9595,7 +9632,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, ? 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 + = new hash_map<tree_operand_hash, tree>; struct_map_to_clause->put (decl, l); if (ptr || attach_detach) { @@ -9629,15 +9667,41 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, flags |= GOVD_SEEN; if (has_attachments) flags |= GOVD_MAP_HAS_ATTACHMENTS; - goto do_add_decl; + + /* If this is a *pointer-to-struct expression, make sure a + firstprivate map of the base-pointer exists. */ + if (component_ref_p + && ((TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1))) + || INDIRECT_REF_P (decl)) + && DECL_P (TREE_OPERAND (decl, 0)) + && !splay_tree_lookup (ctx->variables, + ((splay_tree_key) + TREE_OPERAND (decl, 0)))) + { + decl = TREE_OPERAND (decl, 0); + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + enum gomp_map_kind mkind + = GOMP_MAP_FIRSTPRIVATE_POINTER; + OMP_CLAUSE_SET_MAP_KIND (c2, mkind); + OMP_CLAUSE_DECL (c2) = decl; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = c2; + } + + if (DECL_P (decl)) + goto do_add_decl; } else if (struct_map_to_clause) { tree *osc = struct_map_to_clause->get (decl); tree *sc = NULL, *scp = NULL; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) - || ptr - || attach_detach) + if (n != NULL + && (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) + || ptr + || attach_detach)) n->value |= GOVD_SEEN; sc = &OMP_CLAUSE_CHAIN (*osc); if (*sc != c @@ -9738,6 +9802,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } else if (*sc != c) { + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) + == GS_ERROR) + { + remove = true; + break; + } *list_p = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = *sc; *sc = c; @@ -9873,6 +9944,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } + /* If this was of the form map(*pointer_to_struct), then the + 'pointer_to_struct' DECL should be considered deref'ed. */ + if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC + || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c)) + || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c))) + && INDIRECT_REF_P (orig_decl) + && DECL_P (TREE_OPERAND (orig_decl, 0)) + && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE) + { + tree ptr = TREE_OPERAND (orig_decl, 0); + if (!struct_deref_set || !struct_deref_set->contains (ptr)) + { + if (!struct_deref_set) + struct_deref_set = new hash_set<tree> (); + struct_deref_set->add (ptr); + } + } + if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH @@ -11222,6 +11311,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT + && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)) + { + remove = true; + break; + } if (!DECL_P (decl)) { if ((ctx->region_type & ORT_TARGET) != 0 @@ -11268,10 +11363,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c)); } } - else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - && (code == OMP_TARGET_EXIT_DATA - || code == OACC_EXIT_DATA)) - remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 63a47f62d08..707cc4606c8 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12631,6 +12631,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_ALWAYS_POINTER: case GOMP_MAP_ATTACH: case GOMP_MAP_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: break; case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C new file mode 100644 index 00000000000..f4d40ec8e4b --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-3.C @@ -0,0 +1,36 @@ +// { dg-do compile } +// { dg-options "-fopenmp -fdump-tree-gimple" } + +struct S +{ + int a, b; + void bar (int); +}; + +void +S::bar (int x) +{ + #pragma omp target map (alloc: a, b) + ; + #pragma omp target enter data map (alloc: a, b) +} + +template <int N> +struct T +{ + int a, b; + void bar (int); +}; + +template <int N> +void +T<N>::bar (int x) +{ + #pragma omp target map (alloc: a, b) + ; + #pragma omp target enter data map (alloc: a, b) +} + +template struct T<0>; + +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */ 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..7f83f92ec93 --- /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.* 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\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[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\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-2.C b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C new file mode 100644 index 00000000000..bdf2564cd04 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C @@ -0,0 +1,35 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include <cstdlib> + +#define N 10 +int main (void) +{ + int X, Y; + #pragma omp target map(from: X, Y) + { + int x = 0, y = 0; + + for (int i = 0; i < N; i++) + [&] (int v) { x += v; } (i); + + auto yinc = [&y] { y++; }; + for (int i = 0; i < N; i++) + yinc (); + + X = x; + Y = y; + } + + int Xs = 0; + for (int i = 0; i < N; i++) + Xs += i; + if (X != Xs) + abort (); + + if (Y != N) + abort (); +} + +/* Make sure lambda objects do NOT appear in target maps. */ +/* { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(from:Y \[len: [0-9]+\]\) map\(from:X \[len: [0-9]+\]\)$} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-1.C b/gcc/testsuite/g++.dg/gomp/target-this-1.C new file mode 100644 index 00000000000..de93a3e5e57 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-1.C @@ -0,0 +1,33 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +extern "C" void abort (); + +struct S +{ + int a, b, c, d; + + int sum (void) + { + int val = 0; + val += a + b + this->c + this->d; + return val; + } + + int sum_offload (void) + { + int val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} + +/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C new file mode 100644 index 00000000000..8a76bb836f8 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C @@ -0,0 +1,49 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-do compile } +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } + +extern "C" void abort (); + +struct T +{ + int x, y; + + auto sum_func (int n) + { + auto fn = [=](int m) -> int + { + int v; + v = (x + y) * n + m; + return v; + }; + return fn; + } + + auto sum_func_offload (int n) + { + auto fn = [=](int m) -> int + { + int v; + #pragma omp target map(from:v) + v = (x + y) * n + m; + return v; + }; + return fn; + } + +}; + +int main (void) +{ + T a = { 1, 2 }; + + auto s1 = a.sum_func (3); + auto s2 = a.sum_func_offload (3); + + if (s1 (1) != s2 (1)) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C new file mode 100644 index 00000000000..91cfbd6ef20 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C @@ -0,0 +1,105 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +#include <cstdlib> +#include <cstring> +extern "C" void abort (); + +struct S +{ + int * ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + bool set_ptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr != NULL) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + } + + bool set_refptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr != NULL) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + S s = { ptr1, N, ptr2, N }; + + bool mapped; + int val = 123; + + mapped = s.set_ptr (val); + if (mapped) + abort (); + if (s.ptr != ptr1) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + mapped = s.set_refptr (val); + if (mapped) + abort (); + if (s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N]) + mapped = s.set_ptr (val); + + if (!mapped) + abort (); + if (s.set_refptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != val) + abort (); + + #pragma omp target data map(ptr2[:N]) + mapped = s.set_refptr (val); + + if (!mapped) + abort (); + if (s.set_ptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != val) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C new file mode 100644 index 00000000000..e4b2a71bbb4 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C @@ -0,0 +1,107 @@ +// 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> + +struct T +{ + int *ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + auto set_ptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + }; + return fn; + } + + auto set_refptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + }; + return fn; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + T a = { ptr1, N, ptr2, N }; + + auto p1 = a.set_ptr_func (1); + auto r2 = a.set_refptr_func (2); + + if (p1 ()) + abort (); + if (r2 ()) + abort (); + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N], ptr2[:N]) + { + if (!p1 ()) + abort (); + if (!r2 ()) + abort (); + } + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 1) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 2) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) 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\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-5.C b/gcc/testsuite/g++.dg/gomp/target-this-5.C new file mode 100644 index 00000000000..a9ac74bcf1f --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-5.C @@ -0,0 +1,34 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +extern "C" void abort (); + +template<typename T> +struct S +{ + T a, b, c, d; + + T sum (void) + { + T val = 0; + val += a + b + this->c + this->d; + return val; + } + + T sum_offload (void) + { + T val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S<int> s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} + +/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/this-2.C b/gcc/testsuite/g++.dg/gomp/this-2.C index d03b8a0728e..b521a4faf5e 100644 --- a/gcc/testsuite/g++.dg/gomp/this-2.C +++ b/gcc/testsuite/g++.dg/gomp/this-2.C @@ -9,14 +9,14 @@ struct S void S::bar (int x) { - #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } ; - #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this[0], x) ; - #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update to (this[0], x) + #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update from (this[1], x) } template <int N> @@ -29,14 +29,14 @@ template <int N> void T<N>::bar (int x) { - #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } ; - #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this[0], x) ; - #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update to (this[0], x) + #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update from (this[1], x) } template struct T<0>; diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c new file mode 100644 index 00000000000..3e7921270c9 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/target-3.c @@ -0,0 +1,16 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +struct S +{ + int a, b; +}; + +void foo (struct S *s) +{ + #pragma omp target map (alloc: s->a, s->b) + ; + #pragma omp target enter data map (alloc: s->a, s->b) +} + +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index fcc0796e3a1..a81ba401ef9 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -858,6 +858,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) { case GOMP_MAP_ALLOC: case GOMP_MAP_POINTER: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, "alloc"); break; case GOMP_MAP_IF_PRESENT: @@ -936,6 +937,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ATTACH_DETACH: pp_string (pp, "attach_detach"); break; + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + pp_string (pp, "attach_zero_length_array_section"); + break; default: gcc_unreachable (); } @@ -954,6 +958,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ALWAYS_POINTER: pp_string (pp, " [pointer assign, bias: "); break; + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + pp_string (pp, " [pointer assign, zero-length array section, bias: "); + break; case GOMP_MAP_TO_PSET: pp_string (pp, " [pointer set, len: "); break; @@ -961,6 +968,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_DETACH: case GOMP_MAP_FORCE_DETACH: case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, " [bias: "); break; default: diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 3e42d7123ae..9e7db69f082 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -143,6 +143,11 @@ enum gomp_map_kind No refcount is bumped by this, and the store is done unconditionally. */ GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_FLAG_SPECIAL | 1), + /* Like GOMP_MAP_POINTER, but allow zero-length array section, i.e. set to + NULL if target is not mapped. */ + GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION + = (GOMP_MAP_FLAG_SPECIAL_2 + | GOMP_MAP_FLAG_SPECIAL | 2), /* Forced deallocation of zero length array section. */ GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION = (GOMP_MAP_FLAG_SPECIAL_2 @@ -163,6 +168,12 @@ enum gomp_map_kind GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY | GOMP_MAP_FLAG_FORCE | 1), + /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections + (i.e. set to NULL when array section is not mapped) Currently only used + by OpenMP. */ + GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION + = (GOMP_MAP_DEEP_COPY | 2), + /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), @@ -186,7 +197,8 @@ enum gomp_map_kind ((X) == GOMP_MAP_ALWAYS_POINTER) #define GOMP_MAP_POINTER_P(X) \ - ((X) == GOMP_MAP_POINTER) + ((X) == GOMP_MAP_POINTER \ + || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION) #define GOMP_MAP_ALWAYS_TO_P(X) \ (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM)) diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index ceef643216c..0fb0b783660 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1279,7 +1279,7 @@ extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t); extern void gomp_attach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree, splay_tree_key, uintptr_t, size_t, - struct gomp_coalesce_buf *); + struct gomp_coalesce_buf *, bool); extern void gomp_detach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree_key, uintptr_t, bool, struct gomp_coalesce_buf *); diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 5988db0b886..82d8dacfa1c 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -937,7 +937,7 @@ acc_attach_async (void **hostaddr, int async) } gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr, - 0, NULL); + 0, NULL, false); gomp_mutex_unlock (&acc_dev->lock); } @@ -1141,7 +1141,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH) { gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, - (uintptr_t) h, s, NULL); + (uintptr_t) h, s, NULL, false); /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ } @@ -1159,7 +1159,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, splay_tree_key m = lookup_host (acc_dev, hostaddrs[j], sizeof (void *)); gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m, - (uintptr_t) hostaddrs[j], sizes[j], NULL); + (uintptr_t) hostaddrs[j], sizes[j], NULL, + false); } bool processed = false; diff --git a/libgomp/target.c b/libgomp/target.c index 3c1eee23a44..bb31b1991d1 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -543,7 +543,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, struct gomp_coalesce_buf *cbuf, htab_t *refcount_set) { - assert (kind != GOMP_MAP_ATTACH); + assert (kind != GOMP_MAP_ATTACH + || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); tgt_var->key = oldn; tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind); @@ -616,7 +617,8 @@ get_implicit (bool short_mapkind, void *kinds, int idx) static void gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, uintptr_t host_ptr, uintptr_t target_offset, uintptr_t bias, - struct gomp_coalesce_buf *cbuf) + struct gomp_coalesce_buf *cbuf, + bool allow_zero_length_array_sections) { struct gomp_device_descr *devicep = tgt->device_descr; struct splay_tree_s *mem_map = &devicep->mem_map; @@ -638,16 +640,24 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n == NULL) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("Pointer target of array section wasn't mapped"); - } - cur_node.host_start -= n->host_start; - cur_node.tgt_offset - = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start; - /* At this point tgt_offset is target address of the - array section. Now subtract bias to get what we want - to initialize the pointer with. */ - cur_node.tgt_offset -= bias; + if (allow_zero_length_array_sections) + cur_node.tgt_offset = 0; + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("Pointer target of array section wasn't mapped"); + } + } + else + { + cur_node.host_start -= n->host_start; + cur_node.tgt_offset + = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start; + /* At this point tgt_offset is target address of the + array section. Now subtract bias to get what we want + to initialize the pointer with. */ + cur_node.tgt_offset -= bias; + } gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset), (void *) &cur_node.tgt_offset, sizeof (void *), true, cbuf); @@ -724,7 +734,8 @@ attribute_hidden void gomp_attach_pointer (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, splay_tree mem_map, splay_tree_key n, uintptr_t attach_to, size_t bias, - struct gomp_coalesce_buf *cbufp) + struct gomp_coalesce_buf *cbufp, + bool allow_zero_length_array_sections) { struct splay_tree_key_s s; size_t size, idx; @@ -776,11 +787,21 @@ gomp_attach_pointer (struct gomp_device_descr *devicep, if (!tn) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("pointer target not mapped for attach"); + if (allow_zero_length_array_sections) + { + /* When allowing attachment to zero-length array sections, we + allow attaching to NULL pointers when the target region is not + mapped. */ + data = 0; + } + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("pointer target not mapped for attach"); + } } - - data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; + else + data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; gomp_debug (1, "%s: attaching host %p, target %p (struct base %p) to %p\n", @@ -1038,7 +1059,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, has_firstprivate = true; continue; } - else if ((kind & typemask) == GOMP_MAP_ATTACH) + else if ((kind & typemask) == GOMP_MAP_ATTACH + || ((kind & typemask) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) { tgt->list[i].key = NULL; has_firstprivate = true; @@ -1287,7 +1310,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, (uintptr_t) *(void **) hostaddrs[j], k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp); + sizes[j], cbufp, false); } } i = j - 1; @@ -1416,6 +1439,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, ++i; continue; case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: { cur_node.host_start = (uintptr_t) hostaddrs[i]; cur_node.host_end = cur_node.host_start + sizeof (void *); @@ -1432,9 +1456,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ + bool zlas + = ((kind & typemask) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); gomp_attach_pointer (devicep, aq, mem_map, n, (uintptr_t) hostaddrs[i], sizes[i], - cbufp); + cbufp, zlas); } else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0) { @@ -1545,9 +1572,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, false, cbufp); break; case GOMP_MAP_POINTER: - gomp_map_pointer (tgt, aq, - (uintptr_t) *(void **) k->host_start, - k->tgt_offset, sizes[i], cbufp); + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + gomp_map_pointer + (tgt, aq, (uintptr_t) *(void **) k->host_start, + k->tgt_offset, sizes[i], cbufp, + ((kind & typemask) + == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)); break; case GOMP_MAP_TO_PSET: gomp_copy_host2dev (devicep, aq, @@ -1589,7 +1619,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp); + sizes[j], cbufp, false); } } i = j - 1; diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C new file mode 100644 index 00000000000..d4f9ff3e983 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-23.C @@ -0,0 +1,34 @@ +extern "C" void abort (); + +struct S +{ + int *data; +}; + +int +main (void) +{ + #define SZ 10 + S *s = new S (); + s->data = new int[SZ]; + + for (int i = 0; i < SZ; i++) + s->data[i] = 0; + + #pragma omp target enter data map(to: s) + #pragma omp target enter data map(to: s->data[:SZ]) + #pragma omp target + { + for (int i = 0; i < SZ; i++) + s->data[i] = i; + } + #pragma omp target exit data map(from: s->data[:SZ]) + #pragma omp target exit data map(from: s) + + for (int i = 0; i < SZ; i++) + if (s->data[i] != i) + abort (); + + return 0; +} + 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; +} diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-2.C b/libgomp/testsuite/libgomp.c++/target-lambda-2.C new file mode 100644 index 00000000000..1d3561ffbd7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-lambda-2.C @@ -0,0 +1,30 @@ +#include <cstdlib> + +#define N 10 +int main (void) +{ + int X, Y; + #pragma omp target map(from: X, Y) + { + int x = 0, y = 0; + + for (int i = 0; i < N; i++) + [&] (int v) { x += v; } (i); + + auto yinc = [&y] { y++; }; + for (int i = 0; i < N; i++) + yinc (); + + X = x; + Y = y; + } + + int Xs = 0; + for (int i = 0; i < N; i++) + Xs += i; + if (X != Xs) + abort (); + + if (Y != N) + abort (); +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-1.C b/libgomp/testsuite/libgomp.c++/target-this-1.C new file mode 100644 index 00000000000..a591ea4c564 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-1.C @@ -0,0 +1,29 @@ +extern "C" void abort (); + +struct S +{ + int a, b, c, d; + + int sum (void) + { + int val = 0; + val += a + b + this->c + this->d; + return val; + } + + int sum_offload (void) + { + int val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-2.C b/libgomp/testsuite/libgomp.c++/target-this-2.C new file mode 100644 index 00000000000..8119be8c2c5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-2.C @@ -0,0 +1,47 @@ + +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14" } + +extern "C" void abort (); + +struct T +{ + int x, y; + + auto sum_func (int n) + { + auto fn = [=](int m) -> int + { + int v; + v = (x + y) * n + m; + return v; + }; + return fn; + } + + auto sum_func_offload (int n) + { + auto fn = [=](int m) -> int + { + int v; + #pragma omp target map(from:v) + v = (x + y) * n + m; + return v; + }; + return fn; + } + +}; + +int main (void) +{ + T a = { 1, 2 }; + + auto s1 = a.sum_func (3); + auto s2 = a.sum_func_offload (3); + + if (s1 (1) != s2 (1)) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-3.C b/libgomp/testsuite/libgomp.c++/target-this-3.C new file mode 100644 index 00000000000..e15f69a1623 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-3.C @@ -0,0 +1,99 @@ +#include <stdio.h> +#include <string.h> +extern "C" void abort (); + +struct S +{ + int * ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + bool set_ptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr != NULL) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + } + + bool set_refptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr != NULL) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + S s = { ptr1, N, ptr2, N }; + + bool mapped; + int val = 123; + + mapped = s.set_ptr (val); + if (mapped) + abort (); + if (s.ptr != ptr1) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + mapped = s.set_refptr (val); + if (mapped) + abort (); + if (s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N]) + mapped = s.set_ptr (val); + + if (!mapped) + abort (); + if (s.set_refptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != val) + abort (); + + #pragma omp target data map(ptr2[:N]) + mapped = s.set_refptr (val); + + if (!mapped) + abort (); + if (s.set_ptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != val) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-4.C b/libgomp/testsuite/libgomp.c++/target-this-4.C new file mode 100644 index 00000000000..9f53677a240 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-4.C @@ -0,0 +1,104 @@ + +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14" } +#include <cstdlib> +#include <cstring> + +struct T +{ + int *ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + auto set_ptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + }; + return fn; + } + + auto set_refptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + }; + return fn; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + T a = { ptr1, N, ptr2, N }; + + auto p1 = a.set_ptr_func (1); + auto r2 = a.set_refptr_func (2); + + if (p1 ()) + abort (); + if (r2 ()) + abort (); + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N], ptr2[:N]) + { + if (!p1 ()) + abort (); + if (!r2 ()) + abort (); + } + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 1) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 2) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-5.C b/libgomp/testsuite/libgomp.c++/target-this-5.C new file mode 100644 index 00000000000..e71c566687d --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-5.C @@ -0,0 +1,30 @@ +extern "C" void abort (); + +template<typename T> +struct S +{ + T a, b, c, d; + + T sum (void) + { + T val = 0; + val += a + b + this->c + this->d; + return val; + } + + T sum_offload (void) + { + T val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S<int> s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} ^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [PATCH, v5, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v5] 2021-11-16 12:43 ` [PATCH, v5, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v5] Chung-Lin Tang @ 2021-12-03 16:47 ` Jakub Jelinek 2021-12-09 16:41 ` Chung-Lin Tang 2022-09-07 7:04 ` [committed] openmp: Fix handling of target constructs in static member functions [PR106829] Jakub Jelinek 1 sibling, 1 reply; 6+ messages in thread From: Jakub Jelinek @ 2021-12-03 16:47 UTC (permalink / raw) To: Chung-Lin Tang; +Cc: gcc-patches, Catherine Moore On Tue, Nov 16, 2021 at 08:43:27PM +0800, Chung-Lin Tang wrote: > 2021-11-16 Chung-Lin Tang <cltang@codesourcery.com> > > PR middle-end/92120 > > gcc/cp/ChangeLog: > > * cp-tree.h (finish_omp_target): New declaration. > (finish_omp_target_clauses): Likewise. > * parser.c (cp_parser_omp_clause_map): Adjust call to > cp_parser_omp_var_list_no_open to set 'allow_deref' argument to true. > (cp_parser_omp_target): Factor out code, adjust into calls to new > function finish_omp_target. > * pt.c (tsubst_expr): Add call to finish_omp_target_clauses for > OMP_TARGET case. > * semantics.c (handle_omp_array_sections_1): Add handling to create > 'this->member' from 'member' FIELD_DECL. Remove case of rejecting > 'this' when not in declare simd. > (handle_omp_array_sections): Likewise. > (finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP > map clauses. Handle 'A->member' case in map clauses. Remove case of > rejecting 'this' when not in declare simd. > (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. > (finish_omp_target): New function. > > gcc/c/ChangeLog: > > * c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in > call to c_parser_omp_variable_list to 'true'. > * c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in > array base handling. > (c_finish_omp_clauses): Handle 'A->member' case in map clauses. > > gcc/ChangeLog: > > * gimplify.c ("tree-hash-traits.h"): Add include. > (gimplify_scan_omp_clauses): Change struct_map_to_clause to type > hash_map<tree_operand, tree> *. Adjust struct map handling to handle > cases of *A and A->B expressions. Under !DECL_P case of > GOMP_CLAUSE_MAP handling, add STRIP_NOPS for indir_p case, add to > struct_deref_set for map(*ptr_to_struct) cases. Add MEM_REF case when > handling component_ref_p case. Add unshare_expr and gimplification > when created GOMP_MAP_STRUCT is not a DECL. Add code to add > firstprivate pointer for *pointer-to-struct case. > (gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for > exit data directives code to earlier position. > * omp-low.c (lower_omp_target): > Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and > GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. > * tree-pretty-print.c (dump_omp_clause): Likewise. > > gcc/testsuite/ChangeLog: > > * gcc.dg/gomp/target-3.c: New testcase. > * g++.dg/gomp/target-3.C: New testcase. > * g++.dg/gomp/target-lambda-1.C: New testcase. > * g++.dg/gomp/target-lambda-2.C: New testcase. > * g++.dg/gomp/target-this-1.C: New testcase. > * g++.dg/gomp/target-this-2.C: New testcase. > * g++.dg/gomp/target-this-3.C: New testcase. > * g++.dg/gomp/target-this-4.C: New testcase. > * g++.dg/gomp/target-this-5.C: New testcase. > * g++.dg/gomp/this-2.C: Adjust testcase. > > include/ChangeLog: > > * gomp-constants.h (enum gomp_map_kind): > Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and > GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. > (GOMP_MAP_POINTER_P): > Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION. > > libgomp/ChangeLog: > > * libgomp.h (gomp_attach_pointer): Add bool parameter. > * oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer. > (goacc_enter_data_internal): Likewise. > * target.c (gomp_map_vars_existing): Update assert condition to > include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION. > (gomp_map_pointer): Add 'bool allow_zero_length_array_sections' > parameter, add support for mapping a pointer with NULL target. > (gomp_attach_pointer): Add 'bool allow_zero_length_array_sections' > parameter, add support for attaching a pointer with NULL target. > (gomp_map_vars_internal): Update calls to gomp_map_pointer and > gomp_attach_pointer, add handling for > GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and > GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases. > * testsuite/libgomp.c++/target-23.C: New testcase. > * testsuite/libgomp.c++/target-lambda-1.C: New testcase. > * testsuite/libgomp.c++/target-lambda-2.C: New testcase. > * testsuite/libgomp.c++/target-this-1.C: New testcase. > * testsuite/libgomp.c++/target-this-2.C: New testcase. > * testsuite/libgomp.c++/target-this-3.C: New testcase. > * testsuite/libgomp.c++/target-this-4.C: New testcase. > * testsuite/libgomp.c++/target-this-5.C: New testcase. > +/* 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; > + > + hash_set<tree> local_decls; Can you please add short comments above the members describing what they are for? > +}; > + > +static tree > +finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr) And add function comments above here. > + > +void > +finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr) And here. > + > +tree > +finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p) And here? > + if (allow_zero_length_array_sections) > + { > + /* When allowing attachment to zero-length array sections, we > + allow attaching to NULL pointers when the target region is not > + mapped. */ > + data = 0; > + } No {}s around single statement if body. Otherwise LGTM. Jakub ^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [PATCH, v5, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v5] 2021-12-03 16:47 ` Jakub Jelinek @ 2021-12-09 16:41 ` Chung-Lin Tang 0 siblings, 0 replies; 6+ messages in thread From: Chung-Lin Tang @ 2021-12-09 16:41 UTC (permalink / raw) To: Jakub Jelinek, Thomas Schwinge; +Cc: gcc-patches, Catherine Moore On 2021/12/4 12:47 AM, Jakub Jelinek wrote: > On Tue, Nov 16, 2021 at 08:43:27PM +0800, Chung-Lin Tang wrote: >> 2021-11-16 Chung-Lin Tang <cltang@codesourcery.com> >> >> PR middle-end/92120 >> >> gcc/cp/ChangeLog: >> ... >> + if (allow_zero_length_array_sections) >> + { >> + /* When allowing attachment to zero-length array sections, we >> + allow attaching to NULL pointers when the target region is not >> + mapped. */ >> + data = 0; >> + } > > No {}s around single statement if body. > > Otherwise LGTM. > > Jakub > Thanks for the review and approval, Jakub. Thomas, I pushed another 2766448c5cc3efc4 commit to fix the non-offload config FAILs, just FYI. Chung-Lin ^ permalink raw reply [flat|nested] 6+ messages in thread
* [committed] openmp: Fix handling of target constructs in static member functions [PR106829] 2021-11-16 12:43 ` [PATCH, v5, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v5] Chung-Lin Tang 2021-12-03 16:47 ` Jakub Jelinek @ 2022-09-07 7:04 ` Jakub Jelinek 1 sibling, 0 replies; 6+ messages in thread From: Jakub Jelinek @ 2022-09-07 7:04 UTC (permalink / raw) To: Chung-Lin Tang; +Cc: gcc-patches On Tue, Nov 16, 2021 at 08:43:27PM +0800, Chung-Lin Tang wrote: > 2021-11-16 Chung-Lin Tang <cltang@codesourcery.com> > > PR middle-end/92120 > > gcc/cp/ChangeLog: > > * semantics.c (handle_omp_array_sections_1): Add handling to create ... > (finish_omp_target_clauses): New function. Just calling current_nonlambda_class_type in static member functions returns non-NULL, but something that isn't *this and if unlucky can match part of the IL and can be added to target clauses. if (DECL_NONSTATIC_MEMBER_P (decl) && current_class_ptr) is a guard used elsewhere (in check_accessibility_of_qualified_id). Fixed thusly, bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk, queued for 12.3 backport. 2022-09-07 Jakub Jelinek <jakub@redhat.com> PR c++/106829 * semantics.cc (finish_omp_target_clauses): If current_function_decl isn't a nonstatic member function, don't set data.current_object to non-NULL. * g++.dg/gomp/pr106829.C: New test. --- gcc/cp/semantics.cc.jj 2022-09-03 09:41:34.892005463 +0200 +++ gcc/cp/semantics.cc 2022-09-06 15:00:33.253199294 +0200 @@ -9555,16 +9555,15 @@ finish_omp_target_clauses (location_t lo { omp_target_walk_data data; data.this_expr_accessed = false; + data.current_object = NULL_TREE; - 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_NONSTATIC_MEMBER_P (current_function_decl) && current_class_ptr) + if (tree ct = current_nonlambda_class_type ()) + { + tree object = maybe_dummy_object (ct, NULL); + object = maybe_resolve_dummy (object, true); + data.current_object = object; + } if (DECL_LAMBDA_FUNCTION_P (current_function_decl)) { --- gcc/testsuite/g++.dg/gomp/pr106829.C.jj 2022-09-06 15:03:44.305635408 +0200 +++ gcc/testsuite/g++.dg/gomp/pr106829.C 2022-09-06 15:03:12.987055704 +0200 @@ -0,0 +1,15 @@ +// PR c++/106829 + +namespace std +{ + template <typename> class complex; + template <> struct complex<double> { complex (double); _Complex double d; }; +} +struct S { void static foo (); }; + +void +S::foo () +{ +#pragma omp target + std::complex<double> c = 0.0; +} Jakub ^ permalink raw reply [flat|nested] 6+ messages in thread
end of thread, other threads:[~2022-09-07 7:04 UTC | newest] Thread overview: 6+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2021-06-18 14:25 [PATCH, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v4] Chung-Lin Tang 2021-06-24 13:15 ` Jakub Jelinek 2021-11-16 12:43 ` [PATCH, v5, OpenMP 5.0] Improve OpenMP target support for C++ [PR92120 v5] Chung-Lin Tang 2021-12-03 16:47 ` Jakub Jelinek 2021-12-09 16:41 ` Chung-Lin Tang 2022-09-07 7:04 ` [committed] openmp: Fix handling of target constructs in static member functions [PR106829] Jakub Jelinek
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).