@@ -1275,8 +1275,12 @@ enum c_omp_region_type
C_ORT_ACC = 1 << 1,
C_ORT_DECLARE_SIMD = 1 << 2,
C_ORT_TARGET = 1 << 3,
+ C_ORT_EXIT_DATA = 1 << 4,
+ C_ORT_DECLARE_MAPPER = 1 << 6,
C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD,
C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET,
+ C_ORT_OMP_EXIT_DATA = C_ORT_OMP | C_ORT_EXIT_DATA,
+ C_ORT_OMP_DECLARE_MAPPER = C_ORT_OMP | C_ORT_DECLARE_MAPPER,
C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET
};
@@ -1312,6 +1316,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, enum c_omp_region_type);
namespace omp_addr_tokenizer { struct omp_addr_token; }
typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;
@@ -3938,6 +3938,483 @@ 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;
+}
+
+static enum gomp_map_kind
+omp_split_map_kind (enum gomp_map_kind op, bool *force_p, bool *always_p,
+ bool *present_p)
+{
+ *force_p = *always_p = *present_p = false;
+
+ switch (op)
+ {
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TOFROM:
+ case GOMP_MAP_FORCE_PRESENT:
+ *force_p = true;
+ break;
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ *always_p = true;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ *always_p = true;
+ /* Fallthrough. */
+ case GOMP_MAP_PRESENT_ALLOC:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ *present_p = true;
+ break;
+ default:
+ ;
+ }
+
+ switch (op)
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
+ return GOMP_MAP_ALLOC;
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ return GOMP_MAP_TO;
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ return GOMP_MAP_FROM;
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_FORCE_TOFROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ return GOMP_MAP_TOFROM;
+ default:
+ ;
+ }
+
+ return op;
+}
+
+static enum gomp_map_kind
+omp_join_map_kind (enum gomp_map_kind op, bool force_p, bool always_p,
+ bool present_p)
+{
+ gcc_assert (!force_p || !(always_p || present_p));
+
+ switch (op)
+ {
+ case GOMP_MAP_ALLOC:
+ if (force_p)
+ return GOMP_MAP_FORCE_ALLOC;
+ else if (present_p)
+ return GOMP_MAP_PRESENT_ALLOC;
+ break;
+
+ case GOMP_MAP_TO:
+ if (force_p)
+ return GOMP_MAP_FORCE_TO;
+ else if (always_p && present_p)
+ return GOMP_MAP_ALWAYS_PRESENT_TO;
+ else if (always_p)
+ return GOMP_MAP_ALWAYS_TO;
+ else if (present_p)
+ return GOMP_MAP_PRESENT_TO;
+ break;
+
+ case GOMP_MAP_FROM:
+ if (force_p)
+ return GOMP_MAP_FORCE_FROM;
+ else if (always_p && present_p)
+ return GOMP_MAP_ALWAYS_PRESENT_FROM;
+ else if (always_p)
+ return GOMP_MAP_ALWAYS_FROM;
+ else if (present_p)
+ return GOMP_MAP_PRESENT_FROM;
+ break;
+
+ case GOMP_MAP_TOFROM:
+ if (force_p)
+ return GOMP_MAP_FORCE_TOFROM;
+ else if (always_p && present_p)
+ return GOMP_MAP_ALWAYS_PRESENT_TOFROM;
+ else if (always_p)
+ return GOMP_MAP_ALWAYS_TOFROM;
+ else if (present_p)
+ return GOMP_MAP_PRESENT_TOFROM;
+ break;
+
+ default:
+ ;
+ }
+
+ return op;
+}
+
+/* Map kind decay (OpenMP 5.2, 5.8.8 "declare mapper Directive"). Return the
+ map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS
+ specified on the clause that invokes the mapper. See also
+ fortran/trans-openmp.cc:omp_map_decayed_kind. */
+
+static enum gomp_map_kind
+omp_map_decayed_kind (enum gomp_map_kind mapper_kind,
+ enum gomp_map_kind invoked_as, bool exit_p)
+{
+ if (invoked_as == GOMP_MAP_RELEASE || invoked_as == GOMP_MAP_DELETE)
+ return invoked_as;
+
+ bool force_p, always_p, present_p;
+
+ invoked_as = omp_split_map_kind (invoked_as, &force_p, &always_p, &present_p);
+ gomp_map_kind decay_to;
+
+ switch (mapper_kind)
+ {
+ case GOMP_MAP_ALLOC:
+ if (exit_p && invoked_as == GOMP_MAP_FROM)
+ decay_to = GOMP_MAP_RELEASE;
+ else
+ decay_to = GOMP_MAP_ALLOC;
+ break;
+
+ case GOMP_MAP_TO:
+ if (invoked_as == GOMP_MAP_FROM)
+ decay_to = exit_p ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
+ else if (invoked_as == GOMP_MAP_ALLOC)
+ decay_to = GOMP_MAP_ALLOC;
+ else
+ decay_to = GOMP_MAP_TO;
+ break;
+
+ case GOMP_MAP_FROM:
+ if (invoked_as == GOMP_MAP_ALLOC || invoked_as == GOMP_MAP_TO)
+ decay_to = GOMP_MAP_ALLOC;
+ else
+ decay_to = GOMP_MAP_FROM;
+ break;
+
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_UNSET:
+ decay_to = invoked_as;
+ break;
+
+ default:
+ gcc_unreachable ();
+ }
+
+ return omp_join_map_kind (decay_to, force_p, always_p, present_p);
+}
+
+/* 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 (location_t loc, tree *outlist, tree mapper, tree expr,
+ enum gomp_map_kind outer_kind,
+ enum c_omp_region_type ort)
+{
+ 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)
+ {
+ 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);
+
+ OMP_CLAUSE_LOCATION (unshared) = loc;
+
+ enum gomp_map_kind decayed_kind
+ = omp_map_decayed_kind (clause_kind, outer_kind,
+ (ort & C_ORT_EXIT_DATA) != 0);
+ OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_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)
+ {
+ outlist = omp_instantiate_mapper (loc, outlist, nested_mapper,
+ t, outer_kind, ort);
+ 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, enum c_omp_region_type ort)
+{
+ 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 (OMP_CLAUSE_LOCATION (c),
+ pc, mapper, t, kind, ort);
+ 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. */
@@ -3300,6 +3300,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;
@@ -7160,6 +7163,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:
@@ -9725,6 +9729,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);
@@ -2405,6 +2405,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
@@ -2878,7 +2878,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. */
@@ -4369,6 +4372,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 VAR_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) \
@@ -7750,10 +7758,13 @@ extern tree finish_qualified_id_expr (tree, tree, bool, bool,
extern void simplify_aggr_init_expr (tree *);
extern void finalize_nrv (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> &);
@@ -8338,6 +8349,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))
{
@@ -8604,14 +8610,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);
@@ -5821,10 +5821,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_context)
@@ -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 (VAR_P (t) && 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))
@@ -40726,13 +40726,12 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
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;
@@ -40750,12 +40749,17 @@ 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 present_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);
@@ -40778,6 +40782,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)
{
@@ -40791,6 +40796,71 @@ 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 if (strcmp ("present", p) == 0)
{
@@ -40804,19 +40874,19 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
return list;
}
present_modifier = true;
- }
+ cp_lexer_consume_token (parser->lexer);
+ }
else
{
- cp_parser_error (parser, "%<map%> clause with map-type modifier other"
- " than %<always%>, %<close%> or %<present%>");
+ cp_parser_error (parser, "%<map%> clause with map-type modifier "
+ "other than %<always%>, %<close%>, "
+ "%<mapper%> or %<present%>");
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)
@@ -40872,8 +40942,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;
}
@@ -41684,7 +41776,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:
@@ -45472,7 +45564,10 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
- "#pragma omp target data", pragma_tok);
+ "#pragma omp target data", pragma_tok, false);
+ if (!processing_template_decl)
+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+ clauses = finish_omp_clauses (clauses, C_ORT_OMP);
c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
@@ -45586,7 +45681,11 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
- "#pragma omp target enter data", pragma_tok);
+ "#pragma omp target enter data", pragma_tok,
+ false);
+ if (!processing_template_decl)
+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);
+ clauses = finish_omp_clauses (clauses, C_ORT_OMP);
c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
@@ -45701,7 +45800,11 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
- "#pragma omp target exit data", pragma_tok);
+ "#pragma omp target exit data", pragma_tok,
+ false);
+ if (!processing_template_decl)
+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA);
+ clauses = finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
@@ -45993,6 +46096,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, C_ORT_OMP_TARGET);
clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
c_omp_adjust_map_clauses (clauses, true);
@@ -48251,6 +48356,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) \
@@ -48291,6 +48562,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);
@@ -48304,7 +48581,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;
}
@@ -15479,6 +15479,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;
@@ -18312,8 +18319,10 @@ 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_DECLARE_SIMD && ort != C_ORT_OMP_DECLARE_MAPPER)
{
+ if (ort & C_ORT_OMP)
+ new_clauses = c_omp_instantiate_mappers (new_clauses, ort);
new_clauses = finish_omp_clauses (new_clauses, ort);
if (linear_no_step)
for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
@@ -19732,7 +19741,9 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl)
case OMP_TARGET_UPDATE:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
- tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), C_ORT_OMP, args,
+ tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t),
+ (TREE_CODE (t) == OMP_TARGET_EXIT_DATA
+ ? C_ORT_OMP_EXIT_DATA : C_ORT_OMP), args,
complain, in_decl);
t = copy_node (t);
OMP_STANDALONE_CLAUSES (t) = tmp;
@@ -19880,6 +19891,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_MAPPER, 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;
@@ -27248,7 +27275,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
@@ -6077,6 +6078,101 @@ 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);
+
+ if (TYPE_PTR_P (TREE_TYPE (t)))
+ t = build_array_ref (loc, t, low);
+ else
+ t = error_mark_node;
+ }
+
+ return t;
+}
+
/* Helper function for cp_parser_omp_declare_reduction_exprs
and tsubst_omp_udr.
Remove CLEANUP_STMT for data (omp_priv variable).
@@ -6558,6 +6654,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. */
@@ -8098,6 +8217,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:
@@ -9550,6 +9675,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
@@ -9563,6 +9690,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. */
@@ -9576,6 +9704,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);
@@ -9678,10 +9827,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);
}
@@ -11679,6 +11682,218 @@ 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 we already have clauses pertaining to a struct variable, then we don't
+ want to implicitly invoke a user-defined mapper. */
+ if ((flags & GOVD_EXPLICIT) != 0 && AGGREGATE_TYPE_P (TREE_TYPE (decl)))
+ 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 (mapper_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. */
@@ -12390,6 +12605,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;
@@ -13421,6 +13647,30 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
|| 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_no_se, omp_mapping_group *> *grpmap = NULL;
@@ -17200,6 +17450,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
@@ -18122,6 +18381,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,59 @@
+/* { dg-do compile { target c++ } } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+typedef struct {
+ int a, b, c, d;
+} S;
+
+int main ()
+{
+ S s;
+ #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \
+ map(from: x.c) map(tofrom: x.d)
+
+ #pragma omp target enter data map(to: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(to:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+ #pragma omp target exit data map(from: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: 4\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(from:s\.c \[len: [0-9]+\]\) map\(from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+
+ #pragma omp target enter data map(alloc: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(alloc:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(alloc:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+ #pragma omp target exit data map(release: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(release:s\.c \[len: [0-9]+\]\) map\(release:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+
+ #pragma omp target enter data map(present, to: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(force_present:s\.a \[len: [0-9]+\]\) map\(force_present:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(force_present:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+ #pragma omp target exit data map(present, from: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(force_present:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+
+ #pragma omp target enter data map(always, to: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(always,to:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(always,to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+ #pragma omp target exit data map(always, from: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(always,from:s\.c \[len: [0-9]+\]\) map\(always,from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+
+ #pragma omp target enter data map(always, present, to: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(force_present:s\.a \[len: [0-9]+\]\) map\(always,present,to:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(always,present,to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+ #pragma omp target exit data map(always, present, from: s)
+
+ /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(always,present,from:s\.c \[len: [0-9]+\]\) map\(always,present,from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,39 @@
+/* { dg-do compile { target c++ } } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+typedef struct {
+ int a, b, c, d;
+} S;
+
+int main ()
+{
+ S s = { 0, 0, 0, 0 };
+ #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \
+ map(from: x.c) map(tofrom: x.d)
+
+ #pragma omp target data map(s)
+ /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(to:s\.b \[len: [0-9]+\]\) map\(from:s\.c \[len: [0-9]+\]\) map\(tofrom:s\.d \[len: [0-9]+\]\)} 3 "gimple" } } */
+ {
+ #pragma omp target
+ {
+ s.a++;
+ s.b++;
+ s.c++;
+ s.d++;
+ }
+ }
+
+ #pragma omp target data map(alloc: s)
+ /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(alloc:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(alloc:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */
+ {
+ #pragma omp target
+ {
+ s.a++;
+ s.b++;
+ s.c++;
+ s.d++;
+ }
+ }
+
+ return 0;
+}
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,19 +13,19 @@ foo (void)
#pragma omp target map (to:a)
;
- #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */
;
- #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */
;
- #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */
;
- #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */
;
- #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */
;
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\(alloc: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;
+}
@@ -347,6 +347,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_,
@@ -1015,6 +1015,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
pp_string (pp, "always,present,tofrom");
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 ();
}
@@ -1082,6 +1091,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))
@@ -3851,6 +3877,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]]");
@@ -269,6 +269,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 */
@@ -360,6 +361,7 @@ const char * const omp_clause_code_name[] =
"map",
"has_device_addr",
"doacross",
+ "_mapper_binding_",
"_cache_",
"gang",
"async",
@@ -1291,6 +1291,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). */
@@ -1572,6 +1572,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)
@@ -2019,6 +2026,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
@@ -209,7 +209,13 @@ enum gomp_map_kind
GOMP_MAP_PRESENT_ALLOC = (GOMP_MAP_LAST | 4),
GOMP_MAP_PRESENT_TO = (GOMP_MAP_LAST | 5),
GOMP_MAP_PRESENT_FROM = (GOMP_MAP_LAST | 6),
- GOMP_MAP_PRESENT_TOFROM = (GOMP_MAP_LAST | 7)
+ GOMP_MAP_PRESENT_TOFROM = (GOMP_MAP_LAST | 7),
+ /* 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 | 8),
+ /* Used to record the name of a named mapper. */
+ GOMP_MAP_PUSH_MAPPER_NAME = (GOMP_MAP_LAST | 9),
+ GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 10)
};
#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 patch adds support for OpenMP 5.0 "declare mapper" functionality for C++. This is based on the version of the patch posted for the og13 branch here: https://gcc.gnu.org/pipermail/gcc-patches/2023-June/623353.html The following follow up patches/fixes have also been incorporated into this version: "OpenMP: Expand "declare mapper" mappers for target {enter,exit,} data directives": https://gcc.gnu.org/pipermail/gcc-patches/2023-July/623780.html "OpenMP: Introduce C_ORT_{,OMP_}DECLARE_MAPPER c_omp_region_type types": https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627005.html 2023-09-05 Julian Brown <julian@codesourcery.com> gcc/c-family/ * c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA, C_ORT_DECLARE_MAPPER, C_ORT_OMP_EXIT_DATA, C_ORT_OMP_DECLARE_MAPPER values. (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_split_map_kind, omp_join_map_kind, omp_map_decayed_kind, 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, cxx_omp_mapper_lookup, cxx_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): Handle MAPPER_BINDING clauses. (gimplify_adjust_omp_clauses): Instantiate implicit declared mappers. (gimplify_omp_declare_mapper): New function. (gimplify_expr): Call above function. * langhooks-def.h (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. * c-c++-common/gomp/declare-mapper-15.c: Likewise. * c-c++-common/gomp/declare-mapper-16.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 | 7 + gcc/c-family/c-omp.cc | 477 ++++++++++++++++++ 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 | 305 ++++++++++- gcc/cp/pt.cc | 35 +- gcc/cp/semantics.cc | 191 ++++++- gcc/fortran/parse.cc | 3 + gcc/gimplify.cc | 263 ++++++++++ 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-15.c | 59 +++ .../c-c++-common/gomp/declare-mapper-16.c | 39 ++ .../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 | 10 +- 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 +++ 51 files changed, 2907 insertions(+), 34 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-15.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-16.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