@@ -1257,6 +1257,9 @@ extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree);
extern tree c_omp_check_context_selector (location_t, tree);
extern void c_omp_mark_declare_variant (location_t, tree, tree);
extern void c_omp_adjust_map_clauses (tree, bool);
+template<typename T> struct omp_mapper_list;
+extern void c_omp_find_nested_mappers (struct omp_mapper_list<tree> *, tree);
+extern tree c_omp_instantiate_mappers (tree);
namespace omp_addr_tokenizer { struct omp_addr_token; }
typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;
@@ -3937,6 +3937,306 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
return error_mark_node;
}
+/* Given a mapper function MAPPER_FN, recursively scan through the map clauses
+ for that mapper, and if any of those should use a (named or unnamed) mapper
+ themselves, add it to MLIST. */
+
+void
+c_omp_find_nested_mappers (omp_mapper_list<tree> *mlist, tree mapper_fn)
+{
+ tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
+ tree mapper_name = NULL_TREE;
+
+ if (mapper == error_mark_node)
+ return;
+
+ gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+ for (tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+ clause;
+ clause = OMP_CLAUSE_CHAIN (clause))
+ {
+ tree expr = OMP_CLAUSE_DECL (clause);
+ enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause);
+ tree elem_type;
+
+ if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+ {
+ mapper_name = expr;
+ continue;
+ }
+ else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ mapper_name = NULL_TREE;
+ continue;
+ }
+
+ gcc_assert (TREE_CODE (expr) != TREE_LIST);
+ if (TREE_CODE (expr) == OMP_ARRAY_SECTION)
+ {
+ while (TREE_CODE (expr) == OMP_ARRAY_SECTION)
+ expr = TREE_OPERAND (expr, 0);
+
+ elem_type = TREE_TYPE (expr);
+ }
+ else
+ elem_type = TREE_TYPE (expr);
+
+ /* This might be too much... or not enough? */
+ while (TREE_CODE (elem_type) == ARRAY_TYPE
+ || TREE_CODE (elem_type) == POINTER_TYPE
+ || TREE_CODE (elem_type) == REFERENCE_TYPE)
+ elem_type = TREE_TYPE (elem_type);
+
+ elem_type = TYPE_MAIN_VARIANT (elem_type);
+
+ if (AGGREGATE_TYPE_P (elem_type)
+ && !mlist->contains (mapper_name, elem_type))
+ {
+ tree nested_mapper_fn
+ = lang_hooks.decls.omp_mapper_lookup (mapper_name, elem_type);
+
+ if (nested_mapper_fn)
+ {
+ mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn);
+ c_omp_find_nested_mappers (mlist, nested_mapper_fn);
+ }
+ else if (mapper_name)
+ {
+ error ("mapper %qE not found for type %qT", mapper_name,
+ elem_type);
+ continue;
+ }
+ }
+ }
+}
+
+struct remap_mapper_decl_info
+{
+ tree dummy_var;
+ tree expr;
+};
+
+/* Helper for rewriting DUMMY_VAR into EXPR in a map clause decl. */
+
+static tree
+remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
+{
+ remap_mapper_decl_info *map_info = (remap_mapper_decl_info *) data;
+
+ if (operand_equal_p (*tp, map_info->dummy_var))
+ {
+ *tp = map_info->expr;
+ *walk_subtrees = 0;
+ }
+
+ return NULL_TREE;
+}
+
+/* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to
+ OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in
+ the mapper declaration. */
+
+static tree *
+omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
+ enum gomp_map_kind outer_kind)
+{
+ tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+ tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
+ tree mapper_name = NULL_TREE;
+
+ remap_mapper_decl_info map_info;
+ map_info.dummy_var = dummy_var;
+ map_info.expr = expr;
+
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ tree unshared = unshare_expr (c);
+ enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c);
+ tree t = OMP_CLAUSE_DECL (unshared);
+ tree type = NULL_TREE;
+ bool nonunit_array_with_mapper = false;
+
+ if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+ {
+ mapper_name = t;
+ continue;
+ }
+ else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ mapper_name = NULL_TREE;
+ continue;
+ }
+
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (c);
+ tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
+
+ if (t2 == t)
+ {
+ nonunit_array_with_mapper = true;
+ /* We'd want use the mapper for the element type if this worked:
+ look that one up. */
+ type = TREE_TYPE (TREE_TYPE (t));
+ }
+ else
+ {
+ t = t2;
+ type = TREE_TYPE (t);
+ }
+ }
+ else
+ type = TREE_TYPE (t);
+
+ gcc_assert (type);
+
+ if (type == error_mark_node)
+ continue;
+
+ walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);
+
+ if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)
+ OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
+
+ type = TYPE_MAIN_VARIANT (type);
+
+ tree mapper_fn = lang_hooks.decls.omp_mapper_lookup (mapper_name, type);
+
+ if (mapper_fn && nonunit_array_with_mapper)
+ {
+ sorry ("user-defined mapper with non-unit length array section");
+ continue;
+ }
+ else if (mapper_fn)
+ {
+ tree nested_mapper
+ = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
+ if (nested_mapper != mapper)
+ {
+ if (clause_kind == GOMP_MAP_UNSET)
+ clause_kind = outer_kind;
+
+ outlist = omp_instantiate_mapper (outlist, nested_mapper,
+ t, clause_kind);
+ continue;
+ }
+ }
+ else if (mapper_name)
+ {
+ error ("mapper %qE not found for type %qT", mapper_name, type);
+ continue;
+ }
+
+ *outlist = unshared;
+ outlist = &OMP_CLAUSE_CHAIN (unshared);
+ }
+
+ return outlist;
+}
+
+/* Given a list of CLAUSES, scan each clause and invoke a user-defined mapper
+ appropriate to the type of the data in that clause, if such a mapper is
+ visible in the current parsing context. */
+
+tree
+c_omp_instantiate_mappers (tree clauses)
+{
+ tree c, *pc, mapper_name = NULL_TREE;
+
+ for (pc = &clauses, c = clauses; c; c = *pc)
+ {
+ bool using_mapper = false;
+
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_MAP:
+ {
+ tree t = OMP_CLAUSE_DECL (c);
+ tree type = NULL_TREE;
+ bool nonunit_array_with_mapper = false;
+
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME)
+ mapper_name = OMP_CLAUSE_DECL (c);
+ else
+ mapper_name = NULL_TREE;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (c);
+ tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
+
+ if (t2 == t)
+ {
+ /* !!! Array sections of size >1 with mappers for elements
+ are hard to support. Do something here. */
+ nonunit_array_with_mapper = true;
+ type = TREE_TYPE (TREE_TYPE (t));
+ }
+ else
+ {
+ t = t2;
+ type = TREE_TYPE (t);
+ }
+ }
+ else
+ type = TREE_TYPE (t);
+
+ if (type == NULL_TREE || type == error_mark_node)
+ {
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
+ enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c);
+ if (kind == GOMP_MAP_UNSET)
+ kind = GOMP_MAP_TOFROM;
+
+ type = TYPE_MAIN_VARIANT (type);
+
+ tree mapper_fn
+ = lang_hooks.decls.omp_mapper_lookup (mapper_name, type);
+
+ if (mapper_fn && nonunit_array_with_mapper)
+ {
+ sorry ("user-defined mapper with non-unit length "
+ "array section");
+ using_mapper = true;
+ }
+ else if (mapper_fn)
+ {
+ tree mapper
+ = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
+ pc = omp_instantiate_mapper (pc, mapper, t, kind);
+ using_mapper = true;
+ }
+ else if (mapper_name)
+ {
+ error ("mapper %qE not found for type %qT", mapper_name, type);
+ using_mapper = true;
+ }
+ }
+ break;
+
+ default:
+ ;
+ }
+
+ if (using_mapper)
+ *pc = OMP_CLAUSE_CHAIN (c);
+ else
+ pc = &OMP_CLAUSE_CHAIN (c);
+ }
+
+ return clauses;
+}
+
const struct c_omp_directive c_omp_directives[] = {
/* Keep this alphabetically sorted by the first word. Non-null second/third
if any should precede null ones. */
@@ -3101,6 +3101,9 @@ reduced_constant_expression_p (tree t)
/* Even if we can't lower this yet, it's constant. */
return true;
+ case OMP_DECLARE_MAPPER:
+ return true;
+
case CONSTRUCTOR:
/* And we need to handle PTRMEM_CST wrapped in a CONSTRUCTOR. */
tree field;
@@ -6895,6 +6898,7 @@ cxx_eval_constant_expression (const constexpr_ctx *ctx, tree t,
case LABEL_EXPR:
case CASE_LABEL_EXPR:
case PREDICT_EXPR:
+ case OMP_DECLARE_MAPPER:
return t;
case PARM_DECL:
@@ -9276,6 +9280,11 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now,
error_at (loc, "expression %qE is not a constant expression", t);
return false;
+ case OMP_DECLARE_MAPPER:
+ /* This can be used to initialize VAR_DECLs: it's treated as a magic
+ constant. */
+ return true;
+
case ASM_EXPR:
if (flags & tf_error)
inline_asm_in_constexpr_error (loc);
@@ -2302,6 +2302,12 @@ cxx_omp_finish_clause (tree c, gimple_seq *, bool /* openacc */)
}
}
+tree
+cxx_omp_finish_mapper_clauses (tree clauses)
+{
+ return finish_omp_clauses (clauses, C_ORT_OMP);
+}
+
/* Return true if DECL's DECL_VALUE_EXPR (if any) should be
disregarded in OpenMP construct, because it is going to be
remapped during OpenMP lowering. SHARED is true if DECL
@@ -185,6 +185,15 @@ extern tree cxx_simulate_record_decl (location_t, const char *,
#define LANG_HOOKS_OMP_CLAUSE_DTOR cxx_omp_clause_dtor
#undef LANG_HOOKS_OMP_FINISH_CLAUSE
#define LANG_HOOKS_OMP_FINISH_CLAUSE cxx_omp_finish_clause
+#undef LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES
+#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES cxx_omp_finish_mapper_clauses
+#undef LANG_HOOKS_OMP_MAPPER_LOOKUP
+#define LANG_HOOKS_OMP_MAPPER_LOOKUP cxx_omp_mapper_lookup
+#undef LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE
+#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE \
+ cxx_omp_extract_mapper_directive
+#undef LANG_HOOKS_OMP_MAP_ARRAY_SECTION
+#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION cxx_omp_map_array_section
#undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE
#define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference
#undef LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR
@@ -2860,7 +2860,10 @@ struct GTY(()) lang_decl_base {
/* VAR_DECL or FUNCTION_DECL has keyed decls. */
unsigned module_keyed_decls_p : 1;
- /* 12 spare bits. */
+ /* VAR_DECL being used to represent an OpenMP declared mapper. */
+ unsigned omp_declare_mapper_p : 1;
+
+ /* 10 spare bits. */
};
/* True for DECL codes which have template info and access. */
@@ -4334,6 +4337,11 @@ get_vec_init_expr (tree t)
#define DECL_OMP_DECLARE_REDUCTION_P(NODE) \
(LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_reduction_p)
+/* Nonzero if NODE is an artificial FUNCTION_DECL for
+ #pragma omp declare mapper. */
+#define DECL_OMP_DECLARE_MAPPER_P(NODE) \
+ (DECL_LANG_SPECIFIC (VAR_DECL_CHECK (NODE))->u.base.omp_declare_mapper_p)
+
/* Nonzero if DECL has been declared threadprivate by
#pragma omp threadprivate. */
#define CP_DECL_THREADPRIVATE_P(DECL) \
@@ -7708,10 +7716,13 @@ extern tree finish_qualified_id_expr (tree, tree, bool, bool,
extern void simplify_aggr_init_expr (tree *);
extern void finalize_nrv (tree *, tree, tree);
extern tree omp_reduction_id (enum tree_code, tree, tree);
+extern tree omp_mapper_id (tree, tree);
extern tree cp_remove_omp_priv_cleanup_stmt (tree *, int *, void *);
extern bool cp_check_omp_declare_reduction (tree);
+extern bool cp_check_omp_declare_mapper (tree);
extern void finish_omp_declare_simd_methods (tree);
extern tree finish_omp_clauses (tree, enum c_omp_region_type);
+extern tree omp_instantiate_mappers (tree);
extern tree push_omp_privatization_clauses (bool);
extern void pop_omp_privatization_clauses (tree);
extern void save_omp_privatization_clauses (vec<tree> &);
@@ -8271,6 +8282,10 @@ extern tree cxx_omp_clause_copy_ctor (tree, tree, tree);
extern tree cxx_omp_clause_assign_op (tree, tree, tree);
extern tree cxx_omp_clause_dtor (tree, tree);
extern void cxx_omp_finish_clause (tree, gimple_seq *, bool);
+extern tree cxx_omp_finish_mapper_clauses (tree);
+extern tree cxx_omp_mapper_lookup (tree, tree);
+extern tree cxx_omp_extract_mapper_directive (tree);
+extern tree cxx_omp_map_array_section (location_t, tree);
extern bool cxx_omp_privatize_by_reference (const_tree);
extern bool cxx_omp_disregard_value_expr (tree, bool);
extern void cp_fold_function (tree);
@@ -7311,6 +7311,12 @@ check_initializer (tree decl, tree init, int flags, vec<tree, va_gc> **cleanups)
}
else if (!init && DECL_REALLY_EXTERN (decl))
;
+ else if (flag_openmp
+ && VAR_P (decl)
+ && DECL_LANG_SPECIFIC (decl)
+ && DECL_OMP_DECLARE_MAPPER_P (decl)
+ && TREE_CODE (init) == OMP_DECLARE_MAPPER)
+ return NULL_TREE;
else if (init || type_build_ctor_call (type)
|| TYPE_REF_P (type))
{
@@ -8454,14 +8460,23 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p,
varpool_node::get_create (decl);
}
+ if (flag_openmp
+ && VAR_P (decl)
+ && DECL_LANG_SPECIFIC (decl)
+ && DECL_OMP_DECLARE_MAPPER_P (decl)
+ && init)
+ {
+ gcc_assert (TREE_CODE (init) == OMP_DECLARE_MAPPER);
+ DECL_INITIAL (decl) = init;
+ }
/* Convert the initializer to the type of DECL, if we have not
already initialized DECL. */
- if (!DECL_INITIALIZED_P (decl)
- /* If !DECL_EXTERNAL then DECL is being defined. In the
- case of a static data member initialized inside the
- class-specifier, there can be an initializer even if DECL
- is *not* defined. */
- && (!DECL_EXTERNAL (decl) || init))
+ else if (!DECL_INITIALIZED_P (decl)
+ /* If !DECL_EXTERNAL then DECL is being defined. In the
+ case of a static data member initialized inside the
+ class-specifier, there can be an initializer even if DECL
+ is *not* defined. */
+ && (!DECL_EXTERNAL (decl) || init))
{
cleanups = make_tree_vector ();
init = check_initializer (decl, init, flags, &cleanups);
@@ -5772,10 +5772,15 @@ mark_used (tree decl, tsubst_flags_t complain /* = tf_warning_or_error */)
/* If DECL has a deduced return type, we need to instantiate it now to
find out its type. For OpenMP user defined reductions, we need them
- instantiated for reduction clauses which inline them by hand directly. */
+ instantiated for reduction clauses which inline them by hand directly.
+ OpenMP declared mappers are used implicitly so must be instantiated
+ before they can be detected. */
if (undeduced_auto_decl (decl)
|| (TREE_CODE (decl) == FUNCTION_DECL
- && DECL_OMP_DECLARE_REDUCTION_P (decl)))
+ && DECL_OMP_DECLARE_REDUCTION_P (decl))
+ || (TREE_CODE (decl) == VAR_DECL
+ && DECL_LANG_SPECIFIC (decl)
+ && DECL_OMP_DECLARE_MAPPER_P (decl)))
maybe_instantiate_decl (decl);
if (processing_template_decl || in_template_function ())
@@ -1136,12 +1136,37 @@ dump_global_iord (cxx_pretty_printer *pp, tree t)
pp_printf (pp, p, DECL_SOURCE_FILE (t));
}
+/* Write a representation of OpenMP "declare mapper" T to PP in a manner
+ suitable for error messages. */
+
+static void
+dump_omp_declare_mapper (cxx_pretty_printer *pp, tree t, int flags)
+{
+ pp_string (pp, "#pragma omp declare mapper");
+ if (t == NULL_TREE || t == error_mark_node)
+ return;
+ pp_space (pp);
+ pp_cxx_left_paren (pp);
+ if (OMP_DECLARE_MAPPER_ID (t))
+ {
+ pp_cxx_tree_identifier (pp, OMP_DECLARE_MAPPER_ID (t));
+ pp_colon (pp);
+ }
+ dump_type (pp, TREE_TYPE (t), flags);
+ pp_cxx_right_paren (pp);
+}
+
static void
dump_simple_decl (cxx_pretty_printer *pp, tree t, tree type, int flags)
{
if (TREE_CODE (t) == VAR_DECL && DECL_NTTP_OBJECT_P (t))
return dump_expr (pp, DECL_INITIAL (t), flags);
+ if (TREE_CODE (t) == VAR_DECL
+ && DECL_LANG_SPECIFIC (t)
+ && DECL_OMP_DECLARE_MAPPER_P (t))
+ return dump_omp_declare_mapper (pp, DECL_INITIAL (t), flags);
+
if (flags & TFF_DECL_SPECIFIERS)
{
if (concept_definition_p (t))
@@ -39987,13 +39987,12 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc)
map ( [map-type-modifier[,] ...] map-kind: variable-list )
map-type-modifier:
- always | close */
+ always | close | mapper ( mapper-name ) */
static tree
-cp_parser_omp_clause_map (cp_parser *parser, tree list)
+cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind)
{
tree nlist, c;
- enum gomp_map_kind kind = GOMP_MAP_TOFROM;
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return list;
@@ -40011,11 +40010,16 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
pos++;
+ else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type
+ == CPP_OPEN_PAREN)
+ pos = cp_parser_skip_balanced_tokens (parser, pos + 1);
pos++;
}
bool always_modifier = false;
bool close_modifier = false;
+ bool mapper_modifier = false;
+ tree mapper_name = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -40038,6 +40042,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
return list;
}
always_modifier = true;
+ cp_lexer_consume_token (parser->lexer);
}
else if (strcmp ("close", p) == 0)
{
@@ -40051,20 +40056,83 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
return list;
}
close_modifier = true;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else if (strcmp ("mapper", p) == 0)
+ {
+ cp_lexer_consume_token (parser->lexer);
+
+ matching_parens parens;
+ if (parens.require_open (parser))
+ {
+ if (mapper_modifier)
+ {
+ cp_parser_error (parser, "too many %<mapper%> modifiers");
+ /* Assume it's a well-formed mapper modifier, even if it
+ seems to be in the wrong place. */
+ cp_lexer_consume_token (parser->lexer);
+ parens.require_close (parser);
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/
+ true);
+ return list;
+ }
+
+ tok = cp_lexer_peek_token (parser->lexer);
+ switch (tok->type)
+ {
+ case CPP_NAME:
+ {
+ cp_expr e = cp_parser_identifier (parser);
+ if (e != error_mark_node)
+ mapper_name = e;
+ else
+ goto err;
+ }
+ break;
+
+ case CPP_KEYWORD:
+ if (tok->keyword == RID_DEFAULT)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ break;
+ }
+ /* Fallthrough. */
+
+ default:
+ err:
+ cp_parser_error (parser,
+ "expected identifier or %<default%>");
+ return list;
+ }
+
+ if (!parens.require_close (parser))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/
+ true);
+ return list;
+ }
+
+ mapper_modifier = true;
+ pos += 3;
+ }
}
else
{
cp_parser_error (parser, "%<#pragma omp target%> with "
- "modifier other than %<always%> or "
- "%<close%> on %<map%> clause");
+ "modifier other than %<always%>, %<close%> "
+ "or %<mapper%> on %<map%> clause");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
/*or_comma=*/false,
/*consume_paren=*/true);
return list;
}
-
- cp_lexer_consume_token (parser->lexer);
}
if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
@@ -40110,8 +40178,30 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
NULL, true);
finish_scope ();
+ tree last_new = NULL_TREE;
+
for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ last_new = c;
+ }
+
+ if (mapper_name)
+ {
+ tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME);
+ OMP_CLAUSE_DECL (name) = mapper_name;
+ OMP_CLAUSE_CHAIN (name) = nlist;
+ nlist = name;
+
+ gcc_assert (last_new);
+
+ name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME);
+ OMP_CLAUSE_DECL (name) = null_pointer_node;
+ OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new);
+ OMP_CLAUSE_CHAIN (last_new) = name;
+ }
return nlist;
}
@@ -40927,7 +41017,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
c_name = "detach";
break;
case PRAGMA_OMP_CLAUSE_MAP:
- clauses = cp_parser_omp_clause_map (parser, clauses);
+ clauses = cp_parser_omp_clause_map (parser, clauses, GOMP_MAP_TOFROM);
c_name = "map";
break;
case PRAGMA_OMP_CLAUSE_DEVICE:
@@ -45095,6 +45185,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = nc;
}
+ if (!processing_template_decl)
+ clauses = c_omp_instantiate_mappers (clauses);
clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
c_omp_adjust_map_clauses (clauses, true);
@@ -47359,6 +47451,172 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok,
obstack_free (&declarator_obstack, p);
}
+/* OpenMP 5.0
+ #pragma omp declare mapper([mapper-identifier:]type var) \
+ [clause[[,] clause] ... ] new-line */
+
+static void
+cp_parser_omp_declare_mapper (cp_parser *parser, cp_token *pragma_tok,
+ enum pragma_context)
+{
+ cp_token *token = NULL;
+ tree type = NULL_TREE, vardecl = NULL_TREE, block = NULL_TREE;
+ bool block_scope = false;
+ /* Don't create location wrapper nodes within "declare mapper"
+ directives. */
+ auto_suppress_location_wrappers sentinel;
+ tree mapper_name = NULL_TREE;
+ tree mapper_id, id, placeholder, mapper, maplist = NULL_TREE;
+
+ if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+ goto fail;
+
+ if (current_function_decl)
+ block_scope = true;
+
+ token = cp_lexer_peek_token (parser->lexer);
+
+ if (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ {
+ switch (token->type)
+ {
+ case CPP_NAME:
+ {
+ cp_expr e = cp_parser_identifier (parser);
+ if (e != error_mark_node)
+ mapper_name = e;
+ else
+ goto fail;
+ }
+ break;
+
+ case CPP_KEYWORD:
+ if (token->keyword == RID_DEFAULT)
+ {
+ mapper_name = NULL_TREE;
+ cp_lexer_consume_token (parser->lexer);
+ break;
+ }
+ /* Fallthrough. */
+
+ default:
+ cp_parser_error (parser, "expected identifier or %<default%>");
+ }
+
+ if (!cp_parser_require (parser, CPP_COLON, RT_COLON))
+ goto fail;
+ }
+
+ {
+ const char *saved_message = parser->type_definition_forbidden_message;
+ parser->type_definition_forbidden_message
+ = G_("types may not be defined within %<declare mapper%>");
+ type_id_in_expr_sentinel s (parser);
+ type = cp_parser_type_id (parser);
+ parser->type_definition_forbidden_message = saved_message;
+ }
+
+ if (dependent_type_p (type))
+ mapper_id = omp_mapper_id (mapper_name, NULL_TREE);
+ else
+ mapper_id = omp_mapper_id (mapper_name, type);
+
+ vardecl = build_lang_decl (VAR_DECL, mapper_id, type);
+ DECL_ARTIFICIAL (vardecl) = 1;
+ TREE_STATIC (vardecl) = 1;
+ TREE_PUBLIC (vardecl) = 0;
+ DECL_EXTERNAL (vardecl) = 0;
+ DECL_DECLARED_CONSTEXPR_P (vardecl) = 1;
+ DECL_INITIALIZED_BY_CONSTANT_EXPRESSION_P (vardecl) = 1;
+ DECL_OMP_DECLARE_MAPPER_P (vardecl) = 1;
+
+ keep_next_level (true);
+ block = begin_omp_structured_block ();
+
+ if (block_scope)
+ DECL_CONTEXT (vardecl) = current_function_decl;
+ else if (current_class_type)
+ DECL_CONTEXT (vardecl) = current_class_type;
+ else
+ DECL_CONTEXT (vardecl) = current_namespace;
+
+ if (processing_template_decl)
+ vardecl = push_template_decl (vardecl);
+
+ id = cp_parser_declarator_id (parser, false);
+
+ if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ {
+ finish_omp_structured_block (block);
+ goto fail;
+ }
+
+ placeholder = build_lang_decl (VAR_DECL, id, type);
+ DECL_CONTEXT (placeholder) = DECL_CONTEXT (vardecl);
+ if (processing_template_decl)
+ placeholder = push_template_decl (placeholder);
+ pushdecl (placeholder);
+ cp_finish_decl (placeholder, NULL_TREE, 0, NULL_TREE, 0);
+ DECL_ARTIFICIAL (placeholder) = 1;
+ TREE_USED (placeholder) = 1;
+
+ while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+ {
+ pragma_omp_clause c_kind = cp_parser_omp_clause_name (parser);
+ if (c_kind != PRAGMA_OMP_CLAUSE_MAP)
+ {
+ if (c_kind != PRAGMA_OMP_CLAUSE_NONE)
+ cp_parser_error (parser, "unexpected clause");
+ finish_omp_structured_block (block);
+ goto fail;
+ }
+ maplist = cp_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET);
+ if (maplist == NULL_TREE)
+ break;
+ }
+
+ if (maplist == NULL_TREE)
+ {
+ cp_parser_error (parser, "missing %<map%> clause");
+ finish_omp_structured_block (block);
+ goto fail;
+ }
+
+ mapper = make_node (OMP_DECLARE_MAPPER);
+ TREE_TYPE (mapper) = type;
+ OMP_DECLARE_MAPPER_ID (mapper) = mapper_name;
+ OMP_DECLARE_MAPPER_DECL (mapper) = placeholder;
+ OMP_DECLARE_MAPPER_CLAUSES (mapper) = maplist;
+
+ finish_omp_structured_block (block);
+
+ DECL_INITIAL (vardecl) = mapper;
+
+ if (current_class_type)
+ {
+ if (processing_template_decl)
+ {
+ retrofit_lang_decl (vardecl);
+ SET_DECL_VAR_DECLARED_INLINE_P (vardecl);
+ }
+ finish_static_data_member_decl (vardecl, mapper,
+ /*init_const_expr_p=*/true, NULL_TREE, 0);
+ finish_member_declaration (vardecl);
+ }
+ else if (processing_template_decl && block_scope)
+ add_decl_expr (vardecl);
+ else
+ pushdecl (vardecl);
+
+ cp_check_omp_declare_mapper (vardecl);
+
+ cp_parser_require_pragma_eol (parser, pragma_tok);
+ return;
+
+fail:
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+}
+
/* OpenMP 4.0
#pragma omp declare simd declare-simd-clauses[optseq] new-line
#pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -47399,6 +47657,12 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok,
context);
return false;
}
+ if (strcmp (p, "mapper") == 0)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ cp_parser_omp_declare_mapper (parser, pragma_tok, context);
+ return false;
+ }
if (!flag_openmp) /* flag_openmp_simd */
{
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
@@ -47412,7 +47676,7 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok,
}
}
cp_parser_error (parser, "expected %<simd%>, %<reduction%>, "
- "%<target%> or %<variant%>");
+ "%<target%>, %<mapper%> or %<variant%>");
cp_parser_require_pragma_eol (parser, pragma_tok);
return false;
}
@@ -15220,6 +15220,13 @@ tsubst_decl (tree t, tree args, tsubst_flags_t complain)
TYPE_ALIGN (TREE_TYPE (t)));
}
+ if (flag_openmp
+ && VAR_P (t)
+ && DECL_LANG_SPECIFIC (t)
+ && DECL_OMP_DECLARE_MAPPER_P (t)
+ && strchr (IDENTIFIER_POINTER (DECL_NAME (t)), '~') == NULL)
+ DECL_NAME (r) = omp_mapper_id (DECL_NAME (t), TREE_TYPE (r));
+
layout_decl (r, 0);
}
break;
@@ -18095,6 +18102,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
new_clauses = nreverse (new_clauses);
if (ort != C_ORT_OMP_DECLARE_SIMD)
{
+ if (ort == C_ORT_OMP_TARGET)
+ new_clauses = c_omp_instantiate_mappers (new_clauses);
new_clauses = finish_omp_clauses (new_clauses, ort);
if (linear_no_step)
for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
@@ -19639,6 +19648,22 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
}
break;
+ case OMP_DECLARE_MAPPER:
+ {
+ t = copy_node (t);
+
+ tree decl = OMP_DECLARE_MAPPER_DECL (t);
+ decl = tsubst (decl, args, complain, in_decl);
+ tree type = tsubst (TREE_TYPE (t), args, complain, in_decl);
+ tree clauses = OMP_DECLARE_MAPPER_CLAUSES (t);
+ clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD, args,
+ complain, in_decl);
+ TREE_TYPE (t) = type;
+ OMP_DECLARE_MAPPER_DECL (t) = decl;
+ OMP_DECLARE_MAPPER_CLAUSES (t) = clauses;
+ RETURN (t);
+ }
+
case TRANSACTION_EXPR:
{
int flags = 0;
@@ -26920,7 +26945,9 @@ instantiate_decl (tree d, bool defer_ok, bool expl_inst_class_mem_p)
|| (external_p && VAR_P (d))
/* Handle here a deleted function too, avoid generating
its body (c++/61080). */
- || deleted_p)
+ || deleted_p
+ /* We need the initializer for an OpenMP declare mapper. */
+ || (VAR_P (d) && DECL_LANG_SPECIFIC (d) && DECL_OMP_DECLARE_MAPPER_P (d)))
{
/* The definition of the static data member is now required so
we must substitute the initializer. */
@@ -45,6 +45,7 @@ along with GCC; see the file COPYING3. If not see
#include "gomp-constants.h"
#include "predict.h"
#include "memmodel.h"
+#include "gimplify.h"
/* There routines provide a modular interface to perform many parsing
operations. They may therefore be used during actual parsing, or
@@ -5944,6 +5945,98 @@ omp_reduction_lookup (location_t loc, tree id, tree type, tree *baselinkp,
return id;
}
+/* Return identifier to look up for omp declare mapper. */
+
+tree
+omp_mapper_id (tree mapper_id, tree type)
+{
+ const char *p = NULL;
+ const char *m = NULL;
+
+ if (mapper_id == NULL_TREE)
+ p = "";
+ else if (TREE_CODE (mapper_id) == IDENTIFIER_NODE)
+ p = IDENTIFIER_POINTER (mapper_id);
+ else
+ return error_mark_node;
+
+ if (type != NULL_TREE)
+ m = mangle_type_string (TYPE_MAIN_VARIANT (type));
+
+ const char prefix[] = "omp declare mapper ";
+ size_t lenp = sizeof (prefix);
+ if (strncmp (p, prefix, lenp - 1) == 0)
+ lenp = 1;
+ size_t len = strlen (p);
+ size_t lenm = m ? strlen (m) + 1 : 0;
+ char *name = XALLOCAVEC (char, lenp + len + lenm);
+ memcpy (name, prefix, lenp - 1);
+ memcpy (name + lenp - 1, p, len + 1);
+ if (m)
+ {
+ name[lenp + len - 1] = '~';
+ memcpy (name + lenp + len, m, lenm);
+ }
+ return get_identifier (name);
+}
+
+tree
+cxx_omp_mapper_lookup (tree id, tree type)
+{
+ if (TREE_CODE (type) != RECORD_TYPE
+ && TREE_CODE (type) != UNION_TYPE)
+ return NULL_TREE;
+ id = omp_mapper_id (id, type);
+ return lookup_name (id);
+}
+
+tree
+cxx_omp_extract_mapper_directive (tree vardecl)
+{
+ gcc_assert (TREE_CODE (vardecl) == VAR_DECL);
+
+ /* Instantiate the decl if we haven't already. */
+ mark_used (vardecl);
+ tree body = DECL_INITIAL (vardecl);
+
+ if (TREE_CODE (body) == STATEMENT_LIST)
+ {
+ tree_stmt_iterator tsi = tsi_start (body);
+ gcc_assert (TREE_CODE (tsi_stmt (tsi)) == DECL_EXPR);
+ tsi_next (&tsi);
+ body = tsi_stmt (tsi);
+ }
+
+ gcc_assert (TREE_CODE (body) == OMP_DECLARE_MAPPER);
+
+ return body;
+}
+
+/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but
+ nothing more complicated. */
+
+tree
+cxx_omp_map_array_section (location_t loc, tree t)
+{
+ tree low = TREE_OPERAND (t, 1);
+ tree len = TREE_OPERAND (t, 2);
+
+ if (len && integer_onep (len))
+ {
+ t = TREE_OPERAND (t, 0);
+
+ if (!low)
+ low = integer_zero_node;
+
+ if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
+ t = convert_from_reference (t);
+
+ t = build_array_ref (loc, t, low);
+ }
+
+ return t;
+}
+
/* Helper function for cp_parser_omp_declare_reduction_exprs
and tsubst_omp_udr.
Remove CLEANUP_STMT for data (omp_priv variable).
@@ -6425,6 +6518,29 @@ finish_omp_reduction_clause (tree c, bool *need_default_ctor, bool *need_dtor)
return false;
}
+/* Check an instance of an "omp declare mapper" function. */
+
+bool
+cp_check_omp_declare_mapper (tree udm)
+{
+ tree type = TREE_TYPE (udm);
+ location_t loc = DECL_SOURCE_LOCATION (udm);
+
+ if (type == error_mark_node)
+ return false;
+
+ if (!processing_template_decl
+ && TREE_CODE (type) != RECORD_TYPE
+ && TREE_CODE (type) != UNION_TYPE)
+ {
+ error_at (loc, "%qT is not a struct, union or class type in "
+ "%<#pragma omp declare mapper%>", type);
+ return false;
+ }
+
+ return true;
+}
+
/* Called from finish_struct_1. linear(this) or linear(this:step)
clauses might not be finalized yet because the class has been incomplete
when parsing #pragma omp declare simd methods. Fix those up now. */
@@ -7965,6 +8081,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_MAP:
if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
goto move_implicit;
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ remove = true;
+ break;
+ }
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
@@ -9414,6 +9536,8 @@ struct omp_target_walk_data
/* Local variables declared inside a BIND_EXPR, used to filter out such
variables when recording lambda_objects_accessed. */
hash_set<tree> local_decls;
+
+ omp_mapper_list<tree> *mappers;
};
/* Helper function of finish_omp_target_clauses, called via
@@ -9427,6 +9551,7 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr;
tree current_object = data->current_object;
tree current_closure = data->current_closure;
+ omp_mapper_list<tree> *mlist = data->mappers;
/* References inside of these expression codes shouldn't incur any
form of mapping, so return early. */
@@ -9440,6 +9565,27 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
if (TREE_CODE (t) == OMP_CLAUSE)
return NULL_TREE;
+ if (!processing_template_decl)
+ {
+ tree aggr_type = NULL_TREE;
+
+ if (TREE_CODE (t) == COMPONENT_REF
+ && AGGREGATE_TYPE_P (TREE_TYPE (TREE_OPERAND (t, 0))))
+ aggr_type = TREE_TYPE (TREE_OPERAND (t, 0));
+ else if ((TREE_CODE (t) == VAR_DECL
+ || TREE_CODE (t) == PARM_DECL
+ || TREE_CODE (t) == RESULT_DECL)
+ && AGGREGATE_TYPE_P (TREE_TYPE (t)))
+ aggr_type = TREE_TYPE (t);
+
+ if (aggr_type)
+ {
+ tree mapper_fn = cxx_omp_mapper_lookup (NULL_TREE, aggr_type);
+ if (mapper_fn)
+ mlist->add_mapper (NULL_TREE, aggr_type, mapper_fn);
+ }
+ }
+
if (current_object)
{
tree this_expr = TREE_OPERAND (current_object, 0);
@@ -9542,10 +9688,48 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
else
data.current_closure = NULL_TREE;
- cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data);
-
auto_vec<tree, 16> new_clauses;
+ if (!processing_template_decl)
+ {
+ hash_set<omp_name_type<tree> > seen_types;
+ auto_vec<tree> mapper_fns;
+ omp_mapper_list<tree> mlist (&seen_types, &mapper_fns);
+ data.mappers = &mlist;
+
+ cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r,
+ &data);
+
+ unsigned int i;
+ tree mapper_fn;
+ FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
+ c_omp_find_nested_mappers (&mlist, mapper_fn);
+
+ FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
+ {
+ tree mapper = cxx_omp_extract_mapper_directive (mapper_fn);
+ if (mapper == error_mark_node)
+ continue;
+ tree mapper_name = OMP_DECLARE_MAPPER_ID (mapper);
+ tree decl = OMP_DECLARE_MAPPER_DECL (mapper);
+ if (BASELINK_P (mapper_fn))
+ mapper_fn = BASELINK_FUNCTIONS (mapper_fn);
+
+ tree c = build_omp_clause (loc, OMP_CLAUSE__MAPPER_BINDING_);
+ OMP_CLAUSE__MAPPER_BINDING__ID (c) = mapper_name;
+ OMP_CLAUSE__MAPPER_BINDING__DECL (c) = decl;
+ OMP_CLAUSE__MAPPER_BINDING__MAPPER (c) = mapper_fn;
+
+ new_clauses.safe_push (c);
+ }
+ }
+ else
+ {
+ data.mappers = NULL;
+ cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r,
+ &data);
+ }
+
tree omp_target_this_expr = NULL_TREE;
tree *explicit_this_deref_map = NULL;
if (data.this_expr_accessed)
@@ -27,6 +27,9 @@ along with GCC; see the file COPYING3. If not see
#include "match.h"
#include "parse.h"
#include "tree-core.h"
+#include "tree.h"
+#include "fold-const.h"
+#include "tree-hash-traits.h"
#include "omp-general.h"
/* Current statement label. Zero means no statement label. Because new_st
@@ -219,6 +219,7 @@ struct gimplify_omp_ctx
{
struct gimplify_omp_ctx *outer_context;
splay_tree variables;
+ hash_map<omp_name_type<tree>, tree> *implicit_mappers;
hash_set<tree> *privatized_types;
tree clauses;
/* Iteration variables in an OMP_FOR. */
@@ -452,6 +453,7 @@ new_omp_context (enum omp_region_type region_type)
c = XCNEW (struct gimplify_omp_ctx);
c->outer_context = gimplify_omp_ctxp;
c->variables = splay_tree_new (splay_tree_compare_decl_uid, 0, 0);
+ c->implicit_mappers = new hash_map<omp_name_type<tree>, tree>;
c->privatized_types = new hash_set<tree>;
c->location = input_location;
c->region_type = region_type;
@@ -475,6 +477,7 @@ delete_omp_context (struct gimplify_omp_ctx *c)
{
splay_tree_delete (c->variables);
delete c->privatized_types;
+ delete c->implicit_mappers;
c->loop_iter_var.release ();
XDELETE (c);
}
@@ -11500,6 +11503,227 @@ error_out:
return success;
}
+struct instantiate_mapper_info
+{
+ tree *mapper_clauses_p;
+ struct gimplify_omp_ctx *omp_ctx;
+ gimple_seq *pre_p;
+};
+
+/* Helper function for omp_instantiate_mapper. */
+
+static tree
+remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
+{
+ copy_body_data *id = (copy_body_data *) data;
+
+ if (DECL_P (*tp))
+ {
+ tree replacement = remap_decl (*tp, id);
+ if (*tp != replacement)
+ {
+ *tp = unshare_expr (replacement);
+ *walk_subtrees = 0;
+ }
+ }
+
+ return NULL_TREE;
+}
+
+/* A copy_decl implementation (for use with tree-inline.cc functions) that
+ only transform decls or SSA names that are part of a map we already
+ prepared. */
+
+static tree
+omp_mapper_copy_decl (tree var, copy_body_data *cb)
+{
+ tree *repl = cb->decl_map->get (var);
+
+ if (repl)
+ return *repl;
+
+ return var;
+}
+
+static tree *
+omp_instantiate_mapper (gimple_seq *pre_p,
+ hash_map<omp_name_type<tree>, tree> *implicit_mappers,
+ tree mapperfn, tree expr, enum gomp_map_kind outer_kind,
+ tree *mapper_clauses_p)
+{
+ tree mapper_name = NULL_TREE;
+ tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapperfn);
+ gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+ tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+ tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
+
+ /* The "extraction map" is used to map the mapper variable in the "declare
+ mapper" directive, and also any temporary variables that have been created
+ as part of expanding the mapper function's body (which are expanded as a
+ "bind" expression in the pre_p sequence). */
+ hash_map<tree, tree> extraction_map;
+
+ extraction_map.put (dummy_var, expr);
+ extraction_map.put (expr, expr);
+
+ /* This copy_body_data is only used to remap the decls in the
+ OMP_DECLARE_MAPPER tree node expansion itself. All relevant decls should
+ already be in the current function. */
+ copy_body_data id;
+ memset (&id, 0, sizeof (id));
+ id.src_fn = current_function_decl;
+ id.dst_fn = current_function_decl;
+ id.src_cfun = cfun;
+ id.decl_map = &extraction_map;
+ id.copy_decl = omp_mapper_copy_decl;
+ id.transform_call_graph_edges = CB_CGE_DUPLICATE; // ???
+ id.transform_new_cfg = true; // ???
+
+ for (; clause; clause = OMP_CLAUSE_CHAIN (clause))
+ {
+ enum gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (clause);
+ tree *nested_mapper_p = NULL;
+
+ if (map_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+ {
+ mapper_name = OMP_CLAUSE_DECL (clause);
+ continue;
+ }
+ else if (map_kind == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ mapper_name = NULL_TREE;
+ continue;
+ }
+
+ tree decl = OMP_CLAUSE_DECL (clause);
+ tree unshared, type;
+ bool nonunit_array_with_mapper = false;
+
+ if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (clause);
+ tree tmp = lang_hooks.decls.omp_map_array_section (loc, decl);
+ if (tmp == decl)
+ {
+ unshared = unshare_expr (clause);
+ nonunit_array_with_mapper = true;
+ type = TREE_TYPE (TREE_TYPE (decl));
+ }
+ else
+ {
+ unshared = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_CODE (clause));
+ OMP_CLAUSE_DECL (unshared) = tmp;
+ OMP_CLAUSE_SIZE (unshared)
+ = DECL_P (tmp) ? DECL_SIZE_UNIT (tmp)
+ : TYPE_SIZE_UNIT (TREE_TYPE (tmp));
+ type = TREE_TYPE (tmp);
+ }
+ }
+ else
+ {
+ unshared = unshare_expr (clause);
+ type = TREE_TYPE (decl);
+ }
+
+ walk_tree (&unshared, remap_mapper_decl_1, &id, NULL);
+
+ if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)
+ OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
+
+ decl = OMP_CLAUSE_DECL (unshared);
+ type = TYPE_MAIN_VARIANT (type);
+
+ nested_mapper_p = implicit_mappers->get ({ mapper_name, type });
+
+ if (nested_mapper_p && *nested_mapper_p != mapperfn)
+ {
+ if (nonunit_array_with_mapper)
+ {
+ sorry ("user-defined mapper with non-unit length array section");
+ continue;
+ }
+
+ if (map_kind == GOMP_MAP_UNSET)
+ map_kind = outer_kind;
+
+ mapper_clauses_p
+ = omp_instantiate_mapper (pre_p, implicit_mappers,
+ *nested_mapper_p, decl, map_kind,
+ mapper_clauses_p);
+ continue;
+ }
+
+ *mapper_clauses_p = unshared;
+ mapper_clauses_p = &OMP_CLAUSE_CHAIN (unshared);
+ }
+
+ return mapper_clauses_p;
+}
+
+static int
+omp_instantiate_implicit_mappers (splay_tree_node n, void *data)
+{
+ tree decl = (tree) n->key;
+ instantiate_mapper_info *im_info = (instantiate_mapper_info *) data;
+ gimplify_omp_ctx *ctx = im_info->omp_ctx;
+ tree *mapper_p = NULL;
+ tree type = TREE_TYPE (decl);
+ bool ref_p = false;
+ unsigned flags = n->value;
+
+ if (flags & (GOVD_EXPLICIT | GOVD_LOCAL))
+ return 0;
+ if ((flags & GOVD_SEEN) == 0)
+ return 0;
+
+ if (TREE_CODE (type) == REFERENCE_TYPE)
+ {
+ ref_p = true;
+ type = TREE_TYPE (type);
+ }
+
+ type = TYPE_MAIN_VARIANT (type);
+
+ if (DECL_P (decl) && type && AGGREGATE_TYPE_P (type))
+ {
+ gcc_assert (ctx);
+ mapper_p = ctx->implicit_mappers->get ({ NULL_TREE, type });
+ }
+
+ /* If we already have clauses pertaining to a struct variable, then we don't
+ want to implicitly invoke a user-defined mapper. */
+ bool handled_p = false;
+ splay_tree_node an
+ = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+ if (an)
+ {
+ unsigned flags = an->value;
+ if ((flags & GOVD_EXPLICIT) != 0
+ && AGGREGATE_TYPE_P (TREE_TYPE (decl)))
+ handled_p = true;
+ }
+
+ if (mapper_p && !handled_p)
+ {
+ /* If we have a reference, map the pointed-to object rather than the
+ reference itself. */
+ if (ref_p)
+ decl = build_fold_indirect_ref (decl);
+
+ im_info->mapper_clauses_p
+ = omp_instantiate_mapper (im_info->pre_p, ctx->implicit_mappers,
+ *mapper_p, decl, GOMP_MAP_TOFROM,
+ im_info->mapper_clauses_p);
+ /* Make sure we don't map the same variable implicitly in
+ gimplify_adjust_omp_clauses_1 also. */
+ n->value |= GOVD_EXPLICIT;
+ }
+
+ return 0;
+}
+
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
and previous omp contexts. */
@@ -11511,7 +11735,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
using namespace omp_addr_tokenizer;
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
- tree *prev_list_p = NULL, *orig_list_p = list_p;
+ tree *orig_list_p = list_p;
int handled_depend_iterators = -1;
int nowait = -1;
@@ -11542,73 +11766,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
- if (code == OMP_TARGET
- || code == OMP_TARGET_DATA
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA)
- {
- vec<omp_mapping_group> *groups;
- groups = omp_gather_mapping_groups (list_p);
- if (groups)
- {
- hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
- grpmap = omp_index_mapping_groups (groups);
+ vec<omp_mapping_group> *groups = omp_gather_mapping_groups (list_p);
+ hash_map<tree_operand_hash, omp_mapping_group *> *grpmap = NULL;
+ unsigned grpnum = 0;
+ tree *grp_start_p = NULL, grp_end = NULL_TREE;
- omp_resolve_clause_dependencies (code, groups, grpmap);
- omp_build_struct_sibling_lists (code, region_type, groups, &grpmap,
- list_p);
-
- omp_mapping_group *outlist = NULL;
-
- /* Topological sorting may fail if we have duplicate nodes, which
- we should have detected and shown an error for already. Skip
- sorting in that case. */
- if (seen_error ())
- goto failure;
-
- delete grpmap;
- delete groups;
-
- /* Rebuild now we have struct sibling lists. */
- groups = omp_gather_mapping_groups (list_p);
- grpmap = omp_index_mapping_groups (groups);
-
- outlist = omp_tsort_mapping_groups (groups, grpmap);
- outlist = omp_segregate_mapping_groups (outlist);
- list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
-
- failure:
- delete grpmap;
- delete groups;
- }
- }
- else if (region_type & ORT_ACC)
- {
- vec<omp_mapping_group> *groups;
- groups = omp_gather_mapping_groups (list_p);
- if (groups)
- {
- hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
- grpmap = omp_index_mapping_groups (groups);
-
- oacc_resolve_clause_dependencies (groups, grpmap);
- omp_build_struct_sibling_lists (code, region_type, groups, &grpmap,
- list_p);
-
- delete groups;
- delete grpmap;
- }
- }
+ if (groups)
+ grpmap = omp_index_mapping_groups (groups);
while ((c = *list_p) != NULL)
{
bool remove = false;
bool notice_outer = true;
+ bool map_descriptor;
const char *check_non_private = NULL;
unsigned int flags;
tree decl;
auto_vec<omp_addr_token *, 10> addr_tokens;
+ if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end))
+ {
+ grp_start_p = NULL;
+ grp_end = NULL_TREE;
+ }
+
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_PRIVATE:
@@ -11913,80 +12094,36 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
goto do_add;
case OMP_CLAUSE_MAP:
+ if (!grp_start_p)
+ {
+ grp_start_p = list_p;
+ grp_end = (*groups)[grpnum].grp_end;
+ grpnum++;
+ }
decl = OMP_CLAUSE_DECL (c);
+ if (error_operand_p (decl))
+ {
+ remove = true;
+ break;
+ }
+
if (!omp_parse_expr (addr_tokens, decl))
{
remove = true;
break;
}
- if (error_operand_p (decl))
- remove = true;
- switch (code)
- {
- case OMP_TARGET:
- break;
- case OACC_DATA:
- if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
- break;
- goto check_firstprivate;
- case OACC_ENTER_DATA:
- case OACC_EXIT_DATA:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
- && addr_tokens[0]->type == ARRAY_BASE)
- remove = true;
- /* FALLTHRU */
- case OMP_TARGET_DATA:
- case OMP_TARGET_ENTER_DATA:
- case OMP_TARGET_EXIT_DATA:
- case OACC_HOST_DATA:
- check_firstprivate:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
- /* For target {,enter ,exit }data only the array slice is
- mapped, but not the pointer to it. */
- remove = true;
- break;
- default:
- break;
- }
- /* For Fortran, not only the pointer to the data is mapped but also
- the address of the pointer, the array descriptor etc.; for
- 'exit data' - and in particular for 'delete:' - having an 'alloc:'
- does not make sense. Likewise, for 'update' only transferring the
- data itself is needed as the rest has been handled in previous
- directives. However, for 'exit data', the array descriptor needs
- to be delete; hence, we turn the MAP_TO_PSET into a MAP_DELETE.
-
- NOTE: Generally, it is not safe to perform "enter data" operations
- on arrays where the data *or the descriptor* may go out of scope
- before a corresponding "exit data" operation -- and such a
- descriptor may be synthesized temporarily, e.g. to pass an
- explicit-shape array to a function expecting an assumed-shape
- argument. Performing "enter data" inside the called function
- would thus be problematic. */
if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
- OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*prev_list_p)
- == GOMP_MAP_DELETE
- ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
- else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE)
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
- remove = true;
- else if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
- && OMP_CLAUSE_CHAIN (c)
- && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
- == OMP_CLAUSE_MAP)
- && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ATTACH_DETACH)
- || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
- && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
- (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ && OMP_CLAUSE_CHAIN (c)
+ && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
+ && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ATTACH_DETACH)
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
+ && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
+ (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
if (remove)
@@ -12007,41 +12144,51 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
DECL_NAME (decl));
}
}
- if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
- OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
- : TYPE_SIZE_UNIT (TREE_TYPE (decl));
- if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
- NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+
+ map_descriptor = false;
+
+ /* This condition checks if we're mapping an array descriptor that
+ isn't inside a derived type -- these have special handling, and
+ are not handled as structs in omp_build_struct_sibling_lists.
+ See that function for further details. */
+ if (*grp_start_p != grp_end
+ && OMP_CLAUSE_CHAIN (*grp_start_p)
+ && OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
{
- remove = true;
- break;
- }
- else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
- {
- OMP_CLAUSE_SIZE (c)
- = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
- false);
- if ((region_type & ORT_TARGET) != 0)
- omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
- GOVD_FIRSTPRIVATE | GOVD_SEEN);
+ tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p);
+ if (OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET
+ && DECL_P (OMP_CLAUSE_DECL (grp_mid)))
+ map_descriptor = true;
}
- if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT_UNORD)
- && (addr_tokens[0]->type == STRUCTURE_BASE
- || addr_tokens[0]->type == ARRAY_BASE)
- && addr_tokens[0]->u.structure_base_kind == BASE_DECL)
+ /* Adding the decl for a struct access: we haven't created
+ GOMP_MAP_STRUCT nodes yet, so this statement needs to predict
+ whether they will be created in gimplify_adjust_omp_clauses. */
+ if (c == grp_end
+ && addr_tokens[0]->type == STRUCTURE_BASE
+ && addr_tokens[0]->u.structure_base_kind == BASE_DECL
+ && !map_descriptor)
{
gcc_assert (addr_tokens[1]->type == ACCESS_METHOD);
/* If we got to this struct via a chain of pointers, maybe we
want to map it implicitly instead. */
if (omp_access_chain_p (addr_tokens, 1))
break;
+ omp_mapping_group *wholestruct;
+ if (!(region_type & ORT_ACC)
+ && omp_mapped_by_containing_struct (grpmap,
+ OMP_CLAUSE_DECL (c),
+ &wholestruct))
+ break;
decl = addr_tokens[1]->expr;
+ if (splay_tree_lookup (ctx->variables, (splay_tree_key) decl))
+ break;
+ /* Standalone attach or detach clauses for a struct element
+ should not inhibit implicit mapping of the whole struct. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ break;
flags = GOVD_MAP | GOVD_EXPLICIT;
gcc_assert (addr_tokens[1]->u.access_kind != ACCESS_DIRECT
@@ -12049,14 +12196,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
goto do_add_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))
+ if (!DECL_P (decl))
{
tree d = decl, *pd;
if (TREE_CODE (d) == ARRAY_REF)
@@ -12079,56 +12219,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
pd = &TREE_OPERAND (decl, 0);
decl = TREE_OPERAND (decl, 0);
}
- /* An "attach/detach" operation on an update directive should
- behave as a GOMP_MAP_ALWAYS_POINTER. Beware that
- unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER
- depends on the previous mapping. */
- if (code == OACC_UPDATE
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ if (addr_tokens[0]->type == STRUCTURE_BASE
+ && addr_tokens[0]->u.structure_base_kind == BASE_DECL
+ && addr_tokens[1]->type == ACCESS_METHOD
+ && (addr_tokens[1]->u.access_kind == ACCESS_POINTER
+ || (addr_tokens[1]->u.access_kind
+ == ACCESS_POINTER_OFFSET))
+ && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)))
{
- if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c)))
- == ARRAY_TYPE)
- remove = true;
- else
- {
- gomp_map_kind k = ((code == OACC_EXIT_DATA
- || code == OMP_TARGET_EXIT_DATA)
- ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
- OMP_CLAUSE_SET_MAP_KIND (c, k);
- }
- }
-
- tree cref = decl;
-
- while (TREE_CODE (cref) == ARRAY_REF)
- cref = TREE_OPERAND (cref, 0);
-
- if (TREE_CODE (cref) == INDIRECT_REF)
- cref = TREE_OPERAND (cref, 0);
-
- if (TREE_CODE (cref) == COMPONENT_REF)
- {
- tree base = cref;
- while (base && !DECL_P (base))
- {
- tree innerbase = omp_get_base_pointer (base);
- if (!innerbase)
- break;
- base = innerbase;
- }
- if (base
- && DECL_P (base)
- && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
- && POINTER_TYPE_P (TREE_TYPE (base)))
- {
- splay_tree_node n
- = splay_tree_lookup (ctx->variables,
- (splay_tree_key) base);
- n->value |= GOVD_SEEN;
- }
+ tree base = addr_tokens[1]->expr;
+ splay_tree_node n
+ = splay_tree_lookup (ctx->variables,
+ (splay_tree_key) base);
+ n->value |= GOVD_SEEN;
}
if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
@@ -12239,27 +12343,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
}
}
- else if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
- fb_lvalue) == GS_ERROR)
- {
- remove = true;
- break;
- }
-
- if (!remove
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
- && OMP_CLAUSE_CHAIN (c)
- && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
- && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ALWAYS_POINTER)
- || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ATTACH_DETACH)
- || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_TO_PSET)))
- prev_list_p = list_p;
-
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -12267,43 +12350,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
flags |= GOVD_MAP_ALWAYS_TO;
- if ((code == OMP_TARGET
- || code == OMP_TARGET_DATA
- || code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA)
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- {
- for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
- octx = octx->outer_context)
- {
- splay_tree_node n
- = splay_tree_lookup (octx->variables,
- (splay_tree_key) OMP_CLAUSE_DECL (c));
- /* If this is contained in an outer OpenMP region as a
- firstprivate value, remove the attach/detach. */
- if (n && (n->value & GOVD_FIRSTPRIVATE))
- {
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
- goto do_add;
- }
- }
-
- enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
- ? GOMP_MAP_DETACH
- : GOMP_MAP_ATTACH);
- OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
- }
- else if ((code == OACC_ENTER_DATA
- || code == OACC_EXIT_DATA
- || code == OACC_PARALLEL)
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
- {
- enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA
- ? GOMP_MAP_DETACH
- : GOMP_MAP_ATTACH);
- OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
- }
-
goto do_add;
case OMP_CLAUSE_AFFINITY:
@@ -12392,6 +12438,17 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
goto do_notice;
+ case OMP_CLAUSE__MAPPER_BINDING_:
+ {
+ tree name = OMP_CLAUSE__MAPPER_BINDING__ID (c);
+ tree var = OMP_CLAUSE__MAPPER_BINDING__DECL (c);
+ tree type = TYPE_MAIN_VARIANT (TREE_TYPE (var));
+ tree fndecl = OMP_CLAUSE__MAPPER_BINDING__MAPPER (c);
+ ctx->implicit_mappers->put ({ name, type }, fndecl);
+ remove = true;
+ break;
+ }
+
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
flags = GOVD_EXPLICIT;
@@ -12939,6 +12996,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
list_p = &OMP_CLAUSE_CHAIN (c);
}
+ if (groups)
+ {
+ delete grpmap;
+ delete groups;
+ }
+
ctx->clauses = *orig_list_p;
gimplify_omp_ctxp = ctx;
}
@@ -13406,15 +13469,96 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
+ if (code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ {
+ tree mapper_clauses = NULL_TREE;
+ instantiate_mapper_info im_info;
+
+ im_info.mapper_clauses_p = &mapper_clauses;
+ im_info.omp_ctx = ctx;
+ im_info.pre_p = pre_p;
+
+ splay_tree_foreach (ctx->variables,
+ omp_instantiate_implicit_mappers,
+ (void *) &im_info);
+
+ if (mapper_clauses)
+ {
+ mapper_clauses
+ = lang_hooks.decls.omp_finish_mapper_clauses (mapper_clauses);
+
+ /* Stick the implicitly-expanded mapper clauses at the end of the
+ clause list. */
+ tree *tail = list_p;
+ while (*tail)
+ tail = &OMP_CLAUSE_CHAIN (*tail);
+ *tail = mapper_clauses;
+ }
+
+ vec<omp_mapping_group> *groups;
+ groups = omp_gather_mapping_groups (list_p);
+ hash_map<tree_operand_hash, omp_mapping_group *> *grpmap = NULL;
+
+ if (groups)
+ {
+ grpmap = omp_index_mapping_groups (groups);
+
+ omp_resolve_clause_dependencies (code, groups, grpmap);
+ omp_build_struct_sibling_lists (code, ctx->region_type, groups,
+ &grpmap, list_p);
+
+ omp_mapping_group *outlist = NULL;
+
+ delete grpmap;
+ delete groups;
+
+ /* Rebuild now we have struct sibling lists. */
+ groups = omp_gather_mapping_groups (list_p);
+ grpmap = omp_index_mapping_groups (groups);
+
+ outlist = omp_tsort_mapping_groups (groups, grpmap);
+ outlist = omp_segregate_mapping_groups (outlist);
+ list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
+
+ delete grpmap;
+ delete groups;
+ }
+ }
+ else if (ctx->region_type & ORT_ACC)
+ {
+ vec<omp_mapping_group> *groups;
+ groups = omp_gather_mapping_groups (list_p);
+ if (groups)
+ {
+ hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
+ grpmap = omp_index_mapping_groups (groups);
+
+ oacc_resolve_clause_dependencies (groups, grpmap);
+ omp_build_struct_sibling_lists (code, ctx->region_type, groups,
+ &grpmap, list_p);
+
+ delete groups;
+ delete grpmap;
+ }
+ }
+
tree attach_list = NULL_TREE;
tree *attach_tail = &attach_list;
+ tree *grp_start_p = NULL, grp_end = NULL_TREE;
+
while ((c = *list_p) != NULL)
{
splay_tree_node n;
bool remove = false;
bool move_attach = false;
+ if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end))
+ grp_end = NULL_TREE;
+
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_FIRSTPRIVATE:
@@ -13569,26 +13713,67 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
break;
case OMP_CLAUSE_MAP:
- if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+ decl = OMP_CLAUSE_DECL (c);
+ if (!grp_end)
{
+ grp_start_p = list_p;
+ grp_end = *omp_group_last (grp_start_p);
+ }
+ switch (code)
+ {
+ case OMP_TARGET:
+ break;
+ case OACC_DATA:
+ if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
+ break;
+ goto check_firstprivate;
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ case OMP_TARGET_DATA:
+ case OMP_TARGET_ENTER_DATA:
+ case OMP_TARGET_EXIT_DATA:
+ case OACC_HOST_DATA:
+ check_firstprivate:
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ /* For target {,enter ,exit }data only the array slice is
+ mapped, but not the pointer to it. */
+ remove = true;
+ if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+ remove = true;
+ break;
+ default:
+ break;
+ }
+ if (remove)
+ break;
+ if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+ OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
+ : TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ gimplify_omp_ctxp = ctx->outer_context;
+ if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, is_gimple_val,
+ fb_rvalue) == GS_ERROR)
+ {
+ gimplify_omp_ctxp = ctx;
remove = true;
break;
}
- /* If we have a target region, we can push all the attaches to the
- end of the list (we may have standalone "attach" operations
- synthesized for GOMP_MAP_STRUCT nodes that must be processed after
- the attachment point AND the pointed-to block have been mapped).
- If we have something else, e.g. "enter data", we need to keep
- "attach" nodes together with the previous node they attach to so
- that separate "exit data" operations work properly (see
- libgomp/target.c). */
- if ((ctx->region_type & ORT_TARGET) != 0
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
- || (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
- move_attach = true;
- decl = OMP_CLAUSE_DECL (c);
+ else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
+ {
+ OMP_CLAUSE_SIZE (c)
+ = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
+ false);
+ if ((ctx->region_type & ORT_TARGET) != 0)
+ omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+ GOVD_FIRSTPRIVATE | GOVD_SEEN);
+ }
+ gimplify_omp_ctxp = ctx;
/* Data clauses associated with reductions must be
compatible with present_or_copy. Warn and adjust the clause
if that is not the case. */
@@ -13625,7 +13810,54 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
remove = true;
break;
}
- if (!DECL_P (decl))
+ /* For Fortran, not only the pointer to the data is mapped but also
+ the address of the pointer, the array descriptor etc.; for
+ 'exit data' - and in particular for 'delete:' - having an 'alloc:'
+ does not make sense. Likewise, for 'update' only transferring the
+ data itself is needed as the rest has been handled in previous
+ directives. However, for 'exit data', the array descriptor needs
+ to be deleted; hence, we turn the MAP_TO_PSET into a MAP_DELETE.
+
+ NOTE: Generally, it is not safe to perform "enter data" operations
+ on arrays where the data *or the descriptor* may go out of scope
+ before a corresponding "exit data" operation -- and such a
+ descriptor may be synthesized temporarily, e.g. to pass an
+ explicit-shape array to a function expecting an assumed-shape
+ argument. Performing "enter data" inside the called function
+ would thus be problematic. */
+ if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
+ && grp_start_p
+ && OMP_CLAUSE_CHAIN (*grp_start_p) == c)
+ OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*grp_start_p)
+ == GOMP_MAP_DELETE
+ ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
+ else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
+ {
+ remove = true;
+ break;
+ }
+ else if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ && OMP_CLAUSE_CHAIN (c)
+ && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
+ == OMP_CLAUSE_MAP)
+ && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ATTACH_DETACH)
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
+ && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
+ (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
+ 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))
{
if ((ctx->region_type & ORT_TARGET) != 0
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
@@ -13648,8 +13880,123 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
}
+
+ tree d = decl, *pd;
+ if (TREE_CODE (d) == ARRAY_REF)
+ {
+ while (TREE_CODE (d) == ARRAY_REF)
+ d = TREE_OPERAND (d, 0);
+ if (TREE_CODE (d) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
+ decl = d;
+ }
+ pd = &OMP_CLAUSE_DECL (c);
+ if (d == decl
+ && TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == 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);
+ }
+
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ switch (code)
+ {
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c)))
+ == ARRAY_TYPE)
+ remove = true;
+ else if (code == OACC_ENTER_DATA)
+ goto change_to_attach;
+ /* Fallthrough. */
+ case OMP_TARGET_EXIT_DATA:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DETACH);
+ break;
+ case OACC_UPDATE:
+ /* An "attach/detach" operation on an update directive
+ should behave as a GOMP_MAP_ALWAYS_POINTER. Note that
+ both GOMP_MAP_ATTACH_DETACH and GOMP_MAP_ALWAYS_POINTER
+ kinds depend on the previous mapping (for non-TARGET
+ regions). */
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
+ break;
+ default:
+ change_to_attach:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH);
+ if ((ctx->region_type & ORT_TARGET) != 0)
+ move_attach = true;
+ }
+ else if ((ctx->region_type & ORT_TARGET) != 0
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
+ move_attach = true;
+
+ /* If we have e.g. map(struct: *var), don't gimplify the
+ argument since omp-low.cc wants to see the decl itself. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+ break;
+
+ /* We've already partly gimplified this in
+ gimplify_scan_omp_clauses. Don't do any more. */
+ if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
+ break;
+
+ gimplify_omp_ctxp = ctx->outer_context;
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+ fb_lvalue) == GS_ERROR)
+ remove = true;
+ gimplify_omp_ctxp = ctx;
+
break;
}
+
+ if ((code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ {
+ bool firstprivatize = false;
+
+ for (struct gimplify_omp_ctx *octx = ctx->outer_context; octx;
+ octx = octx->outer_context)
+ {
+ splay_tree_node n
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) OMP_CLAUSE_DECL (c));
+ /* If this is contained in an outer OpenMP region as a
+ firstprivate value, remove the attach/detach. */
+ if (n && (n->value & GOVD_FIRSTPRIVATE))
+ {
+ firstprivatize = true;
+ break;
+ }
+ }
+
+ enum gomp_map_kind map_kind;
+ if (firstprivatize)
+ map_kind = GOMP_MAP_FIRSTPRIVATE_POINTER;
+ else if (code == OMP_TARGET_EXIT_DATA)
+ map_kind = GOMP_MAP_DETACH;
+ else
+ map_kind = GOMP_MAP_ATTACH;
+ OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+ }
+ else if ((ctx->region_type & ORT_ACC) != 0
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ {
+ enum gomp_map_kind map_kind = (code == OACC_EXIT_DATA
+ ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH);
+ OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+ }
+
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if ((ctx->region_type & ORT_TARGET) != 0
&& !(n->value & GOVD_SEEN)
@@ -13724,6 +14071,21 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
|| ((n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE))
== 0));
}
+
+ /* If we have a target region, we can push all the attaches to the
+ end of the list (we may have standalone "attach" operations
+ synthesized for GOMP_MAP_STRUCT nodes that must be processed after
+ the attachment point AND the pointed-to block have been mapped).
+ If we have something else, e.g. "enter data", we need to keep
+ "attach" nodes together with the previous node they attach to so
+ that separate "exit data" operations work properly (see
+ libgomp/target.c). */
+ if ((ctx->region_type & ORT_TARGET) != 0
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
+ move_attach = true;
+
break;
case OMP_CLAUSE_TO:
@@ -16930,6 +17292,15 @@ gimplify_omp_ordered (tree expr, gimple_seq body)
return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr));
}
+/* Gimplify an OMP_DECLARE_MAPPER node (by just removing it). */
+
+static enum gimplify_status
+gimplify_omp_declare_mapper (tree *expr_p)
+{
+ *expr_p = NULL_TREE;
+ return GS_ALL_DONE;
+}
+
/* Convert the GENERIC expression tree *EXPR_P to GIMPLE. If the
expression produces a value to be used as an operand inside a GIMPLE
statement, the value will be stored back in *EXPR_P. This value will
@@ -17852,6 +18223,10 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = gimplify_omp_atomic (expr_p, pre_p);
break;
+ case OMP_DECLARE_MAPPER:
+ ret = gimplify_omp_declare_mapper (expr_p);
+ break;
+
case TRANSACTION_EXPR:
ret = gimplify_transaction (expr_p, pre_p);
break;
@@ -85,6 +85,10 @@ extern enum omp_clause_defaultmap_kind lhd_omp_predetermined_mapping (tree);
extern tree lhd_omp_assignment (tree, tree, tree);
extern void lhd_omp_finish_clause (tree, gimple_seq *, bool);
extern tree lhd_omp_array_size (tree, gimple_seq *);
+extern tree lhd_omp_finish_mapper_clauses (tree);
+extern tree lhd_omp_mapper_lookup (tree, tree);
+extern tree lhd_omp_extract_mapper_directive (tree);
+extern tree lhd_omp_map_array_section (location_t, tree);
struct gimplify_omp_ctx;
extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
tree);
@@ -273,6 +277,11 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
#define LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR NULL
#define LANG_HOOKS_OMP_CLAUSE_DTOR hook_tree_tree_tree_null
#define LANG_HOOKS_OMP_FINISH_CLAUSE lhd_omp_finish_clause
+#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES lhd_omp_finish_mapper_clauses
+#define LANG_HOOKS_OMP_MAPPER_LOOKUP lhd_omp_mapper_lookup
+#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE \
+ lhd_omp_extract_mapper_directive
+#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION lhd_omp_map_array_section
#define LANG_HOOKS_OMP_ALLOCATABLE_P hook_bool_tree_false
#define LANG_HOOKS_OMP_SCALAR_P lhd_omp_scalar_p
#define LANG_HOOKS_OMP_SCALAR_TARGET_P hook_bool_tree_false
@@ -307,6 +316,10 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR, \
LANG_HOOKS_OMP_CLAUSE_DTOR, \
LANG_HOOKS_OMP_FINISH_CLAUSE, \
+ LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, \
+ LANG_HOOKS_OMP_MAPPER_LOOKUP, \
+ LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, \
+ LANG_HOOKS_OMP_MAP_ARRAY_SECTION, \
LANG_HOOKS_OMP_ALLOCATABLE_P, \
LANG_HOOKS_OMP_SCALAR_P, \
LANG_HOOKS_OMP_SCALAR_TARGET_P, \
@@ -642,6 +642,41 @@ lhd_omp_array_size (tree, gimple_seq *)
return NULL_TREE;
}
+/* Finalize clause list C after expanding custom mappers for implicitly-mapped
+ variables. */
+
+tree
+lhd_omp_finish_mapper_clauses (tree c)
+{
+ return c;
+}
+
+/* Look up an OpenMP "declare mapper" mapper. */
+
+tree
+lhd_omp_mapper_lookup (tree, tree)
+{
+ return NULL_TREE;
+}
+
+/* Given the representation used by the front-end to contain a mapper
+ directive, return the statement for the directive itself. */
+
+tree
+lhd_omp_extract_mapper_directive (tree)
+{
+ return error_mark_node;
+}
+
+/* Return a simplified form for OMP_ARRAY_SECTION argument, or
+ error_mark_node if impossible. */
+
+tree
+lhd_omp_map_array_section (location_t, tree)
+{
+ return error_mark_node;
+}
+
/* Return true if DECL is a scalar variable (for the purpose of
implicit firstprivatization & mapping). Only if alloc_ptr_ok
are allocatables and pointers accepted. */
@@ -313,6 +313,22 @@ struct lang_hooks_for_decls
/* Do language specific checking on an implicitly determined clause. */
void (*omp_finish_clause) (tree clause, gimple_seq *pre_p, bool);
+ /* Finish language-specific processing on mapping nodes after expanding
+ user-defined mappers. */
+ tree (*omp_finish_mapper_clauses) (tree clauses);
+
+ /* Find a mapper in the current parsing context, given a NAME (or
+ NULL_TREE) and TYPE. */
+ tree (*omp_mapper_lookup) (tree name, tree type);
+
+ /* Return the statement for the mapper directive definition, from the
+ representation used to contain it (e.g. an inline function
+ declaration). */
+ tree (*omp_extract_mapper_directive) (tree fndecl);
+
+ /* Return a simplified form for OMP_ARRAY_SECTION argument. */
+ tree (*omp_map_array_section) (location_t, tree t);
+
/* Return true if DECL is an allocatable variable (for the purpose of
implicit mapping). */
bool (*omp_allocatable_p) (tree decl);
@@ -152,6 +152,92 @@ get_openacc_privatization_dump_flags ()
extern tree omp_build_component_ref (tree obj, tree field);
+template <typename T>
+struct omp_name_type
+{
+ tree name;
+ T type;
+};
+
+template <>
+struct default_hash_traits <omp_name_type<tree> >
+ : typed_noop_remove <omp_name_type<tree> >
+{
+ GTY((skip)) typedef omp_name_type<tree> value_type;
+ GTY((skip)) typedef omp_name_type<tree> compare_type;
+
+ static hashval_t
+ hash (omp_name_type<tree> p)
+ {
+ return p.name ? iterative_hash_expr (p.name, TYPE_UID (p.type))
+ : TYPE_UID (p.type);
+ }
+
+ static const bool empty_zero_p = true;
+
+ static bool
+ is_empty (omp_name_type<tree> p)
+ {
+ return p.type == NULL;
+ }
+
+ static bool
+ is_deleted (omp_name_type<tree>)
+ {
+ return false;
+ }
+
+ static bool
+ equal (const omp_name_type<tree> &a, const omp_name_type<tree> &b)
+ {
+ if (a.name == NULL_TREE && b.name == NULL_TREE)
+ return a.type == b.type;
+ else if (a.name == NULL_TREE || b.name == NULL_TREE)
+ return false;
+ else
+ return a.name == b.name && a.type == b.type;
+ }
+
+ static void
+ mark_empty (omp_name_type<tree> &e)
+ {
+ e.type = NULL;
+ }
+};
+
+template <typename T>
+struct omp_mapper_list
+{
+ hash_set<omp_name_type<T>> *seen_types;
+ vec<tree> *mappers;
+
+ omp_mapper_list (hash_set<omp_name_type<T>> *s, vec<tree> *m)
+ : seen_types (s), mappers (m) { }
+
+ void add_mapper (tree name, T type, tree mapperfn)
+ {
+ /* We can't hash a NULL_TREE... */
+ if (!name)
+ name = void_node;
+
+ omp_name_type<T> n_t = { name, type };
+
+ if (seen_types->contains (n_t))
+ return;
+
+ seen_types->add (n_t);
+ mappers->safe_push (mapperfn);
+ }
+
+ bool contains (tree name, T type)
+ {
+ if (!name)
+ name = void_node;
+
+ return seen_types->contains ({ name, type });
+ }
+};
+
namespace omp_addr_tokenizer {
/* These are the ways of accessing a variable that have special-case handling
new file mode 100644
@@ -0,0 +1,22 @@
+/* { dg-do compile { target c++ } } */
+
+struct XYZ {
+ int a;
+ int *b;
+ int c;
+};
+
+#pragma omp declare mapper(struct XYZ t)
+/* { dg-error "missing 'map' clause" "" { target c } .-1 } */
+/* { dg-error "missing 'map' clause before end of line" "" { target c++ } .-2 } */
+
+struct ABC {
+ int *a;
+ int b;
+ int c;
+};
+
+#pragma omp declare mapper(struct ABC d) firstprivate(d.b)
+/* { dg-error "unexpected clause" "" { target c } .-1 } */
+/* { dg-error "expected end of line before '\\(' token" "" { target c } .-2 } */
+/* { dg-error "unexpected clause before '\\(' token" "" { target c++ } .-3 } */
new file mode 100644
@@ -0,0 +1,30 @@
+// { dg-do compile { target c++ } }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+#include <stdlib.h>
+
+// Test named mapper invocation.
+
+struct S {
+ int *ptr;
+ int size;
+};
+
+int main (int argc, char *argv[])
+{
+ int N = 1024;
+#pragma omp declare mapper (mapN:struct S s) map(to:s.ptr, s.size) \
+ map(s.ptr[:N])
+
+ struct S s;
+ s.ptr = (int *) malloc (sizeof (int) * N);
+
+#pragma omp target map(mapper(mapN), tofrom: s)
+// { dg-final { scan-tree-dump {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} "gimple" } }
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,78 @@
+/* { dg-do compile { target c++ } } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+/* Check mapper binding clauses. */
+
+struct Y {
+ int z;
+};
+
+struct Z {
+ int z;
+};
+
+#pragma omp declare mapper (struct Y y) map(tofrom: y)
+#pragma omp declare mapper (struct Z z) map(tofrom: z)
+
+int foo (void)
+{
+ struct Y yy;
+ struct Z zz;
+ int dummy;
+
+#pragma omp target data map(dummy)
+ {
+ #pragma omp target
+ {
+ yy.z++;
+ zz.z++;
+ }
+ yy.z++;
+ }
+ return yy.z;
+}
+
+struct P
+{
+ struct Z *zp;
+};
+
+int bar (void)
+{
+ struct Y yy;
+ struct Z zz;
+ struct P pp;
+ struct Z t;
+ int dummy;
+
+ pp.zp = &t;
+
+#pragma omp declare mapper (struct Y y) map(tofrom: y.z)
+#pragma omp declare mapper (struct Z z) map(tofrom: z.z)
+
+#pragma omp target data map(dummy)
+ {
+ #pragma omp target
+ {
+ yy.z++;
+ zz.z++;
+ }
+ yy.z++;
+ }
+
+ #pragma omp declare mapper(struct P x) map(to:x.zp) map(tofrom:*x.zp)
+
+ #pragma omp target
+ {
+ zz = *pp.zp;
+ }
+
+ return zz.z;
+}
+
+/* { dg-final { scan-tree-dump-times {mapper_binding\(struct Y,omp declare mapper ~1Y\) mapper_binding\(struct Z,omp declare mapper ~1Z\)} 2 "original" { target c++ } } } */
+/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,omp declare mapper ~1Z\) mapper_binding\(struct P,omp declare mapper ~1P\)} "original" { target c++ } } } */
+
+/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\)\) mapper_binding\(struct Y,#pragma omp declare mapper \(struct Y y\) map\(tofrom:y\)\)} "original" { target c } } } */
+/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\.z\)\) mapper_binding\(struct Y,#pragma omp declare mapper \(struct Y y\) map\(tofrom:y\.z\)\)} "original" { target c } } } */
+/* { dg-final { scan-tree-dump {mapper_binding\(struct P,#pragma omp declare mapper \(struct P x\) map\(tofrom:\(x\.zp\)\[0:1\]\) map\(to:x.zp\)\) mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\.z\)\)} "original" { target c } } } */
new file mode 100644
@@ -0,0 +1,26 @@
+/* { dg-do compile { target c++ } } */
+
+typedef struct S_ {
+ int *myarr;
+ int size;
+} S;
+
+#pragma omp declare mapper (named: struct S_ v) map(to:v.size, v.myarr) \
+ map(tofrom: v.myarr[0:v.size])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-2 } */
+/* { dg-note "'#pragma omp declare mapper \\(named: S_\\)' previously defined here" "" { target c++ } .-3 } */
+
+#pragma omp declare mapper (named: S v) map(to:v.size, v.myarr) \
+ map(tofrom: v.myarr[0:v.size])
+/* { dg-error "redeclaration of 'named' '#pragma omp declare mapper' for type 'S' \\\{aka 'struct S_'\\\}" "" { target c } .-2 } */
+/* { dg-error "redefinition of '#pragma omp declare mapper \\(named: S\\)'" "" { target c++ } .-3 } */
+
+#pragma omp declare mapper (struct S_ v) map(to:v.size, v.myarr) \
+ map(tofrom: v.myarr[0:v.size])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-2 } */
+/* { dg-note "'#pragma omp declare mapper \\(S_\\)' previously defined here" "" { target c++ } .-3 } */
+
+#pragma omp declare mapper (S v) map(to:v.size, v.myarr) \
+ map(tofrom: v.myarr[0:v.size])
+/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'S' \\\{aka 'struct S_'\\\}" "" { target c } .-2 } */
+/* { dg-error "redefinition of '#pragma omp declare mapper \\(S\\)'" "" { target c++ } .-3 } */
new file mode 100644
@@ -0,0 +1,23 @@
+/* { dg-do compile { target c++ } } */
+
+int x = 5;
+
+struct Q {
+ int *arr1;
+ int *arr2;
+ int *arr3;
+};
+
+#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
+
+struct R {
+ int *arr1;
+ int *arr2;
+ int *arr3;
+};
+
+#pragma omp declare mapper (struct R myr) map(myr.arr3[0:y])
+/* { dg-error "'y' undeclared" "" { target c } .-1 } */
+/* { dg-error "'y' was not declared in this scope" "" { target c++ } .-2 } */
+
+int y = 7;
new file mode 100644
@@ -0,0 +1,29 @@
+/* { dg-do compile { target c++ } } */
+
+struct Q {
+ int *arr1;
+ int *arr2;
+ int *arr3;
+};
+
+int foo (void)
+{
+ int x = 5;
+ #pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
+ return x;
+}
+
+struct R {
+ int *arr1;
+ int *arr2;
+ int *arr3;
+};
+
+int bar (void)
+{
+ #pragma omp declare mapper (struct R myr) map(myr.arr3[0:y])
+ /* { dg-error "'y' undeclared" "" { target c } .-1 } */
+ /* { dg-error "'y' was not declared in this scope" "" { target c++ } .-2 } */
+ int y = 7;
+ return y;
+}
new file mode 100644
@@ -0,0 +1,43 @@
+/* { dg-do compile { target c++ } } */
+
+struct Q {
+ int *arr1;
+ int *arr2;
+ int *arr3;
+ int len;
+};
+
+struct R {
+ struct Q qarr[5];
+};
+
+struct R2 {
+ struct Q *qptr;
+};
+
+#pragma omp declare mapper (struct Q myq) map(myq.arr1[0:myq.len]) \
+ map(myq.arr2[0:myq.len]) \
+ map(myq.arr3[0:myq.len])
+
+#pragma omp declare mapper (struct R myr) map(myr.qarr[2:3])
+
+#pragma omp declare mapper (struct R2 myr2) map(myr2.qptr[2:3])
+
+int main (int argc, char *argv[])
+{
+ struct R r;
+ struct R2 r2;
+ int N = 256;
+
+#pragma omp target
+/* { dg-message "sorry, unimplemented: user-defined mapper with non-unit length array section" "" { target *-*-* } .-1 } */
+ {
+ for (int i = 2; i < 5; i++)
+ for (int j = 0; j < N; j++)
+ {
+ r.qarr[i].arr1[j]++;
+ r2.qptr[i].arr2[j]++;
+ }
+ }
+}
+
new file mode 100644
@@ -0,0 +1,34 @@
+/* { dg-do compile { target c++ } } */
+
+int x = 5;
+
+struct Q {
+ int *arr1;
+ int *arr2;
+ int *arr3;
+};
+
+int y = 5;
+
+#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */
+/* { dg-note "'#pragma omp declare mapper \\(Q\\)' previously defined here" "" { target c++ } .-2 } */
+
+#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:y])
+/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'struct Q'" "" { target c } .-1 } */
+/* { dg-error "redefinition of '#pragma omp declare mapper \\(Q\\)'" "" { target c++ } .-2 } */
+
+struct R {
+ int *arr1;
+};
+
+void foo (void)
+{
+#pragma omp declare mapper (struct R myr) map(myr.arr1[0:x])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */
+/* { dg-note "'#pragma omp declare mapper \\(R\\)' previously declared here" "" { target c++ } .-2 } */
+
+#pragma omp declare mapper (struct R myr) map(myr.arr1[0:y])
+/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'struct R'" "" { target c } .-1 } */
+/* { dg-error "redeclaration of '#pragma omp declare mapper \\(R\\)'" "" { target c++ } .-2 } */
+}
@@ -13,10 +13,12 @@ foo (void)
#pragma omp target map (to:a)
;
- #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+ #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */
+/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */
;
- #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+ #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */
+/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */
;
#pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */
new file mode 100644
@@ -0,0 +1,58 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+// "omp declare mapper" support -- check expansion in gimple.
+
+struct S {
+ int *ptr;
+ int size;
+};
+
+#define N 64
+
+#pragma omp declare mapper (S w) map(w.size, w.ptr, w.ptr[:w.size])
+#pragma omp declare mapper (foo:S w) map(to:w.size, w.ptr) map(w.ptr[:w.size])
+
+int main (int argc, char *argv[])
+{
+ S s;
+ s.ptr = new int[N];
+ s.size = N;
+
+#pragma omp declare mapper (bar:S w) map(w.size, w.ptr, w.ptr[:w.size])
+
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+#pragma omp target map(tofrom: s)
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+#pragma omp target map(mapper(default), tofrom: s)
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+#pragma omp target map(mapper(foo), alloc: s)
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+#pragma omp target map(mapper(bar), tofrom: s)
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+ return 0;
+}
+
+// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(tofrom:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 4 "gimple" } }
+// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } }
new file mode 100644
@@ -0,0 +1,30 @@
+// { dg-do compile }
+
+// Error-checking tests for "omp declare mapper".
+
+struct S {
+ int *ptr;
+ int size;
+};
+
+struct Z {
+ int z;
+};
+
+int main (int argc, char *argv[])
+{
+#pragma omp declare mapper (S v) map(v.size, v.ptr[:v.size]) // { dg-note "'#pragma omp declare mapper \\(S\\)' previously declared here" }
+
+ /* This one's a duplicate. */
+#pragma omp declare mapper (default: S v) map (to: v.size) map (v) // { dg-error "redeclaration of '#pragma omp declare mapper \\(S\\)'" }
+
+ /* ...and this one doesn't use a "base language identifier" for the mapper
+ name. */
+#pragma omp declare mapper (case: S v) map (to: v.size) // { dg-error "expected identifier or 'default' before 'case'" }
+ // { dg-error "expected ':' before 'case'" "" { target *-*-* } .-1 }
+
+ /* A non-struct/class/union type isn't supposed to work. */
+#pragma omp declare mapper (name:Z [5]foo) map (foo[0].z) // { dg-error "'Z \\\[5\\\]' is not a struct, union or class type in '#pragma omp declare mapper'" }
+
+ return 0;
+}
@@ -348,6 +348,10 @@ enum omp_clause_code {
/* OpenMP clause: doacross ({source,sink}:vec). */
OMP_CLAUSE_DOACROSS,
+ /* OpenMP mapper binding: record implicit mappers in scope for aggregate
+ types used within an offload region. */
+ OMP_CLAUSE__MAPPER_BINDING_,
+
/* Internal structure to hold OpenACC cache directive's variable-list.
#pragma acc cache (variable-list). */
OMP_CLAUSE__CACHE_,
@@ -994,6 +994,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, "attach_zero_length_array_section");
break;
+ case GOMP_MAP_UNSET:
+ pp_string (pp, "unset");
+ break;
+ case GOMP_MAP_PUSH_MAPPER_NAME:
+ pp_string (pp, "push_mapper");
+ break;
+ case GOMP_MAP_POP_MAPPER_NAME:
+ pp_string (pp, "pop_mapper");
+ break;
default:
gcc_unreachable ();
}
@@ -1057,6 +1066,23 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
spc, flags, false);
goto print_clause_size;
+ case OMP_CLAUSE__MAPPER_BINDING_:
+ pp_string (pp, "mapper_binding(");
+ if (OMP_CLAUSE__MAPPER_BINDING__ID (clause))
+ {
+ dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__ID (clause), spc,
+ flags, false);
+ pp_comma (pp);
+ }
+ dump_generic_node (pp,
+ TREE_TYPE (OMP_CLAUSE__MAPPER_BINDING__DECL (clause)),
+ spc, flags, false);
+ pp_comma (pp);
+ dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__MAPPER (clause), spc,
+ flags, false);
+ pp_right_paren (pp);
+ break;
+
case OMP_CLAUSE_NUM_TEAMS:
pp_string (pp, "num_teams(");
if (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (clause))
@@ -3833,6 +3859,21 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
is_expr = false;
break;
+ case OMP_DECLARE_MAPPER:
+ pp_string (pp, "#pragma omp declare mapper (");
+ if (OMP_DECLARE_MAPPER_ID (node))
+ {
+ dump_generic_node (pp, OMP_DECLARE_MAPPER_ID (node), spc, flags,
+ false);
+ pp_colon (pp);
+ }
+ dump_generic_node (pp, TREE_TYPE (node), spc, flags, false);
+ pp_space (pp);
+ dump_generic_node (pp, OMP_DECLARE_MAPPER_DECL (node), spc, flags, false);
+ pp_right_paren (pp);
+ dump_omp_clauses (pp, OMP_DECLARE_MAPPER_CLAUSES (node), spc, flags);
+ break;
+
case TRANSACTION_EXPR:
if (TRANSACTION_EXPR_OUTER (node))
pp_string (pp, "__transaction_atomic [[outer]]");
@@ -295,6 +295,7 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_MAP */
1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
1, /* OMP_CLAUSE_DOACROSS */
+ 3, /* OMP_CLAUSE__MAPPER_BINDING_ */
2, /* OMP_CLAUSE__CACHE_ */
2, /* OMP_CLAUSE_GANG */
1, /* OMP_CLAUSE_ASYNC */
@@ -386,6 +387,7 @@ const char * const omp_clause_code_name[] =
"map",
"has_device_addr",
"doacross",
+ "_mapper_binding_",
"_cache_",
"gang",
"async",
@@ -1247,6 +1247,13 @@ DEFTREECODE (OMP_SECTION, "omp_section", tcc_statement, 1)
Operand 0: OMP_MASTER_BODY: Master section body. */
DEFTREECODE (OMP_MASTER, "omp_master", tcc_statement, 1)
+/* OpenMP - #pragma omp declare mapper ([id:] type var) [clause1 ... clauseN]
+ Operand 0: Identifier.
+ Operand 1: Variable decl.
+ Operand 2: List of clauses.
+ The type of the construct is used for the type to be mapped. */
+DEFTREECODE (OMP_DECLARE_MAPPER, "omp_declare_mapper", tcc_statement, 3)
+
/* OpenACC - #pragma acc cache (variable1 ... variableN)
Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into
OMP_CLAUSE__CACHE_ clauses). */
@@ -1532,6 +1532,13 @@ class auto_suppress_location_wrappers
#define OMP_TARGET_EXIT_DATA_CLAUSES(NODE)\
TREE_OPERAND (OMP_TARGET_EXIT_DATA_CHECK (NODE), 0)
+#define OMP_DECLARE_MAPPER_ID(NODE) \
+ TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 0)
+#define OMP_DECLARE_MAPPER_DECL(NODE) \
+ TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 1)
+#define OMP_DECLARE_MAPPER_CLAUSES(NODE) \
+ TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 2)
+
#define OMP_SCAN_BODY(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0)
#define OMP_SCAN_CLAUSES(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1)
@@ -1972,6 +1979,18 @@ class auto_suppress_location_wrappers
#define OMP_CLAUSE__SCANTEMP__CONTROL(NODE) \
TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__SCANTEMP_))
+#define OMP_CLAUSE__MAPPER_BINDING__ID(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
+ OMP_CLAUSE__MAPPER_BINDING_), 0)
+
+#define OMP_CLAUSE__MAPPER_BINDING__DECL(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
+ OMP_CLAUSE__MAPPER_BINDING_), 1)
+
+#define OMP_CLAUSE__MAPPER_BINDING__MAPPER(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
+ OMP_CLAUSE__MAPPER_BINDING_), 2)
+
/* SSA_NAME accessors. */
/* Whether SSA_NAME NODE is a virtual operand. This simply caches the
@@ -188,7 +188,13 @@ enum gomp_map_kind
/* An attach or detach operation. Rewritten to the appropriate type during
gimplification, depending on directive (i.e. "enter data" or
parallel/kernels region vs. "exit data"). */
- GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3)
+ GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3),
+ /* Unset, used for "declare mapper" maps with no explicit data movement
+ specified. These use the movement specified at the invocation site. */
+ GOMP_MAP_UNSET = (GOMP_MAP_LAST | 4),
+ /* Used to record the name of a named mapper. */
+ GOMP_MAP_PUSH_MAPPER_NAME = (GOMP_MAP_LAST | 5),
+ GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 6)
};
#define GOMP_MAP_COPY_TO_P(X) \
new file mode 100644
@@ -0,0 +1,87 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+#define N 64
+
+struct points
+{
+ double *x;
+ double *y;
+ double *z;
+ size_t len;
+};
+
+#pragma omp declare mapper(points p) map(to:p.x, p.y, p.z) \
+ map(p.x[0:p.len]) \
+ map(p.y[0:p.len]) \
+ map(p.z[0:p.len])
+
+struct shape
+{
+ points tmp;
+ points *pts;
+ int metadata[128];
+};
+
+#pragma omp declare mapper(shape s) map(tofrom:s.pts, *s.pts) map(alloc:s.tmp)
+
+void
+alloc_points (points *pts, size_t sz)
+{
+ pts->x = new double[sz];
+ pts->y = new double[sz];
+ pts->z = new double[sz];
+ pts->len = sz;
+ for (int i = 0; i < sz; i++)
+ pts->x[i] = pts->y[i] = pts->z[i] = 0;
+}
+
+int main (int argc, char *argv[])
+{
+ shape myshape;
+ points mypts;
+
+ myshape.pts = &mypts;
+
+ alloc_points (&myshape.tmp, N);
+ myshape.pts = new points;
+ alloc_points (myshape.pts, N);
+
+ #pragma omp target map(myshape)
+ {
+ for (int i = 0; i < N; i++)
+ {
+ myshape.pts->x[i]++;
+ myshape.pts->y[i]++;
+ myshape.pts->z[i]++;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (myshape.pts->x[i] == 1);
+ assert (myshape.pts->y[i] == 1);
+ assert (myshape.pts->z[i] == 1);
+ }
+
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ {
+ myshape.pts->x[i]++;
+ myshape.pts->y[i]++;
+ myshape.pts->z[i]++;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (myshape.pts->x[i] == 2);
+ assert (myshape.pts->y[i] == 2);
+ assert (myshape.pts->z[i] == 2);
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,55 @@
+// { dg-do run }
+
+#include <cassert>
+
+#define N 256
+
+struct doublebuf
+{
+ int buf_a[N][N];
+ int buf_b[N][N];
+};
+
+#pragma omp declare mapper(lo:doublebuf b) map(b.buf_a[0:N/2][0:N]) \
+ map(b.buf_b[0:N/2][0:N])
+
+#pragma omp declare mapper(hi:doublebuf b) map(b.buf_a[N/2:N/2][0:N]) \
+ map(b.buf_b[N/2:N/2][0:N])
+
+int main (int argc, char *argv[])
+{
+ doublebuf db;
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ db.buf_a[i][j] = db.buf_b[i][j] = 0;
+
+ #pragma omp target map(mapper(lo), tofrom:db)
+ {
+ for (int i = 0; i < N / 2; i++)
+ for (int j = 0; j < N; j++)
+ {
+ db.buf_a[i][j]++;
+ db.buf_b[i][j]++;
+ }
+ }
+
+ #pragma omp target map(mapper(hi), tofrom:db)
+ {
+ for (int i = N / 2; i < N; i++)
+ for (int j = 0; j < N; j++)
+ {
+ db.buf_a[i][j]++;
+ db.buf_b[i][j]++;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ {
+ assert (db.buf_a[i][j] == 1);
+ assert (db.buf_b[i][j] == 1);
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,63 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+struct S {
+ int *myarr;
+};
+
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20])
+
+namespace A {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100])
+}
+
+namespace B {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100])
+}
+
+namespace A
+{
+ void incr_a (S my_s)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 100; i++)
+ my_s.myarr[i]++;
+ }
+ }
+}
+
+namespace B
+{
+ void incr_b (S my_s)
+ {
+#pragma omp target
+ {
+ for (int i = 100; i < 200; i++)
+ my_s.myarr[i]++;
+ }
+ }
+}
+
+int main (int argc, char *argv[])
+{
+ S my_s;
+
+ my_s.myarr = (int *) calloc (200, sizeof (int));
+
+#pragma omp target
+ {
+ for (int i = 0; i < 20; i++)
+ my_s.myarr[i]++;
+ }
+
+ A::incr_a (my_s);
+ B::incr_b (my_s);
+
+ for (int i = 0; i < 200; i++)
+ assert (my_s.myarr[i] == (i < 20) ? 2 : 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,63 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+struct S {
+ int *myarr;
+};
+
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20])
+
+namespace A {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100])
+}
+
+namespace B {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100])
+}
+
+namespace A
+{
+ void incr_a (S &my_s)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 100; i++)
+ my_s.myarr[i]++;
+ }
+ }
+}
+
+namespace B
+{
+ void incr_b (S &my_s)
+ {
+#pragma omp target
+ {
+ for (int i = 100; i < 200; i++)
+ my_s.myarr[i]++;
+ }
+ }
+}
+
+int main (int argc, char *argv[])
+{
+ S my_s;
+
+ my_s.myarr = (int *) calloc (200, sizeof (int));
+
+#pragma omp target
+ {
+ for (int i = 0; i < 20; i++)
+ my_s.myarr[i]++;
+ }
+
+ A::incr_a (my_s);
+ B::incr_b (my_s);
+
+ for (int i = 0; i < 200; i++)
+ assert (my_s.myarr[i] == (i < 20) ? 2 : 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,52 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+ int *myarr;
+ int len;
+};
+
+class C
+{
+ S smemb;
+#pragma omp declare mapper (custom:S s) map(to:s.myarr) \
+ map(tofrom:s.myarr[0:s.len])
+
+public:
+ C(int l)
+ {
+ smemb.myarr = new int[l];
+ smemb.len = l;
+ for (int i = 0; i < l; i++)
+ smemb.myarr[i] = 0;
+ }
+ void bump();
+ void check();
+};
+
+void
+C::bump ()
+{
+#pragma omp target map(mapper(custom), tofrom: smemb)
+ {
+ for (int i = 0; i < smemb.len; i++)
+ smemb.myarr[i]++;
+ }
+}
+
+void
+C::check ()
+{
+ for (int i = 0; i < smemb.len; i++)
+ assert (smemb.myarr[i] == 1);
+}
+
+int main (int argc, char *argv[])
+{
+ C test (100);
+ test.bump ();
+ test.check ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+// { dg-do run }
+
+#include <cassert>
+
+template <typename T>
+void adjust (T param)
+{
+#pragma omp declare mapper (T x) map(to:x.len, x.base) \
+ map(tofrom:x.base[0:x.len])
+
+#pragma omp target
+ for (int i = 0; i < param.len; i++)
+ param.base[i]++;
+}
+
+struct S {
+ int len;
+ int *base;
+};
+
+int main (int argc, char *argv[])
+{
+ S a;
+
+ a.len = 100;
+ a.base = new int[a.len];
+
+ for (int i = 0; i < a.len; i++)
+ a.base[i] = 0;
+
+ adjust (a);
+
+ for (int i = 0; i < a.len; i++)
+ assert (a.base[i] == 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,48 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+ int *myarr;
+};
+
+struct T
+{
+ S *s;
+};
+
+#pragma omp declare mapper (s100: S x) map(to: x.myarr) \
+ map(tofrom: x.myarr[0:100])
+
+void
+bump (T t)
+{
+ /* Here we have an implicit/default mapper invoking a named mapper. We
+ need to make sure that can be located properly at gimplification
+ time. */
+#pragma omp declare mapper (T t) map(to:t.s) map(mapper(s100), tofrom: t.s[0])
+
+#pragma omp target
+ for (int i = 0; i < 100; i++)
+ t.s->myarr[i]++;
+}
+
+int main (int argc, char *argv[])
+{
+ S my_s;
+ T my_t;
+
+ my_s.myarr = new int[100];
+ my_t.s = &my_s;
+
+ for (int i = 0; i < 100; i++)
+ my_s.myarr[i] = 0;
+
+ bump (my_t);
+
+ for (int i = 0; i < 100; i++)
+ assert (my_s.myarr[i] == 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,61 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+ int *myarr;
+ int len;
+};
+
+template<typename T>
+class C
+{
+ T memb;
+#pragma omp declare mapper (T t) map(to:t.len, t.myarr) \
+ map(tofrom:t.myarr[0:t.len])
+
+public:
+ C(int sz);
+ ~C();
+ void bump();
+ void check();
+};
+
+template<typename T>
+C<T>::C(int sz)
+{
+ memb.myarr = new int[sz];
+ for (int i = 0; i < sz; i++)
+ memb.myarr[i] = 0;
+ memb.len = sz;
+}
+
+template<typename T>
+C<T>::~C()
+{
+ delete[] memb.myarr;
+}
+
+template<typename T>
+void C<T>::bump()
+{
+#pragma omp target map(memb)
+ for (int i = 0; i < memb.len; i++)
+ memb.myarr[i]++;
+}
+
+template<typename T>
+void C<T>::check()
+{
+ for (int i = 0; i < memb.len; i++)
+ assert (memb.myarr[i] == 1);
+}
+
+int main(int argc, char *argv[])
+{
+ C<S> c_int(100);
+ c_int.bump();
+ c_int.check();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,60 @@
+/* { dg-do run { target c++ } } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+
+typedef struct {
+ int *arr;
+ int size;
+} B;
+
+#pragma omp declare mapper (mapB : B myb) map(to: myb.size, myb.arr) \
+ map(tofrom: myb.arr[0:myb.size])
+
+struct A {
+ int *arr1;
+ B *arr2;
+ int arr3[N];
+};
+
+int
+main (int argc, char *argv[])
+{
+ struct A var;
+
+ memset (&var, 0, sizeof var);
+ var.arr1 = (int *) calloc (N, sizeof (int));
+ var.arr2 = (B *) malloc (sizeof (B));
+ var.arr2->arr = (int *) calloc (N, sizeof (float));
+ var.arr2->size = N;
+
+ {
+ #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
+ map(tofrom: x.arr1[0:N]) \
+ map(mapper(mapB), tofrom: x.arr2[0:1])
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ {
+ var.arr1[i]++;
+ var.arr2->arr[i]++;
+ }
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (var.arr1[i] == 1);
+ assert (var.arr2->arr[i] == 1);
+ assert (var.arr3[i] == 0);
+ }
+
+ free (var.arr1);
+ free (var.arr2->arr);
+ free (var.arr2);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,59 @@
+/* { dg-do run { target c++ } } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+
+typedef struct B_tag {
+ int *arr;
+ int size;
+} B;
+
+#pragma omp declare mapper (B myb) map(to: myb.size, myb.arr) \
+ map(tofrom: myb.arr[0:myb.size])
+
+struct A {
+ int *arr1;
+ B *arr2;
+ int arr3[N];
+};
+
+int
+main (int argc, char *argv[])
+{
+ struct A var;
+
+ memset (&var, 0, sizeof var);
+ var.arr1 = (int *) calloc (N, sizeof (int));
+ var.arr2 = (B *) malloc (sizeof (B));
+ var.arr2->arr = (int *) calloc (N, sizeof (int));
+ var.arr2->size = N;
+
+ {
+ #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
+ map(tofrom: x.arr1[0:N]) map(tofrom: x.arr2[0:1])
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ {
+ var.arr1[i]++;
+ var.arr2->arr[i]++;
+ }
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (var.arr1[i] == 1);
+ assert (var.arr2->arr[i] == 1);
+ assert (var.arr3[i] == 0);
+ }
+
+ free (var.arr1);
+ free (var.arr2->arr);
+ free (var.arr2);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,87 @@
+/* { dg-do run { target c++ } } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+
+typedef struct {
+ int *arr;
+ int size;
+} B;
+
+#pragma omp declare mapper (samename : B myb) map(to: myb.size, myb.arr) \
+ map(tofrom: myb.arr[0:myb.size])
+
+typedef struct {
+ int *arr;
+ int size;
+} C;
+
+
+struct A {
+ int *arr1;
+ B *arr2;
+ C *arr3;
+};
+
+int
+main (int argc, char *argv[])
+{
+ struct A var;
+
+ memset (&var, 0, sizeof var);
+ var.arr1 = (int *) calloc (N, sizeof (int));
+ var.arr2 = (B *) malloc (sizeof (B));
+ var.arr2->arr = (int *) calloc (N, sizeof (int));
+ var.arr2->size = N;
+ var.arr3 = (C *) malloc (sizeof (C));
+ var.arr3->arr = (int *) calloc (N, sizeof (int));
+ var.arr3->size = N;
+
+ {
+ #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
+ map(tofrom: x.arr1[0:N]) \
+ map(mapper(samename), tofrom: x.arr2[0:1])
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ {
+ var.arr1[i]++;
+ var.arr2->arr[i]++;
+ }
+ }
+ }
+
+ {
+ #pragma omp declare mapper (samename : C myc) map(to: myc.size, myc.arr) \
+ map(tofrom: myc.arr[0:myc.size])
+ #pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr3) \
+ map(tofrom: x.arr1[0:N]) \
+ map(mapper(samename), tofrom: *x.arr3)
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ {
+ var.arr1[i]++;
+ var.arr3->arr[i]++;
+ }
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (var.arr1[i] == 2);
+ assert (var.arr2->arr[i] == 1);
+ assert (var.arr3->arr[i] == 1);
+ }
+
+ free (var.arr1);
+ free (var.arr2->arr);
+ free (var.arr2);
+ free (var.arr3->arr);
+ free (var.arr3);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,55 @@
+/* { dg-do run { target c++ } } */
+
+#include <assert.h>
+
+struct T {
+ int a;
+ int b;
+ int c;
+};
+
+void foo (void)
+{
+ struct T x;
+ x.a = x.b = x.c = 0;
+
+#pragma omp target
+ {
+ x.a++;
+ x.c++;
+ }
+
+ assert (x.a == 1);
+ assert (x.b == 0);
+ assert (x.c == 1);
+}
+
+// An identity mapper. This should do the same thing as the default!
+#pragma omp declare mapper (struct T v) map(v)
+
+void bar (void)
+{
+ struct T x;
+ x.a = x.b = x.c = 0;
+
+#pragma omp target
+ {
+ x.b++;
+ }
+
+#pragma omp target map(x)
+ {
+ x.a++;
+ }
+
+ assert (x.a == 1);
+ assert (x.b == 1);
+ assert (x.c == 0);
+}
+
+int main (int argc, char *argv[])
+{
+ foo ();
+ bar ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,57 @@
+/* { dg-do run { target c++ } } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+struct Z {
+ int *arr;
+};
+
+void baz (struct Z *zarr, int len)
+{
+#pragma omp declare mapper (struct Z myvar) map(to: myvar.arr) \
+ map(tofrom: myvar.arr[0:len])
+ zarr[0].arr = (int *) calloc (len, sizeof (int));
+ zarr[5].arr = (int *) calloc (len, sizeof (int));
+
+#pragma omp target map(zarr, *zarr)
+ {
+ for (int i = 0; i < len; i++)
+ zarr[0].arr[i]++;
+ }
+
+#pragma omp target map(zarr, zarr[5])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[5].arr[i]++;
+ }
+
+#pragma omp target map(zarr[5])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[5].arr[i]++;
+ }
+
+#pragma omp target map(zarr, zarr[5:1])
+ {
+ for (int i = 0; i < len; i++)
+ zarr[5].arr[i]++;
+ }
+
+ for (int i = 0; i < len; i++)
+ assert (zarr[0].arr[i] == 1);
+
+ for (int i = 0; i < len; i++)
+ assert (zarr[5].arr[i] == 3);
+
+ free (zarr[5].arr);
+ free (zarr[0].arr);
+}
+
+int
+main (int argc, char *argv[])
+{
+ struct Z myzarr[10];
+ baz (myzarr, 256);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,62 @@
+/* { dg-do run { target c++ } } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <assert.h>
+
+#define N 64
+
+struct A {
+ int *arr1;
+ float *arr2;
+ int arr3[N];
+};
+
+int
+main (int argc, char *argv[])
+{
+ struct A var;
+
+ memset (&var, 0, sizeof var);
+ var.arr1 = (int *) calloc (N, sizeof (int));
+ var.arr2 = (float *) calloc (N, sizeof (float));
+
+ {
+ #pragma omp declare mapper (struct A x) map(to: x.arr1) \
+ map(tofrom: x.arr1[0:N])
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ var.arr1[i]++;
+ }
+ }
+
+ {
+ #pragma omp declare mapper (struct A x) map(to: x.arr2) \
+ map(tofrom: x.arr2[0:N])
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ var.arr2[i]++;
+ }
+ }
+
+ {
+ #pragma omp declare mapper (struct A x) map(tofrom: x.arr3[0:N])
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ var.arr3[i]++;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (var.arr1[i] == 1);
+ assert (var.arr2[i] == 1);
+ assert (var.arr3[i] == 1);
+ }
+
+ free (var.arr1);
+ free (var.arr2);
+}
This is a new version of the patch to support OpenMP 5.0 "declare mapper" functionality for C++. As with the previously-posted version, arrays of structs whose elements would be mapped via a user-defined mapper remain unsupported. (Previous versions were posted here: https://gcc.gnu.org/pipermail/gcc-patches/2022-September/601560.html https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591983.html) This version of the patch uses a magic VAR_DECL instead of a magic FUNCTION_DECL for representing mappers, which simplifies parsing somewhat, and slightly reduces the number of places that need special-case handling in the FE. We use the DECL_INITIAL of the VAR_DECL to hold the OMP_DECLARE_MAPPER definition. To make types agree, we use the type of the object to be mapped for both the var decl and the OMP_DECLARE_MAPPER node itself. Hence the OMP_DECLARE_MAPPER looks like a magic constant of struct type in this particular case. The magic var decl can go in all the places that the "declare mapper" function decl went previously: at the top level of the program, within a class definition (including template classes), and within a function definition (including template functions). In the class case we conceptually use the C++-17-ism of definining the var decl "inline static", equivalent to e.g.: [template ...] class bla { static inline omp declare mapper ... = #define omp declare mapper ..." }; (though of course we don't restrict the "declare mapper"-in-class syntax to C++-17.) The new representation necessitates some changes to template instantiation. In particular, declare mappers may trigger implicitly, so we must make sure they are instantiated before they are needed (see changes to mark_used, etc.). I've rearranged the processing done by the gimplify_scan_omp_clauses and gimplify_adjust_omp_clauses functions so the order of the phases can remain intact in the presence of declared mappers. To do this, most gimplification of clauses in gimplify_scan_omp_clauses has been moved to gimplify_adjust_omp_clauses. This allows e.g. struct sibling-list handling and topological clause sorting to work with the non-gimplified form of clauses in the latter function -- including those that arise from mapper expansion. This seems to work well now. Relative to the last-posted version, this patch brings forward various refactoring that was previously done by the C and Fortran "declare mapper" support patches -- aiming to reduce churn. E.g. nested mapper finding and mapper instantiation has been moved to c-family/c-omp.cc so it can be shared between C and C++, and omp_name_type<T> in omp-general.h (used as the key to hash mapper definitions) is already templatized ready for Fortran support. This patch does not synthesize default mappers that map each of a struct's elements individually: whole-struct mappings are still done by copying the block of memory containing the struct. That works fine apart from cases where a struct has a member that is a reference (to a pointer). We could fix that by synthesizing a mapper for such cases (only), but that hasn't been attempted yet. (I think that means Jakub's concerns about blow-up of element mappings won't be a problem until that's done.) New tests added in {gcc,libgomp}/c-c++-common have been restricted to C++ for now, as the equivalent C parser changes to "declare mapper" support are still TBD (as are any Fortran FE adjustments that might be necessary due to various changes here and elsewhere). This patch depends on the in-review or approved-but-pending-other-patches patches: "OpenMP/OpenACC: Rework clause expansion and nested struct handling" https://gcc.gnu.org/pipermail/gcc-patches/2022-September/601837.html (aka https://gcc.gnu.org/pipermail/gcc-patches/2022-October/603792.html) [approved, but breaks Fortran mapping a bit without...] "OpenMP: Pointers and member mappings" https://gcc.gnu.org/pipermail/gcc-patches/2022-October/603794.html [...this one, which is unreviewed] "OpenMP/OpenACC: Unordered/non-constant component offset runtime diagnostic" https://gcc.gnu.org/pipermail/gcc-patches/2022-October/603793.html [unreviewed also -- "QoI" improvement to above] "OpenMP: lvalue parsing for map/to/from clauses (C++)" https://gcc.gnu.org/pipermail/gcc-patches/2022-November/605367.html [reviewed, but needs a little revision] 2022-11-30 Julian Brown <julian@codesourcery.com> gcc/c-family/ * c-common.h (omp_mapper_list): Add forward declaration. (c_omp_find_nested_mappers, c_omp_instantiate_mappers): Add prototypes. * c-omp.cc (c_omp_find_nested_mappers): New function. (remap_mapper_decl_info): New struct. (remap_mapper_decl_1, omp_instantiate_mapper, c_omp_instantiate_mappers): New functions. gcc/cp/ * constexpr.cc (reduced_constant_expression_p): Add OMP_DECLARE_MAPPER case. (cxx_eval_constant_expression, potential_constant_expression_1): Likewise. * cp-gimplify.cc (cxx_omp_finish_mapper_clauses): New function. * cp-objcp-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks. * cp-tree.h (lang_decl_base): Add omp_declare_mapper_p field. Recount spare bits comment. (DECL_OMP_DECLARE_MAPPER_P): New macro. (omp_mapper_id, cp_check_omp_declare_mapper, omp_instantiate_mappers, cxx_omp_finish_mapper_clauses, cxx_omp_mapper_lookup, cxx_omp_extract_mapper_directive, cxx_omp_map_array_section: Add prototypes. * decl.cc (check_initializer): Add OpenMP declare mapper support. (cp_finish_decl): Set DECL_INITIAL for OpenMP declare mapper var decls as appropriate. * decl2.cc (mark_used): Instantiate OpenMP "declare mapper" magic var decls. * error.cc (dump_omp_declare_mapper): New function. (dump_simple_decl): Use above. * parser.cc (cp_parser_omp_clause_map): Add KIND parameter. Support "mapper" modifier. (cp_parser_omp_all_clauses): Add KIND argument to cp_parser_omp_clause_map call. (cp_parser_omp_target): Call omp_instantiate_mappers before finish_omp_clauses. (cp_parser_omp_declare_mapper): New function. (cp_parser_omp_declare): Add "declare mapper" support. * pt.cc (tsubst_decl): Adjust name of "declare mapper" magic var decls once we know their type. (tsubst_omp_clauses): Call omp_instantiate_mappers before finish_omp_clauses, for target regions. (tsubst_expr): Support OMP_DECLARE_MAPPER nodes. (instantiate_decl): Instantiate initialiser (i.e definition) for OpenMP declare mappers. * semantics.cc (gimplify.h): Include. (omp_mapper_id, omp_mapper_lookup, omp_extract_mapper_directive, cxx_omp_map_array_section, cp_check_omp_declare_mapper): New functions. (finish_omp_clauses): Delete GOMP_MAP_PUSH_MAPPER_NAME and GOMP_MAP_POP_MAPPER_NAME artificial clauses. (omp_target_walk_data): Add MAPPERS field. (finish_omp_target_clauses_r): Scan for uses of struct/union/class type variables. (finish_omp_target_clauses): Create artificial mapper binding clauses for used structs/unions/classes in offload region. gcc/fortran/ * parse.cc (tree.h, fold-const.h, tree-hash-traits.h): Add includes (for additions to omp-general.h). gcc/ * gimplify.cc (gimplify_omp_ctx): Add IMPLICIT_MAPPERS field. (new_omp_context): Initialise IMPLICIT_MAPPERS hash map. (delete_omp_context): Delete IMPLICIT_MAPPERS hash map. (instantiate_mapper_info): New structs. (remap_mapper_decl_1, omp_mapper_copy_decl, omp_instantiate_mapper, omp_instantiate_implicit_mappers): New functions. (gimplify_scan_omp_clauses): Remove PREV_LIST_P tracking and use mapping groups instead. Remove most gimplification of OMP_CLAUSE_MAP nodes from here, but still populate ctx->variables splay tree. Handle MAPPER_BINDING clauses. (gimplify_adjust_omp_clauses): Move most gimplification of OMP clauses here. Instantiate implicit declared mappers. (gimplify_omp_declare_mapper): New function. (gimplify_expr): Call above function. * langhooks.def (lhd_omp_finish_mapper_clauses, lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): Add prototypes. (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define macros. (LANG_HOOK_DECLS): Add above macros. * langhooks.cc (lhd_omp_finish_mapper_clauses, lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): New dummy functions. * langhooks.h (lang_hooks_for_decls): Add OMP_FINISH_MAPPER_CLAUSES, OMP_MAPPER_LOOKUP, OMP_EXTRACT_MAPPER_DIRECTIVE, OMP_MAP_ARRAY_SECTION hooks. * omp-general.h (omp_name_type<T>): Add templatized struct, hash type traits (for omp_name_type<tree> specialization). (omp_mapper_list<T>): Add struct. * tree-core.h (omp_clause_code): Add OMP_CLAUSE__MAPPER_BINDING_. * tree-pretty-print.cc (dump_omp_clause): Support GOMP_MAP_UNSET, GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping clauses. Support OMP_CLAUSE__MAPPER_BINDING_ and OMP_DECLARE_MAPPER. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Add OMP_CLAUSE__MAPPER_BINDING_. * tree.def (OMP_DECLARE_MAPPER): New tree code. * tree.h (OMP_DECLARE_MAPPER_ID, OMP_DECLARE_MAPPER_DECL, OMP_DECLARE_MAPPER_CLAUSES): New defines. (OMP_CLAUSE__MAPPER_BINDING__ID, OMP_CLAUSE__MAPPER_BINDING__DECL, OMP_CLAUSE__MAPPER_BINDING__MAPPER): New defines. include/ * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_UNSET, GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping clause types. gcc/testsuite/ * c-c++-common/gomp/map-6.c: Update error scan output. * c-c++-common/gomp/declare-mapper-3.c: New test (only enabled for C++ for now). * c-c++-common/gomp/declare-mapper-4.c: Likewise. * c-c++-common/gomp/declare-mapper-5.c: Likewise. * c-c++-common/gomp/declare-mapper-6.c: Likewise. * c-c++-common/gomp/declare-mapper-7.c: Likewise. * c-c++-common/gomp/declare-mapper-8.c: Likewise. * c-c++-common/gomp/declare-mapper-9.c: Likewise. * c-c++-common/gomp/declare-mapper-12.c: Likewise. * g++.dg/gomp/declare-mapper-1.C: New test. * g++.dg/gomp/declare-mapper-2.C: New test. libgomp/ * testsuite/libgomp.c++/declare-mapper-1.C: New test. * testsuite/libgomp.c++/declare-mapper-2.C: New test. * testsuite/libgomp.c++/declare-mapper-3.C: New test. * testsuite/libgomp.c++/declare-mapper-4.C: New test. * testsuite/libgomp.c++/declare-mapper-5.C: New test. * testsuite/libgomp.c++/declare-mapper-6.C: New test. * testsuite/libgomp.c++/declare-mapper-7.C: New test. * testsuite/libgomp.c++/declare-mapper-8.C: New test. * testsuite/libgomp.c-c++-common/declare-mapper-9.c: New test (only enabled for C++ for now). * testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise. --- gcc/c-family/c-common.h | 3 + gcc/c-family/c-omp.cc | 300 ++++++ gcc/cp/constexpr.cc | 9 + gcc/cp/cp-gimplify.cc | 6 + gcc/cp/cp-objcp-common.h | 9 + gcc/cp/cp-tree.h | 17 +- gcc/cp/decl.cc | 27 +- gcc/cp/decl2.cc | 9 +- gcc/cp/error.cc | 25 + gcc/cp/parser.cc | 284 +++++- gcc/cp/pt.cc | 29 +- gcc/cp/semantics.cc | 188 +++- gcc/fortran/parse.cc | 3 + gcc/gimplify.cc | 931 ++++++++++++------ gcc/langhooks-def.h | 13 + gcc/langhooks.cc | 35 + gcc/langhooks.h | 16 + gcc/omp-general.h | 86 ++ .../c-c++-common/gomp/declare-mapper-12.c | 22 + .../c-c++-common/gomp/declare-mapper-3.c | 30 + .../c-c++-common/gomp/declare-mapper-4.c | 78 ++ .../c-c++-common/gomp/declare-mapper-5.c | 26 + .../c-c++-common/gomp/declare-mapper-6.c | 23 + .../c-c++-common/gomp/declare-mapper-7.c | 29 + .../c-c++-common/gomp/declare-mapper-8.c | 43 + .../c-c++-common/gomp/declare-mapper-9.c | 34 + gcc/testsuite/c-c++-common/gomp/map-6.c | 6 +- gcc/testsuite/g++.dg/gomp/declare-mapper-1.C | 58 ++ gcc/testsuite/g++.dg/gomp/declare-mapper-2.C | 30 + gcc/tree-core.h | 4 + gcc/tree-pretty-print.cc | 41 + gcc/tree.cc | 2 + gcc/tree.def | 7 + gcc/tree.h | 19 + include/gomp-constants.h | 8 +- .../testsuite/libgomp.c++/declare-mapper-1.C | 87 ++ .../testsuite/libgomp.c++/declare-mapper-2.C | 55 ++ .../testsuite/libgomp.c++/declare-mapper-3.C | 63 ++ .../testsuite/libgomp.c++/declare-mapper-4.C | 63 ++ .../testsuite/libgomp.c++/declare-mapper-5.C | 52 + .../testsuite/libgomp.c++/declare-mapper-6.C | 37 + .../testsuite/libgomp.c++/declare-mapper-7.C | 48 + .../testsuite/libgomp.c++/declare-mapper-8.C | 61 ++ .../libgomp.c-c++-common/declare-mapper-10.c | 60 ++ .../libgomp.c-c++-common/declare-mapper-11.c | 59 ++ .../libgomp.c-c++-common/declare-mapper-12.c | 87 ++ .../libgomp.c-c++-common/declare-mapper-13.c | 55 ++ .../libgomp.c-c++-common/declare-mapper-14.c | 57 ++ .../libgomp.c-c++-common/declare-mapper-9.c | 62 ++ 49 files changed, 2993 insertions(+), 303 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-4.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-5.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-6.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-7.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-8.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-9.c create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-2.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-1.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-2.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-3.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-4.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-5.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-6.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-7.C create mode 100644 libgomp/testsuite/libgomp.c++/declare-mapper-8.C create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-10.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-11.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-12.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-13.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-14.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-9.c