@@ -1281,6 +1281,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;
@@ -3939,6 +3939,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. */
@@ -3273,6 +3273,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;
@@ -7088,6 +7091,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:
@@ -9499,6 +9503,11 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now,
"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, fundef_p);
@@ -2319,6 +2319,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
@@ -184,6 +184,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
@@ -2882,7 +2882,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. */
@@ -4356,6 +4359,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) \
@@ -7716,10 +7724,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> &);
@@ -8279,6 +8290,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);
@@ -7441,6 +7441,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))
{
@@ -8584,14 +8590,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);
@@ -5785,10 +5785,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 ())
@@ -1137,12 +1137,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))
@@ -40551,13 +40551,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;
@@ -40575,11 +40574,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);
@@ -40602,6 +40606,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)
{
@@ -40615,20 +40620,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)
@@ -40674,8 +40742,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;
}
@@ -41486,7 +41576,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:
@@ -45654,6 +45744,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);
@@ -47902,6 +47994,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) \
@@ -47942,6 +48200,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);
@@ -47955,7 +48219,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;
}
@@ -15320,6 +15320,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;
@@ -18146,6 +18153,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))
@@ -19698,6 +19707,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;
@@ -26949,7 +26974,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
@@ -5972,6 +5973,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).
@@ -6453,6 +6546,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. */
@@ -7993,6 +8109,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:
@@ -9442,6 +9564,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
@@ -9455,6 +9579,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. */
@@ -9468,6 +9593,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);
@@ -9570,10 +9716,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);
}
@@ -11539,6 +11542,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. */
@@ -11550,7 +11774,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;
@@ -11581,75 +11805,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;
- bool enter_exit = (code == OMP_TARGET_ENTER_DATA
- || code == OMP_TARGET_EXIT_DATA);
-
- /* 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, enter_exit);
- 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:
@@ -11954,99 +12133,37 @@ 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. */
-
- tree desc;
- if (DECL_P (decl)
- || (prev_list_p
- && (desc = OMP_CLAUSE_CHAIN (*prev_list_p))
- && OMP_CLAUSE_CODE (desc) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_RELEASE
- || OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_DELETE)
- && DECL_P (OMP_CLAUSE_DECL (desc))))
- {
- /* We only do the transformations here for non-component
- mappings. Array descriptors for component mappings are
- handled in omp_accumulate_sibling_list. */
- if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
- {
- enum gomp_map_kind k
- = (OMP_CLAUSE_MAP_KIND (*prev_list_p) == GOMP_MAP_DELETE
- ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
- OMP_CLAUSE_SET_MAP_KIND (c, k);
- }
- 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_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
- }
+ 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 (remove)
break;
@@ -12066,41 +12183,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
@@ -12108,14 +12235,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)
@@ -12138,56 +12258,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))
@@ -12298,27 +12382,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;
@@ -12326,43 +12389,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:
@@ -12451,6 +12477,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;
@@ -12998,6 +13035,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;
}
@@ -13465,15 +13508,99 @@ 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);
+
+ bool enter_exit = (code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA);
+
+ outlist = omp_tsort_mapping_groups (groups, grpmap, enter_exit);
+ 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:
@@ -13628,26 +13755,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. */
@@ -13684,7 +13852,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)
@@ -13707,8 +13922,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)
@@ -13783,6 +14113,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:
@@ -16990,6 +17335,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
@@ -17912,6 +18266,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))
@@ -3825,6 +3851,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]]");
@@ -271,6 +271,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 */
@@ -362,6 +363,7 @@ const char * const omp_clause_code_name[] =
"map",
"has_device_addr",
"doacross",
+ "_mapper_binding_",
"_cache_",
"gang",
"async",
@@ -1250,6 +1250,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). */
@@ -1536,6 +1536,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)
@@ -1976,6 +1983,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. 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 -- 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. 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 | 958 ++++++++++++------ 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, 2998 insertions(+), 325 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