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);
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))
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;
{
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))
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_)
{
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
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),
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);
}
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);
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);
= 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 */
{
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);
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;
}
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);
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)
{
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))
}
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))
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
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))
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_)
{
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)
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
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)
{
#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"
{
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;
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)
&& 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
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))
{
indir_p = true;
decl = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (decl);
}
if (TREE_CODE (decl) == INDIRECT_REF
&& DECL_P (TREE_OPERAND (decl, 0))
}
}
}
- 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
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
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)
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);
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)
? 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)
{
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
}
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;
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
}
}
}
+ 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
= 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
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:
--- /dev/null
+// { 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" } } */
--- /dev/null
+// 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" } } */
--- /dev/null
+// 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" } } */
--- /dev/null
+// { 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" } } */
--- /dev/null
+// 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" } } */
--- /dev/null
+// { 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" } } */
--- /dev/null
+// 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" } } */
--- /dev/null
+// { 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" } } */
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>
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>;
--- /dev/null
+/* { 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" } } */
{
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:
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 ();
}
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;
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:
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
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),
((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))
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 *);
}
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);
}
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'). */
}
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;
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);
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;
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);
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;
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",
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;
(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;
++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 *);
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)
{
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,
k->tgt_offset
+ ((uintptr_t) hostaddrs[j]
- k->host_start),
- sizes[j], cbufp);
+ sizes[j], cbufp, false);
}
}
i = j - 1;
--- /dev/null
+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;
+}
+
--- /dev/null
+#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;
+}
--- /dev/null
+#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 ();
+}
--- /dev/null
+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;
+}
--- /dev/null
+
+// 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;
+}
--- /dev/null
+#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;
+}
--- /dev/null
+
+// 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;
+}
--- /dev/null
+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;
+}
projects / test_jj.git / commitdiff