@@ -1252,6 +1252,9 @@ extern tree c_omp_check_context_selector (location_t, tree);
extern void c_omp_mark_declare_variant (location_t, tree, tree);
extern const char *c_omp_map_clause_name (tree, bool);
extern void c_omp_adjust_map_clauses (tree, bool);
+struct omp_mapper_list;
+extern void c_omp_find_nested_mappers (struct omp_mapper_list *, tree);
+extern tree c_omp_instantiate_mappers (tree);
class c_omp_address_inspector
{
@@ -3396,6 +3396,306 @@ c_omp_address_inspector::get_attachment_point (tree expr)
return get_origin (baseptr);
}
+/* 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 *mlist, tree mapper_fn)
+{
+ tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
+ tree mapper_name = NULL_TREE;
+
+ if (mapper == error_mark_node)
+ return;
+
+ gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+ for (tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+ clause;
+ clause = OMP_CLAUSE_CHAIN (clause))
+ {
+ tree expr = OMP_CLAUSE_DECL (clause);
+ enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause);
+ tree elem_type;
+
+ if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+ {
+ mapper_name = expr;
+ continue;
+ }
+ else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ mapper_name = NULL_TREE;
+ continue;
+ }
+
+ gcc_assert (TREE_CODE (expr) != TREE_LIST);
+ if (TREE_CODE (expr) == OMP_ARRAY_SECTION)
+ {
+ while (TREE_CODE (expr) == OMP_ARRAY_SECTION)
+ expr = TREE_OPERAND (expr, 0);
+
+ elem_type = TREE_TYPE (expr);
+ }
+ else
+ elem_type = TREE_TYPE (expr);
+
+ /* This might be too much... or not enough? */
+ while (TREE_CODE (elem_type) == ARRAY_TYPE
+ || TREE_CODE (elem_type) == POINTER_TYPE
+ || TREE_CODE (elem_type) == REFERENCE_TYPE)
+ elem_type = TREE_TYPE (elem_type);
+
+ elem_type = TYPE_MAIN_VARIANT (elem_type);
+
+ if (AGGREGATE_TYPE_P (elem_type)
+ && !mlist->contains (mapper_name, elem_type))
+ {
+ tree nested_mapper_fn
+ = lang_hooks.decls.omp_mapper_lookup (mapper_name, elem_type);
+
+ if (nested_mapper_fn)
+ {
+ mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn);
+ c_omp_find_nested_mappers (mlist, nested_mapper_fn);
+ }
+ else if (mapper_name)
+ {
+ error ("mapper %qE not found for type %qT", mapper_name,
+ elem_type);
+ continue;
+ }
+ }
+ }
+}
+
+struct remap_mapper_decl_info
+{
+ tree dummy_var;
+ tree expr;
+};
+
+/* Helper for rewriting DUMMY_VAR into EXPR in a map clause decl. */
+
+static tree
+remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
+{
+ remap_mapper_decl_info *map_info = (remap_mapper_decl_info *) data;
+
+ if (operand_equal_p (*tp, map_info->dummy_var))
+ {
+ *tp = map_info->expr;
+ *walk_subtrees = 0;
+ }
+
+ return NULL_TREE;
+}
+
+/* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to
+ OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in
+ the mapper declaration. */
+
+static tree *
+omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
+ enum gomp_map_kind outer_kind)
+{
+ tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+ tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
+ tree mapper_name = NULL_TREE;
+
+ remap_mapper_decl_info map_info;
+ map_info.dummy_var = dummy_var;
+ map_info.expr = expr;
+
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ tree unshared = unshare_expr (c);
+ enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c);
+ tree t = OMP_CLAUSE_DECL (unshared);
+ tree type = NULL_TREE;
+ bool nonunit_array_with_mapper = false;
+
+ if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+ {
+ mapper_name = t;
+ continue;
+ }
+ else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ mapper_name = NULL_TREE;
+ continue;
+ }
+
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (c);
+ tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
+
+ if (t2 == t)
+ {
+ nonunit_array_with_mapper = true;
+ /* We'd want use the mapper for the element type if this worked:
+ look that one up. */
+ type = TREE_TYPE (TREE_TYPE (t));
+ }
+ else
+ {
+ t = t2;
+ type = TREE_TYPE (t);
+ }
+ }
+ else
+ type = TREE_TYPE (t);
+
+ gcc_assert (type);
+
+ if (type == error_mark_node)
+ continue;
+
+ walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);
+
+ if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)
+ OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
+
+ type = TYPE_MAIN_VARIANT (type);
+
+ tree mapper_fn = lang_hooks.decls.omp_mapper_lookup (mapper_name, type);
+
+ if (mapper_fn && nonunit_array_with_mapper)
+ {
+ sorry ("user-defined mapper with non-unit length array section");
+ continue;
+ }
+ else if (mapper_fn)
+ {
+ tree nested_mapper
+ = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
+ if (nested_mapper != mapper)
+ {
+ if (clause_kind == GOMP_MAP_UNSET)
+ clause_kind = outer_kind;
+
+ outlist = omp_instantiate_mapper (outlist, nested_mapper,
+ t, clause_kind);
+ continue;
+ }
+ }
+ else if (mapper_name)
+ {
+ error ("mapper %qE not found for type %qT", mapper_name, type);
+ continue;
+ }
+
+ *outlist = unshared;
+ outlist = &OMP_CLAUSE_CHAIN (unshared);
+ }
+
+ return outlist;
+}
+
+/* Given a list of CLAUSES, scan each clause and invoke a user-defined mapper
+ appropriate to the type of the data in that clause, if such a mapper is
+ visible in the current parsing context. */
+
+tree
+c_omp_instantiate_mappers (tree clauses)
+{
+ tree c, *pc, mapper_name = NULL_TREE;
+
+ for (pc = &clauses, c = clauses; c; c = *pc)
+ {
+ bool using_mapper = false;
+
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_MAP:
+ {
+ tree t = OMP_CLAUSE_DECL (c);
+ tree type = NULL_TREE;
+ bool nonunit_array_with_mapper = false;
+
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME)
+ mapper_name = OMP_CLAUSE_DECL (c);
+ else
+ mapper_name = NULL_TREE;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
+ if (TREE_CODE (t) == OMP_ARRAY_SECTION)
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (c);
+ tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
+
+ if (t2 == t)
+ {
+ /* !!! Array sections of size >1 with mappers for elements
+ are hard to support. Do something here. */
+ nonunit_array_with_mapper = true;
+ type = TREE_TYPE (TREE_TYPE (t));
+ }
+ else
+ {
+ t = t2;
+ type = TREE_TYPE (t);
+ }
+ }
+ else
+ type = TREE_TYPE (t);
+
+ if (type == NULL_TREE || type == error_mark_node)
+ {
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
+ enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c);
+ if (kind == GOMP_MAP_UNSET)
+ kind = GOMP_MAP_TOFROM;
+
+ type = TYPE_MAIN_VARIANT (type);
+
+ tree mapper_fn
+ = lang_hooks.decls.omp_mapper_lookup (mapper_name, type);
+
+ if (mapper_fn && nonunit_array_with_mapper)
+ {
+ sorry ("user-defined mapper with non-unit length "
+ "array section");
+ using_mapper = true;
+ }
+ else if (mapper_fn)
+ {
+ tree mapper
+ = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
+ pc = omp_instantiate_mapper (pc, mapper, t, kind);
+ using_mapper = true;
+ }
+ else if (mapper_name)
+ {
+ error ("mapper %qE not found for type %qT", mapper_name, type);
+ using_mapper = true;
+ }
+ }
+ break;
+
+ default:
+ ;
+ }
+
+ if (using_mapper)
+ *pc = OMP_CLAUSE_CHAIN (c);
+ else
+ pc = &OMP_CLAUSE_CHAIN (c);
+ }
+
+ return clauses;
+}
+
static const struct c_omp_directive omp_directives[] = {
/* Keep this alphabetically sorted by the first word. Non-null second/third
if any should precede null ones. */
@@ -12458,6 +12458,175 @@ c_check_omp_declare_reduction_r (tree *tp, int *, void *data)
return NULL_TREE;
}
+/* Return identifier to look up for omp declare reduction. */
+
+tree
+c_omp_mapper_id (tree mapper_id)
+{
+ const char *p = NULL;
+
+ const char prefix[] = "omp declare mapper ";
+
+ if (mapper_id == NULL_TREE)
+ p = "<default>";
+ else if (TREE_CODE (mapper_id) == IDENTIFIER_NODE)
+ p = IDENTIFIER_POINTER (mapper_id);
+ else
+ return error_mark_node;
+
+ size_t lenp = sizeof (prefix);
+ size_t len = strlen (p);
+ char *name = XALLOCAVEC (char, lenp + len);
+ memcpy (name, prefix, lenp - 1);
+ memcpy (name + lenp - 1, p, len + 1);
+ return get_identifier (name);
+}
+
+/* Lookup MAPPER_ID in the current scope, or create an artificial
+ VAR_DECL, bind it into the current scope and return it. */
+
+tree
+c_omp_mapper_decl (tree mapper_id)
+{
+ struct c_binding *b = I_SYMBOL_BINDING (mapper_id);
+ if (b != NULL && B_IN_CURRENT_SCOPE (b))
+ return b->decl;
+
+ tree decl = build_decl (BUILTINS_LOCATION, VAR_DECL,
+ mapper_id, integer_type_node);
+ DECL_ARTIFICIAL (decl) = 1;
+ DECL_EXTERNAL (decl) = 1;
+ TREE_STATIC (decl) = 1;
+ TREE_PUBLIC (decl) = 0;
+ bind (mapper_id, decl, current_scope, true, false, BUILTINS_LOCATION);
+ return decl;
+}
+
+/* Lookup MAPPER_ID in the first scope where it has entry for TYPE. */
+
+tree
+c_omp_mapper_lookup (tree mapper_id, tree type)
+{
+ if (TREE_CODE (type) != RECORD_TYPE
+ && TREE_CODE (type) != UNION_TYPE)
+ return NULL_TREE;
+
+ mapper_id = c_omp_mapper_id (mapper_id);
+
+ struct c_binding *b = I_SYMBOL_BINDING (mapper_id);
+ while (b)
+ {
+ tree t;
+ for (t = DECL_INITIAL (b->decl); t; t = TREE_CHAIN (t))
+ if (comptypes (TREE_PURPOSE (t), type))
+ return TREE_VALUE (t);
+ b = b->shadowed;
+ }
+ return NULL_TREE;
+}
+
+/* For C, we record a pointer to the mapper itself without wrapping it in an
+ artificial function or similar. So, just return it. */
+
+tree
+c_omp_extract_mapper_directive (tree mapper)
+{
+ return mapper;
+}
+
+/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but
+ nothing more complicated. */
+
+tree
+c_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;
+
+ t = build_array_ref (loc, t, low);
+ }
+
+ return t;
+}
+
+/* Helper function for below function. */
+
+static tree
+c_omp_scan_mapper_bindings_r (tree *tp, int *walk_subtrees, void *ptr)
+{
+ tree t = *tp;
+ omp_mapper_list *mlist = (omp_mapper_list *) ptr;
+ tree aggr_type = NULL_TREE;
+
+ if (TREE_CODE (t) == SIZEOF_EXPR
+ || TREE_CODE (t) == ALIGNOF_EXPR)
+ {
+ *walk_subtrees = 0;
+ return NULL_TREE;
+ }
+
+ if (TREE_CODE (t) == OMP_CLAUSE)
+ return 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 = c_omp_mapper_lookup (NULL_TREE, aggr_type);
+ if (mapper_fn)
+ mlist->add_mapper (NULL_TREE, aggr_type, mapper_fn);
+ }
+
+ return NULL_TREE;
+}
+
+/* Scan an offload region's body, and record uses of struct- or union-typed
+ variables. Add _mapper_binding_ fake clauses to *CLAUSES_PTR. */
+
+void
+c_omp_scan_mapper_bindings (location_t loc, tree *clauses_ptr, tree body)
+{
+ hash_set<omp_name_type> seen_types;
+ auto_vec<tree> mappers;
+ omp_mapper_list mlist (&seen_types, &mappers);
+
+ walk_tree_without_duplicates (&body, c_omp_scan_mapper_bindings_r, &mlist);
+
+ unsigned int i;
+ tree mapper;
+ FOR_EACH_VEC_ELT (mappers, i, mapper)
+ c_omp_find_nested_mappers (&mlist, mapper);
+
+ FOR_EACH_VEC_ELT (mappers, i, mapper)
+ {
+ if (mapper == error_mark_node)
+ continue;
+ tree mapper_name = OMP_DECLARE_MAPPER_ID (mapper);
+ tree decl = OMP_DECLARE_MAPPER_DECL (mapper);
+
+ 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;
+
+ OMP_CLAUSE_CHAIN (c) = *clauses_ptr;
+ *clauses_ptr = c;
+ }
+}
bool
c_check_in_current_scope (tree decl)
@@ -122,6 +122,18 @@ along with GCC; see the file COPYING3. If not see
#undef LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP
#define LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP c_omp_clause_copy_ctor
+#undef LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES
+#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES c_omp_finish_mapper_clauses
+
+#undef LANG_HOOKS_OMP_MAPPER_LOOKUP
+#define LANG_HOOKS_OMP_MAPPER_LOOKUP c_omp_mapper_lookup
+
+#undef LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE
+#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE c_omp_extract_mapper_directive
+
+#undef LANG_HOOKS_OMP_MAP_ARRAY_SECTION
+#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION c_omp_map_array_section
+
#undef LANG_HOOKS_TREE_INLINING_VAR_MOD_TYPE_P
#define LANG_HOOKS_TREE_INLINING_VAR_MOD_TYPE_P c_vla_unspec_p
#endif /* GCC_C_OBJC_COMMON */
@@ -16263,10 +16263,9 @@ c_parser_omp_clause_depend (c_parser *parser, tree list)
always | close */
static tree
-c_parser_omp_clause_map (c_parser *parser, tree list)
+c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind)
{
location_t clause_loc = c_parser_peek_token (parser)->location;
- enum gomp_map_kind kind = GOMP_MAP_TOFROM;
tree nl, c;
matching_parens parens;
@@ -16285,11 +16284,27 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
pos++;
+ else if ((c_parser_peek_nth_token_raw (parser, pos + 1)->type
+ == CPP_OPEN_PAREN)
+ && ((c_parser_peek_nth_token_raw (parser, pos + 2)->type
+ == CPP_NAME)
+ || ((c_parser_peek_nth_token_raw (parser, pos + 2)->type
+ == CPP_KEYWORD)
+ && (c_parser_peek_nth_token_raw (parser,
+ pos + 2)->keyword
+ == RID_DEFAULT)))
+ && (c_parser_peek_nth_token_raw (parser, pos + 3)->type
+ == CPP_CLOSE_PAREN)
+ && (c_parser_peek_nth_token_raw (parser, pos + 4)->type
+ == CPP_COMMA))
+ pos += 4;
pos++;
}
int always_modifier = 0;
int close_modifier = 0;
+ int mapper_modifier = 0;
+ tree mapper_name = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
c_token *tok = c_parser_peek_token (parser);
@@ -16310,6 +16325,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
return list;
}
always_modifier++;
+ c_parser_consume_token (parser);
}
else if (strcmp ("close", p) == 0)
{
@@ -16320,6 +16336,60 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
return list;
}
close_modifier++;
+ c_parser_consume_token (parser);
+ }
+ else if (strcmp ("mapper", p) == 0)
+ {
+ c_parser_consume_token (parser);
+
+ matching_parens mparens;
+ if (mparens.require_open (parser))
+ {
+ if (mapper_modifier)
+ {
+ c_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. */
+ c_parser_consume_token (parser);
+ mparens.require_close (parser);
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+
+ tok = c_parser_peek_token (parser);
+
+ switch (tok->type)
+ {
+ case CPP_NAME:
+ {
+ mapper_name = tok->value;
+ c_parser_consume_token (parser);
+ }
+ break;
+
+ case CPP_KEYWORD:
+ if (tok->keyword == RID_DEFAULT)
+ {
+ c_parser_consume_token (parser);
+ break;
+ }
+ /* Fallthrough. */
+
+ default:
+ error_at (tok->location,
+ "expected identifier or %<default%>");
+ return list;
+ }
+
+ if (!mparens.require_close (parser))
+ {
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+
+ mapper_modifier++;
+ pos += 3;
+ }
}
else
{
@@ -16329,8 +16399,6 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
parens.skip_until_found_close (parser);
return list;
}
-
- c_parser_consume_token (parser);
}
if (c_parser_next_token_is (parser, CPP_NAME)
@@ -16363,8 +16431,30 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
true);
+ tree last_new = NULL_TREE;
+
for (c = nl; 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) = nl;
+ nl = 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;
+ }
parens.skip_until_found_close (parser);
return nl;
@@ -17157,7 +17247,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
c_name = "depend";
break;
case PRAGMA_OMP_CLAUSE_MAP:
- clauses = c_parser_omp_clause_map (parser, clauses);
+ clauses = c_parser_omp_clause_map (parser, clauses, GOMP_MAP_TOFROM);
c_name = "map";
break;
case PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR:
@@ -21157,7 +21247,7 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
{
location_t loc = c_parser_peek_token (parser)->location;
c_parser_consume_pragma (parser);
- tree *pc = NULL, stmt, block;
+ tree *pc = NULL, stmt, block, body, clauses;
if (context != pragma_stmt && context != pragma_compound)
{
@@ -21312,10 +21402,9 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
stmt = make_node (OMP_TARGET);
TREE_TYPE (stmt) = void_type_node;
- OMP_TARGET_CLAUSES (stmt)
- = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
- "#pragma omp target", false);
- for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+ clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+ "#pragma omp target", false);
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION)
{
tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
@@ -21324,14 +21413,19 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = nc;
}
- OMP_TARGET_CLAUSES (stmt)
- = c_finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET);
- c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
+ clauses = c_omp_instantiate_mappers (clauses);
+ clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
+ c_omp_adjust_map_clauses (clauses, true);
- pc = &OMP_TARGET_CLAUSES (stmt);
keep_next_level ();
block = c_begin_compound_stmt (true);
- add_stmt (c_parser_omp_structured_block (parser, if_p));
+ body = c_parser_omp_structured_block (parser, if_p);
+
+ c_omp_scan_mapper_bindings (loc, &clauses, body);
+
+ add_stmt (body);
+ OMP_TARGET_CLAUSES (stmt) = clauses;
+ pc = &OMP_TARGET_CLAUSES (stmt);
OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true);
SET_EXPR_LOCATION (stmt, loc);
@@ -22545,6 +22639,172 @@ c_parser_omp_declare_reduction (c_parser *parser, enum pragma_context context)
}
+/* OpenMP 5.0
+ #pragma omp declare mapper ([mapper-identifier :] type var) \
+ [clause [ [,] clause ] ... ] new-line */
+
+static void
+c_parser_omp_declare_mapper (c_parser *parser, enum pragma_context context)
+{
+ tree type, mapper_name = NULL_TREE, var = NULL_TREE, fndecl, stmt, stmtlist;
+ tree maplist = NULL_TREE, mapper_id, mapper_decl, t;
+ c_token *token;
+ bool nested;
+
+ if (context == pragma_struct || context == pragma_param)
+ {
+ error ("%<#pragma omp declare reduction%> not at file or block scope");
+ goto fail;
+ }
+
+ if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+ goto fail;
+
+ token = c_parser_peek_token (parser);
+
+ if (c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ {
+ switch (token->type)
+ {
+ case CPP_NAME:
+ mapper_name = token->value;
+ c_parser_consume_token (parser);
+ break;
+ case CPP_KEYWORD:
+ if (token->keyword == RID_DEFAULT)
+ {
+ mapper_name = NULL_TREE;
+ c_parser_consume_token (parser);
+ break;
+ }
+ /* Fallthrough. */
+ default:
+ error_at (token->location, "expected identifier or %<default%>");
+ c_parser_skip_to_pragma_eol (parser, false);
+ return;
+ }
+
+ if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+ goto fail;
+ }
+
+ mapper_id = c_omp_mapper_id (mapper_name);
+ mapper_decl = c_omp_mapper_decl (mapper_id);
+
+ {
+ location_t loc = c_parser_peek_token (parser)->location;
+ struct c_type_name *ctype = c_parser_type_name (parser);
+ type = groktypename (ctype, NULL, NULL);
+ if (type == error_mark_node)
+ goto fail;
+ if (TREE_CODE (type) != RECORD_TYPE
+ && TREE_CODE (type) != UNION_TYPE)
+ {
+ error_at (loc, "%qT is not a struct or union type in "
+ "%<#pragma omp declare mapper%>", type);
+ c_parser_skip_to_pragma_eol (parser, false);
+ return;
+ }
+ for (tree t = DECL_INITIAL (mapper_decl); t; t = TREE_CHAIN (t))
+ if (comptypes (TREE_PURPOSE (t), type))
+ {
+ error_at (loc, "redeclaration of %qs %<#pragma omp declare "
+ "mapper%> for type %qT", IDENTIFIER_POINTER (mapper_id)
+ + sizeof ("omp declare mapper ") - 1,
+ type);
+ tree prevmapper = TREE_VALUE (t);
+ /* Hmm, this location might not be very accurate. */
+ location_t ploc
+ = DECL_SOURCE_LOCATION (OMP_DECLARE_MAPPER_DECL (prevmapper));
+ error_at (ploc, "previous %<#pragma omp declare mapper%>");
+ c_parser_skip_to_pragma_eol (parser, false);
+ return;
+ }
+ }
+
+ token = c_parser_peek_token (parser);
+ if (token->type == CPP_NAME)
+ {
+ var = build_decl (token->location, VAR_DECL, token->value, type);
+ c_parser_consume_token (parser);
+ DECL_ARTIFICIAL (var) = 1;
+ }
+ else
+ {
+ error_at (token->location, "expected identifier");
+ goto fail;
+ }
+
+ if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
+ goto fail;
+
+ nested = current_function_decl != NULL_TREE;
+ if (nested)
+ c_push_function_context ();
+
+ fndecl = build_decl (BUILTINS_LOCATION, FUNCTION_DECL, mapper_id,
+ default_function_type);
+ current_function_decl = fndecl;
+ allocate_struct_function (fndecl, true);
+ push_scope ();
+ stmtlist = push_stmt_list ();
+ pushdecl (var);
+ DECL_CONTEXT (var) = fndecl;
+
+ while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL))
+ {
+ location_t here;
+ pragma_omp_clause c_kind;
+ here = c_parser_peek_token (parser)->location;
+ c_kind = c_parser_omp_clause_name (parser);
+ if (c_kind != PRAGMA_OMP_CLAUSE_MAP)
+ {
+ error_at (here, "unexpected clause");
+ goto fail;
+ }
+ maplist = c_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET);
+ }
+
+ if (maplist == NULL_TREE)
+ {
+ error_at (input_location, "missing %<map%> clause");
+ goto fail;
+ }
+
+ stmt = make_node (OMP_DECLARE_MAPPER);
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_DECLARE_MAPPER_ID (stmt) = mapper_name;
+ OMP_DECLARE_MAPPER_TYPE (stmt) = type;
+ OMP_DECLARE_MAPPER_DECL (stmt) = var;
+ OMP_DECLARE_MAPPER_CLAUSES (stmt) = maplist;
+
+ add_stmt (stmt);
+
+ pop_stmt_list (stmtlist);
+ pop_scope ();
+
+ if (cfun->language != NULL)
+ {
+ ggc_free (cfun->language);
+ cfun->language = NULL;
+ }
+ set_cfun (NULL);
+ current_function_decl = NULL_TREE;
+
+ if (nested)
+ c_pop_function_context ();
+
+ c_parser_skip_to_pragma_eol (parser);
+
+ t = tree_cons (type, stmt, DECL_INITIAL (mapper_decl));
+ DECL_INITIAL (mapper_decl) = t;
+
+ return;
+
+ fail:
+ c_parser_skip_to_pragma_eol (parser);
+}
+
/* OpenMP 4.0
#pragma omp declare simd declare-simd-clauses[optseq] new-line
#pragma omp declare reduction (reduction-id : typename-list : expression) \
@@ -22574,6 +22834,12 @@ c_parser_omp_declare (c_parser *parser, enum pragma_context context)
c_parser_omp_declare_reduction (parser, context);
return false;
}
+ if (strcmp (p, "mapper") == 0)
+ {
+ c_parser_consume_token (parser);
+ c_parser_omp_declare_mapper (parser, context);
+ return false;
+ }
if (!flag_openmp) /* flag_openmp_simd */
{
c_parser_skip_to_pragma_eol (parser, false);
@@ -761,6 +761,10 @@ extern tree c_finish_omp_task (location_t, tree, tree);
extern void c_finish_omp_cancel (location_t, tree);
extern void c_finish_omp_cancellation_point (location_t, tree);
extern tree c_finish_omp_clauses (tree, enum c_omp_region_type);
+extern tree c_omp_finish_mapper_clauses (tree);
+extern tree c_omp_mapper_lookup (tree, tree);
+extern tree c_omp_extract_mapper_directive (tree);
+extern tree c_omp_map_array_section (location_t, tree);
extern tree c_build_va_arg (location_t, tree, location_t, tree);
extern tree c_finish_transaction (location_t, tree, int);
extern bool c_tree_equal (tree, tree);
@@ -812,6 +816,10 @@ extern tree c_omp_reduction_id (enum tree_code, tree);
extern tree c_omp_reduction_decl (tree);
extern tree c_omp_reduction_lookup (tree, tree);
extern tree c_check_omp_declare_reduction_r (tree *, int *, void *);
+extern tree c_omp_mapper_id (tree);
+extern tree c_omp_mapper_decl (tree);
+extern void c_omp_scan_mapper_bindings (location_t, tree *, tree);
+extern tree c_omp_instantiate_mappers (tree);
extern bool c_check_in_current_scope (tree);
extern void c_pushtag (location_t, tree, tree);
extern void c_bind (location_t, tree, bool);
@@ -14877,6 +14877,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
t = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME))
+ {
+ remove = true;
+ break;
+ }
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
if (handle_omp_array_sections (c, ort))
@@ -15642,6 +15649,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
return clauses;
}
+/* Do processing necessary to make CLAUSES well-formed, where CLAUSES result
+ from implicit instantiation of user-defined mappers (in gimplify.cc). */
+
+tree
+c_omp_finish_mapper_clauses (tree clauses)
+{
+ return c_finish_omp_clauses (clauses, C_ORT_OMP);
+}
+
/* Return code to initialize DST with a copy constructor from SRC.
C doesn't have copy constructors nor assignment operators, only for
_Atomic vars we need to perform __atomic_load from src into a temporary
@@ -184,6 +184,13 @@ extern tree cxx_simulate_record_decl (location_t, const char *,
#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_MAPPABLE_TYPE
@@ -8218,6 +8218,9 @@ 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);
@@ -44665,7 +44665,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
OMP_CLAUSE_CHAIN (c) = nc;
}
if (!processing_template_decl)
- clauses = omp_instantiate_mappers (clauses);
+ clauses = c_omp_instantiate_mappers (clauses);
clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
c_omp_adjust_map_clauses (clauses, true);
@@ -17865,7 +17865,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
if (ort != C_ORT_OMP_DECLARE_SIMD)
{
if (ort == C_ORT_OMP_TARGET)
- new_clauses = omp_instantiate_mappers (new_clauses);
+ new_clauses = c_omp_instantiate_mappers (new_clauses);
new_clauses = finish_omp_clauses (new_clauses, ort);
if (linear_no_step)
for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
@@ -5992,8 +5992,8 @@ omp_mapper_id (tree mapper_id, tree type)
return get_identifier (name);
}
-static tree
-omp_mapper_lookup (tree id, tree type)
+tree
+cxx_omp_mapper_lookup (tree id, tree type)
{
if (TREE_CODE (type) != RECORD_TYPE
&& TREE_CODE (type) != UNION_TYPE)
@@ -6002,8 +6002,8 @@ omp_mapper_lookup (tree id, tree type)
return lookup_name (id);
}
-static tree
-omp_extract_mapper_directive (tree fndecl)
+tree
+cxx_omp_extract_mapper_directive (tree fndecl)
{
if (BASELINK_P (fndecl))
/* See through BASELINK nodes to the underlying function. */
@@ -6027,6 +6027,31 @@ omp_extract_mapper_directive (tree fndecl)
return body;
}
+/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but
+ nothing more complicated. */
+
+tree
+cxx_omp_map_array_section (location_t loc, tree t)
+{
+ tree low = TREE_OPERAND (t, 1);
+ tree len = TREE_OPERAND (t, 2);
+
+ if (len && integer_onep (len))
+ {
+ t = TREE_OPERAND (t, 0);
+
+ if (!low)
+ low = integer_zero_node;
+
+ if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
+ t = convert_from_reference (t);
+
+ t = build_array_ref (loc, t, low);
+ }
+
+ return t;
+}
+
/* Helper function for cp_parser_omp_declare_reduction_exprs
and tsubst_omp_udr.
Remove CLEANUP_STMT for data (omp_priv variable).
@@ -6793,242 +6818,6 @@ cp_oacc_check_attachments (tree c)
return false;
}
-struct remap_mapper_decl_info
-{
- tree dummy_var;
- tree expr;
-};
-
-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 tree *
-omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
- enum gomp_map_kind outer_kind)
-{
- tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
- tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
- tree mapper_name = NULL_TREE;
-
- remap_mapper_decl_info map_info;
- map_info.dummy_var = dummy_var;
- map_info.expr = expr;
-
- for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
- {
- tree unshared = unshare_expr (c);
- enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c);
- tree t = OMP_CLAUSE_DECL (unshared);
- tree type = NULL_TREE;
- bool nonunit_array_with_mapper = false;
-
- if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
- {
- mapper_name = t;
- continue;
- }
- else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
- {
- mapper_name = NULL_TREE;
- continue;
- }
-
- if (TREE_CODE (t) == OMP_ARRAY_SECTION)
- {
- tree low = TREE_OPERAND (t, 1);
- tree len = TREE_OPERAND (t, 2);
-
- if (len && integer_onep (len))
- {
- t = TREE_OPERAND (t, 0);
-
- if (POINTER_TYPE_P (TREE_TYPE (t))
- || TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
- type = TREE_TYPE (TREE_TYPE (t));
-
- if (!low)
- low = integer_zero_node;
-
- if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
- t = convert_from_reference (t);
-
- t = build_array_ref (OMP_CLAUSE_LOCATION (c), t, low);
- }
- else
- {
- type = TREE_TYPE (t);
- nonunit_array_with_mapper = true;
- }
- }
- else
- type = TREE_TYPE (t);
-
- gcc_assert (type);
-
- if (type == error_mark_node)
- continue;
-
- walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);
-
- if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)
- OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
-
- type = TYPE_MAIN_VARIANT (type);
-
- tree mapper_fn = 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 = omp_extract_mapper_directive (mapper_fn);
- if (nested_mapper != mapper)
- {
- if (clause_kind == GOMP_MAP_UNSET)
- clause_kind = outer_kind;
-
- outlist = omp_instantiate_mapper (outlist, nested_mapper,
- t, clause_kind);
- continue;
- }
- }
- else if (mapper_name)
- {
- error ("mapper %qE not found for type %qT", mapper_name, type);
- continue;
- }
-
- *outlist = unshared;
- outlist = &OMP_CLAUSE_CHAIN (unshared);
- }
-
- return outlist;
-}
-
-tree
-omp_instantiate_mappers (tree clauses)
-{
- tree c, *pc, mapper_name = NULL_TREE;
-
- for (pc = &clauses, c = clauses; c; c = *pc)
- {
- bool using_mapper = false;
-
- switch (OMP_CLAUSE_CODE (c))
- {
- case OMP_CLAUSE_MAP:
- {
- tree t = OMP_CLAUSE_DECL (c);
- tree type = NULL_TREE;
- bool nonunit_array_with_mapper = false;
-
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
- {
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME)
- mapper_name = OMP_CLAUSE_DECL (c);
- else
- mapper_name = NULL_TREE;
- pc = &OMP_CLAUSE_CHAIN (c);
- continue;
- }
-
- gcc_assert (TREE_CODE (t) != TREE_LIST);
-
- if (TREE_CODE (t) == OMP_ARRAY_SECTION)
- {
- tree low = TREE_OPERAND (t, 1);
- tree len = TREE_OPERAND (t, 2);
-
- if (len && integer_onep (len))
- {
- t = TREE_OPERAND (t, 0);
-
- if (!TREE_TYPE (t))
- {
- pc = &OMP_CLAUSE_CHAIN (c);
- continue;
- }
-
- if (POINTER_TYPE_P (TREE_TYPE (t))
- || TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
- type = TREE_TYPE (TREE_TYPE (t));
-
- if (!low)
- low = integer_zero_node;
- }
- else
- {
- /* !!! Array sections of size >1 with mappers for elements
- are hard to support. Do something here. */
- nonunit_array_with_mapper = true;
- 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 = 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 = omp_extract_mapper_directive (mapper_fn);
- pc = omp_instantiate_mapper (pc, mapper, t, kind);
- using_mapper = true;
- }
- else if (mapper_name)
- {
- error ("mapper %qE not found for type %qT", mapper_name, type);
- using_mapper = true;
- }
- }
- break;
-
- default:
- ;
- }
-
- if (using_mapper)
- *pc = OMP_CLAUSE_CHAIN (c);
- else
- pc = &OMP_CLAUSE_CHAIN (c);
- }
-
- return clauses;
-}
-
/* For all elements of CLAUSES, validate them vs OpenMP constraints.
Remove any elements from the list that are invalid. */
@@ -9640,108 +9429,6 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses)
return add_stmt (stmt);
}
-struct mapper_list
-{
- hash_set<omp_name_type> *seen_types;
- vec<tree> *mappers;
-
- mapper_list (hash_set<omp_name_type> *s, vec<tree> *m)
- : seen_types (s), mappers (m) { }
-
- void add_mapper (tree name, tree type, tree mapperfn)
- {
- /* We can't hash a NULL_TREE... */
- if (!name)
- name = void_node;
-
- omp_name_type n_t = { name, type };
-
- if (seen_types->contains (n_t))
- return;
-
- seen_types->add (n_t);
- mappers->safe_push (mapperfn);
- }
-
- bool contains (tree name, tree type)
- {
- if (!name)
- name = void_node;
-
- return seen_types->contains ({ name, type });
- }
-};
-
-static void
-find_nested_mappers (mapper_list *mlist, tree mapper_fn)
-{
- tree mapper = 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); //TREE_CHAIN (expr);
-
- 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
- = omp_mapper_lookup (mapper_name, elem_type);
-
- if (nested_mapper_fn)
- {
- mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn);
- find_nested_mappers (mlist, nested_mapper_fn);
- }
- else if (mapper_name)
- {
- error ("mapper %qE not found for type %qT", mapper_name,
- elem_type);
- continue;
- }
- }
- }
-}
-
/* Used to walk OpenMP target directive body. */
struct omp_target_walk_data
@@ -9768,7 +9455,7 @@ struct omp_target_walk_data
variables when recording lambda_objects_accessed. */
hash_set<tree> local_decls;
- mapper_list *mappers;
+ omp_mapper_list *mappers;
};
/* Helper function of finish_omp_target_clauses, called via
@@ -9782,7 +9469,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;
- mapper_list *mlist = data->mappers;
+ omp_mapper_list *mlist = data->mappers;
tree aggr_type = NULL_TREE;
/* References inside of these expression codes shouldn't incur any
@@ -9808,7 +9495,7 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
if (aggr_type)
{
- tree mapper_fn = omp_mapper_lookup (NULL_TREE, 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);
}
@@ -9918,7 +9605,7 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
hash_set<omp_name_type> seen_types;
auto_vec<tree> mapper_fns;
- mapper_list mlist (&seen_types, &mapper_fns);
+ omp_mapper_list mlist (&seen_types, &mapper_fns);
data.mappers = &mlist;
cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data);
@@ -9926,13 +9613,13 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
unsigned int i;
tree mapper_fn;
FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
- find_nested_mappers (&mlist, mapper_fn);
+ c_omp_find_nested_mappers (&mlist, mapper_fn);
auto_vec<tree, 16> new_clauses;
FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
{
- tree mapper = omp_extract_mapper_directive (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);
@@ -10464,24 +10464,35 @@ omp_instantiate_mapper (hash_map<omp_name_type, tree> *implicit_mappers,
continue;
}
- tree decl = OMP_CLAUSE_DECL (clause), unshared;
+ tree decl = OMP_CLAUSE_DECL (clause), unshared, type;
+ bool nonunit_array_with_mapper = false;
- if (TREE_CODE (decl) == OMP_ARRAY_SECTION
- && TREE_OPERAND (decl, 2)
- && integer_onep (TREE_OPERAND (decl, 2)))
+ if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
{
- unshared = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
- OMP_CLAUSE_CODE (clause));
- tree low = TREE_OPERAND (decl, 1);
- if (!low || integer_zerop (low))
- OMP_CLAUSE_DECL (unshared)
- = build_fold_indirect_ref (TREE_OPERAND (decl, 0));
+ 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
- OMP_CLAUSE_DECL (unshared) = decl;
- OMP_CLAUSE_SIZE (unshared) = OMP_CLAUSE_SIZE (clause);
+ {
+ 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);
+ {
+ unshared = unshare_expr (clause);
+ type = TREE_TYPE (decl);
+ }
walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);
@@ -10489,12 +10500,18 @@ omp_instantiate_mapper (hash_map<omp_name_type, tree> *implicit_mappers,
OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
decl = OMP_CLAUSE_DECL (unshared);
- tree type = TYPE_MAIN_VARIANT (TREE_TYPE (decl));
+ type = TYPE_MAIN_VARIANT (type);
tree *nested_mapper_p = implicit_mappers->get ({ mapper_name, type });
if (nested_mapper_p && *nested_mapper_p != mapper)
{
+ if (nonunit_array_with_mapper)
+ {
+ sorry ("user-defined mapper with non-unit length array section");
+ continue;
+ }
+
if (clause_kind == GOMP_MAP_UNSET)
clause_kind = outer_kind;
@@ -11505,16 +11522,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree var = OMP_CLAUSE__MAPPER_BINDING__DECL (c);
tree type = TYPE_MAIN_VARIANT (TREE_TYPE (var));
tree fndecl = OMP_CLAUSE__MAPPER_BINDING__MAPPER (c);
- tree mapper = DECL_SAVED_TREE (fndecl);
- if (TREE_CODE (mapper) == BIND_EXPR)
- mapper = BIND_EXPR_BODY (mapper);
- if (TREE_CODE (mapper) == STATEMENT_LIST)
- {
- tree_stmt_iterator tsi = tsi_start (mapper);
- gcc_assert (TREE_CODE (tsi_stmt (tsi)) == DECL_EXPR);
- tsi_next (&tsi);
- mapper = tsi_stmt (tsi);
- }
+ tree mapper
+ = lang_hooks.decls.omp_extract_mapper_directive (fndecl);
gcc_assert (mapper != NULL_TREE
&& TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
ctx->implicit_mappers->put ({ name, type }, mapper);
@@ -85,6 +85,9 @@ 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_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);
@@ -272,6 +275,10 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
#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
@@ -306,6 +313,9 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
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, \
@@ -643,6 +643,32 @@ 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. */
@@ -310,6 +310,18 @@ struct lang_hooks_for_decls
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);
@@ -201,4 +201,36 @@ struct default_hash_traits <omp_name_type>
}
};
+struct omp_mapper_list
+{
+ hash_set<omp_name_type> *seen_types;
+ vec<tree> *mappers;
+
+ omp_mapper_list (hash_set<omp_name_type> *s, vec<tree> *m)
+ : seen_types (s), mappers (m) { }
+
+ void add_mapper (tree name, tree type, tree mapperfn)
+ {
+ /* We can't hash a NULL_TREE... */
+ if (!name)
+ name = void_node;
+
+ omp_name_type n_t = { name, type };
+
+ if (seen_types->contains (n_t))
+ return;
+
+ seen_types->add (n_t);
+ mappers->safe_push (mapperfn);
+ }
+
+ bool contains (tree name, tree type)
+ {
+ if (!name)
+ name = void_node;
+
+ return seen_types->contains ({ name, type });
+ }
+};
+
#endif /* GCC_OMP_GENERAL_H */
new file mode 100644
@@ -0,0 +1,22 @@
+/* { dg-do compile } */
+
+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 } */
similarity index 75%
rename from gcc/testsuite/g++.dg/gomp/declare-mapper-3.C
rename to gcc/testsuite/c-c++-common/gomp/declare-mapper-3.c
@@ -1,6 +1,8 @@
// { dg-do compile }
// { dg-additional-options "-fdump-tree-gimple" }
+#include <stdlib.h>
+
// Test named mapper invocation.
struct S {
@@ -11,10 +13,11 @@ struct S {
int main (int argc, char *argv[])
{
int N = 1024;
-#pragma omp declare mapper (mapN:S s) map(to:s.ptr, s.size) map(s.ptr[:N])
+#pragma omp declare mapper (mapN:struct S s) map(to:s.ptr, s.size) \
+ map(s.ptr[:N])
- S s;
- s.ptr = new int[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\(to: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" } }
new file mode 100644
@@ -0,0 +1,78 @@
+/* { dg-do compile } */
+/* { 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 } */
+
+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 "previous 'pragma omp declare mapper' declaration" "" { 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 "redeclaration of 'pragma omp declare mapper'" "" { 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 "previous 'pragma omp declare mapper' declaration" "" { 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 "redeclaration of 'pragma omp declare mapper'" "" { target c++ } .-3 } */
new file mode 100644
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+
+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 } */
+/* { dg-error "expected '\\)' before '\\\]' token" "" { target c++ } .-3 } */
+
+int y = 7;
new file mode 100644
@@ -0,0 +1,30 @@
+/* { dg-do compile } */
+
+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 } */
+ /* { dg-error "expected '\\)' before '\\\]' token" "" { target c++ } .-3 } */
+ int y = 7;
+ return y;
+}
new file mode 100644
@@ -0,0 +1,43 @@
+/* { dg-do compile } */
+
+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 } */
+
+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 "previous 'pragma omp declare mapper' declaration" "" { 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 "redeclaration of 'pragma omp declare mapper'" "" { 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 "previous 'pragma omp declare mapper' declaration" "" { 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'" "" { target c++ } .-2 } */
+}
deleted file mode 100644
@@ -1,74 +0,0 @@
-// { dg-do compile }
-// { dg-additional-options "-fdump-tree-original" }
-
-// Check mapper binding clauses.
-
-struct Y {
- int z;
-};
-
-struct Z {
- int z;
-};
-
-#pragma omp declare mapper (Y y) map(tofrom: y)
-#pragma omp declare mapper (Z z) map(tofrom: z)
-
-int foo (void)
-{
- Y yy;
- Z zz;
- int dummy;
-
-#pragma omp target data map(dummy)
- {
- #pragma omp target
- {
- yy.z++;
- zz.z++;
- }
- yy.z++;
- }
- return yy.z;
-}
-
-struct P
-{
- Z *zp;
-};
-
-int bar (void)
-{
- Y yy;
- Z zz;
- P pp;
- Z t;
- int dummy;
-
- pp.zp = &t;
-
-#pragma omp declare mapper (Y y) map(tofrom: y.z)
-#pragma omp declare mapper (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(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" } }
-// { dg-final { scan-tree-dump {mapper_binding\(struct Z,omp declare mapper ~1Z\) mapper_binding\(struct P,omp declare mapper ~1P\)} "original" } }
new file mode 100644
@@ -0,0 +1,61 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+// "omp declare mapper" support -- check expansion in gimple.
+
+#include <stdlib.h>
+
+struct S {
+ int *ptr;
+ int size;
+};
+
+#define N 64
+
+#pragma omp declare mapper (struct S w) map(w.size, w.ptr, w.ptr[:w.size])
+#pragma omp declare mapper (foo:struct S w) map(to:w.size, w.ptr) \
+ map(w.ptr[:w.size])
+
+int main (int argc, char *argv[])
+{
+ struct S s;
+ s.ptr = (int *) malloc (sizeof (int) * N);
+ s.size = N;
+
+#pragma omp declare mapper (bar:struct 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\(tofrom: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" { target c++ } } } */
+/* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(to:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" { target c++ } } } */
new file mode 100644
@@ -0,0 +1,33 @@
+// { dg-do compile }
+
+// Error-checking tests for "omp declare mapper".
+
+typedef struct {
+ int *ptr;
+ int size;
+} S;
+
+typedef struct {
+ int z;
+} Z;
+
+int main (int argc, char *argv[])
+{
+#pragma omp declare mapper (S v) map(v.size, v.ptr[:v.size])
+/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */
+
+ /* This one's a duplicate. */
+#pragma omp declare mapper (default: S v) map (to: v.size) map (v)
+/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'S'" "" { target c } .-1 } */
+
+ /* ...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'" "" { target c } .-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 or union type in '#pragma omp declare mapper'" "" { target c } .-1 } */
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,58 @@
+#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,57 @@
+#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,85 @@
+#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 } */
+
+#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 } */
+
+#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,60 @@
+#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 "declare mapper" directives (and the "mapper" modifier on "map" clauses) for C. As for C++, arrays of custom-mapped objects are not supported yet. I've taken hints from the existing C support for "declare reduction" directives: this works a little differently from C++ for things such as looking up user-defined reductions (or user-defined mappers, in our case). Some support functions have been pulled out of the C++ FE and shared with the C implementation: several language hooks have been added to facilitate that, given the above differences. (Fortran FE support is TBD.) 2022-03-17 Julian Brown <julian@codesourcery.com> gcc/c-family/ * c-common.h (omp_mapper_list, c_omp_find_nested_mappers, c_omp_instantiate_mappers): Add forward declarations/prototypes. * c-omp.cc (c_omp_find_nested_mappers): New function. (remap_mapper_decl_info): New struct. (remap_mapper_decl_1, omp_instantiate_mapper, c_omp_instantiate_mappers): Add functions. * c-decl.cc (c_omp_mapper_id, c_omp_mapper_decl, c_omp_mapper_lookup, c_omp_extract_mapper_directive, c_omp_map_array_section, c_omp_scan_mapper_bindings_r, c_omp_scan_mapper_bindings): New functions. gcc/c/ * c-objc-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 for C. * c-parser.cc (c_parser_omp_clause_map): Add KIND parameter. Handle mapper modifier. (c_parser_omp_all_clauses): Update call to c_parser_omp_clause_map with new kind argument. (c_parser_omp_target): Instantiate explicit mappers and record bindings for implicit mappers. (c_parser_omp_declare_mapper): Parse "declare mapper" directives. (c_parser_omp_declare): Support "declare mapper". * c-tree.h (c_omp_finish_mapper_clauses, c_omp_mapper_lookup, c_omp_extract_mapper_directive, c_omp_map_array_section, c_omp_mapper_id, c_omp_mapper_decl, c_omp_scan_mapper_bindings, c_omp_instantiate_mappers): Add prototypes. * c-typeck.cc (c_finish_omp_clauses): Handle GOMP_MAP_PUSH_MAPPER_NAME and GOMP_MAP_POP_MAPPER_NAME. (c_omp_finish_mapper_clauses): New function (langhook). gcc/cp/ * cp-objcp-common.h (LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks for C++. * cp-tree.h (cxx_omp_mapper_lookup, cxx_omp_extract_mapper_directive, cxx_omp_map_array_section): Add prototypes. * parser.cc (cp_parser_omp_target): Use new name for c_omp_instantiate_mappers. * pt.cc (tsubst_omp_clauses): Use new name for c_omp_instantiate_mappers. (omp_mapper_lookup): Rename to... (cxx_omp_mapper_lookup): This. (omp_extract_mapper_directive): Rename to... (cxx_omp_extract_mapper_directive): This. (cxx_omp_map_array_section): New function. (remap_mapper_decl_info, remap_mapper_decl_1, omp_instantiate_mapper, omp_instantiate_mappers, mapper_list, find_nested_mappers): Remove. (omp_target_walk_data): Rename mapper_list to omp_mapper_list. (finish_omp_target_clauses_r): Likewise. Use renamed cxx_omp_mapper_lookup, cxx_omp_extract_mapper_directive and c_omp_find_nested_mappers. (finish_omp_target_clauses): Likewise. gcc/ * gimplify.cc (omp_instantiate_mapper): Use omp_map_array_section langhook to handle (singleton only, for now) array sections. Diagnose attempts to use length >1 array sections with custom mappers. (gimplify_scan_omp_clauses): Use omp_extract_mapper_directive langhook. * langhooks-def.h (lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): Add prototypes. (LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define lang hooks. (LANG_HOOKS_DECLS): Add LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION. * langhooks.cc (lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): New default definitions of langhooks. * langhooks.h (lang_hooks_for_decls): Add omp_mapper_lookup, omp_extract_mapper_directive, omp_map_array_section. * omp-general.h (omp_mapper_list): New. gcc/testsuite/ * g++.dg/gomp/declare-mapper-3.C: Remove from here. * c-c++-common/gomp/declare-mapper-3.c: Move test here, make C-compatible. * g++.dg/gomp/declare-mapper-4.C: Remove from here. * c-c++-common/gomp/declare-mapper-4.c: Move test here, make C-compatible. * c-c++-common/gomp/declare-mapper-5.c: New test. * c-c++-common/gomp/declare-mapper-6.c: New test. * c-c++-common/gomp/declare-mapper-7.c: New test. * c-c++-common/gomp/declare-mapper-8.c: New test. * c-c++-common/gomp/declare-mapper-9.c: New test. * gcc.dg/gomp/declare-mapper-10.c: New test. * gcc.dg/gomp/declare-mapper-11.c: New test. * c-c++-common/gomp/declare-mapper-12.c: New test. libgomp/ * testsuite/libgomp.c-c++-common/declare-mapper-9.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-10.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-11.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-12.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-13.c: New test. * testsuite/libgomp.c-c++-common/declare-mapper-14.c: New test. --- gcc/c-family/c-common.h | 3 + gcc/c-family/c-omp.cc | 300 ++++++++++++++ gcc/c/c-decl.cc | 169 ++++++++ gcc/c/c-objc-common.h | 12 + gcc/c/c-parser.cc | 298 +++++++++++++- gcc/c/c-tree.h | 8 + gcc/c/c-typeck.cc | 16 + gcc/cp/cp-objcp-common.h | 7 + gcc/cp/cp-tree.h | 3 + gcc/cp/parser.cc | 2 +- gcc/cp/pt.cc | 2 +- gcc/cp/semantics.cc | 383 ++---------------- gcc/gimplify.cc | 57 +-- gcc/langhooks-def.h | 10 + gcc/langhooks.cc | 26 ++ gcc/langhooks.h | 12 + gcc/omp-general.h | 32 ++ .../c-c++-common/gomp/declare-mapper-12.c | 22 + .../gomp/declare-mapper-3.c} | 9 +- .../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 | 24 ++ .../c-c++-common/gomp/declare-mapper-7.c | 30 ++ .../c-c++-common/gomp/declare-mapper-8.c | 43 ++ .../c-c++-common/gomp/declare-mapper-9.c | 34 ++ gcc/testsuite/g++.dg/gomp/declare-mapper-4.C | 74 ---- gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c | 61 +++ gcc/testsuite/gcc.dg/gomp/declare-mapper-11.c | 33 ++ .../libgomp.c-c++-common/declare-mapper-10.c | 58 +++ .../libgomp.c-c++-common/declare-mapper-11.c | 57 +++ .../libgomp.c-c++-common/declare-mapper-12.c | 85 ++++ .../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 | 60 +++ 34 files changed, 1679 insertions(+), 467 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-12.c rename gcc/testsuite/{g++.dg/gomp/declare-mapper-3.C => c-c++-common/gomp/declare-mapper-3.c} (75%) 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 delete mode 100644 gcc/testsuite/g++.dg/gomp/declare-mapper-4.C create mode 100644 gcc/testsuite/gcc.dg/gomp/declare-mapper-10.c create mode 100644 gcc/testsuite/gcc.dg/gomp/declare-mapper-11.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