OpenE2K
/
gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

openmp: Improve OpenMP target support for C++ (PR92120) 

This patch implements several C++ specific mapping capabilities introduced for
OpenMP 5.0, including implicit mapping of this[:1] for non-static member
functions, zero-length array section mapping of pointer-typed members,
lambda captured variable access in target regions, and use of lambda objects
inside target regions.

Several adjustments to the C/C++ front-ends to allow more member-access syntax
as valid is also included.

	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.
This commit is contained in:
Chung-Lin Tang 2021-12-08 22:28:03 +08:00
parent dbf8bd3c2f
commit 0ab29cf0bb
31 changed files with 1765 additions and 142 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
gcc/c/c-parser.c
View File

@ -16194,7 +16194,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);

23
gcc/c/c-typeck.c
View File

@ -13241,6 +13241,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))
@ -14085,6 +14090,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;
@ -14886,6 +14892,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))
@ -14952,6 +14963,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_)
{
@ -15024,6 +15043,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
@ -15080,8 +15100,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),

2
gcc/cp/cp-tree.h
View File

@ -7671,6 +7671,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);

69
gcc/cp/parser.c
View File

@ -39315,7 +39315,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);
@ -44105,8 +44105,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);
@ -44211,16 +44209,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 */
{
@ -44255,13 +44247,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);
@ -44270,45 +44259,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);
tree body = cp_parser_omp_structured_block (parser, if_p);
SET_EXPR_LOCATION (stmt, pragma_tok->location);
add_stmt (stmt);
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;
}

5
gcc/cp/pt.c
View File

@ -18975,6 +18975,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);

581
gcc/cp/semantics.c
View File

@ -5049,15 +5049,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)
{
@ -5073,18 +5074,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))
@ -5599,6 +5588,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))
@ -6611,6 +6602,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
@ -7862,6 +7854,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))
@ -7934,9 +7931,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_)
{
@ -7983,6 +7985,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)
@ -8009,19 +8018,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
@ -9177,6 +9179,533 @@ 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
{
/* Holds the 'this' expression found in current function. */
tree current_object;
/* True if the 'this' expression was accessed in the target body. */
bool this_expr_accessed;
/* For non-static functions, record which pointer-typed members were
accessed, and the whole expression. */
hash_map<tree, tree> ptr_members_accessed;
/* Record which lambda objects were accessed in target body. */
hash_set<tree> lambda_objects_accessed;
/* For lambda functions, the __closure object expression of the current
function, and the set of captured variables accessed in target body. */
tree current_closure;
hash_set<tree> closure_vars_accessed;
/* Local variables declared inside a BIND_EXPR, used to filter out such
variables when recording lambda_objects_accessed. */
hash_set<tree> local_decls;
};
/* Helper function of finish_omp_target_clauses, called via
cp_walk_tree_without_duplicates. Traverse body of OpenMP target
directive *TP, and fill out omp_target_walk_data passed in *PTR. */
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;
}
/* Helper function for finish_omp_target, and also from tsubst_expr.
Create additional clauses for mapping of non-static members, lambda objects,
etc. */
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;
}
/* Called from cp_parser_omp_target. Create additional implicit clauses for
OpenMP target directives, and do sanity checks. */
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)
{

125
gcc/gimplify.c
View File

@ -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"
@ -8921,7 +8922,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;
@ -9365,7 +9366,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)
@ -9381,12 +9389,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
@ -9397,6 +9408,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))
@ -9405,6 +9417,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))
@ -9416,8 +9429,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
@ -9487,7 +9503,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
@ -9544,7 +9563,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)
@ -9570,7 +9592,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);
@ -9581,7 +9607,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)
@ -9589,7 +9626,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)
{
@ -9623,15 +9661,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
@ -9732,6 +9796,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;
@ -9867,6 +9938,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
@ -11216,6 +11305,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
@ -11262,10 +11357,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

2
gcc/omp-low.c
View File

@ -12635,6 +12635,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:

36
gcc/testsuite/g++.dg/gomp/target-3.C Normal file
View File

@ -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" } } */

94
gcc/testsuite/g++.dg/gomp/target-lambda-1.C Normal file
View File

@ -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" } } */

35
gcc/testsuite/g++.dg/gomp/target-lambda-2.C Normal file
View File

@ -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" } } */

33
gcc/testsuite/g++.dg/gomp/target-this-1.C Normal file
View File

@ -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" } } */

49
gcc/testsuite/g++.dg/gomp/target-this-2.C Normal file
View File

@ -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" } } */

105
gcc/testsuite/g++.dg/gomp/target-this-3.C Normal file
View File

@ -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" } } */

107
gcc/testsuite/g++.dg/gomp/target-this-4.C Normal file
View File

@ -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" } } */

34
gcc/testsuite/g++.dg/gomp/target-this-5.C Normal file
View File

@ -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" } } */

24
gcc/testsuite/g++.dg/gomp/this-2.C
View File

@ -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>;

16
gcc/testsuite/gcc.dg/gomp/target-3.c Normal file
View File

@ -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" } } */

8
gcc/tree-pretty-print.c
View File

@ -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:

14
include/gomp-constants.h
View File

@ -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))

2
libgomp/libgomp.h
View File

@ -1283,7 +1283,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 *);

7
libgomp/oacc-mem.c
View File

@ -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;

74
libgomp/target.c
View File

@ -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");
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;
}
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,19 @@ 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 +1057,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 +1308,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 +1437,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 +1454,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 +1570,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 +1617,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;

34
libgomp/testsuite/libgomp.c++/target-23.C Normal file
View File

@ -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;
}

86
libgomp/testsuite/libgomp.c++/target-lambda-1.C Normal file
View File

@ -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;
}

30
libgomp/testsuite/libgomp.c++/target-lambda-2.C Normal file
View File

@ -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 ();
}

29
libgomp/testsuite/libgomp.c++/target-this-1.C Normal file
View File

@ -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;
}

47
libgomp/testsuite/libgomp.c++/target-this-2.C Normal file
View File

@ -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;
}

99
libgomp/testsuite/libgomp.c++/target-this-3.C Normal file
View File

@ -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;
}

104
libgomp/testsuite/libgomp.c++/target-this-4.C Normal file
View File

@ -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;
}

30
libgomp/testsuite/libgomp.c++/target-this-5.C Normal file
View File

@ -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;
}