@@ -2284,6 +2284,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
@@ -182,6 +182,8 @@ 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_PRIVATIZE_BY_REFERENCE
#define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference
#undef LANG_HOOKS_OMP_MAPPABLE_TYPE
@@ -2891,6 +2891,7 @@ struct GTY(()) lang_decl_fn {
unsigned this_thunk_p : 1;
unsigned omp_declare_reduction_p : 1;
+ unsigned omp_declare_mapper_p : 1;
unsigned has_dependent_explicit_spec_p : 1;
unsigned immediate_fn_p : 1;
unsigned maybe_deleted : 1;
@@ -4295,6 +4296,11 @@ get_vec_init_expr (tree t)
#define DECL_OMP_DECLARE_REDUCTION_P(NODE) \
(LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_reduction_p)
+/* Nonzero if NODE is an artificial FUNCTION_DECL for
+ #pragma omp declare mapper. */
+#define DECL_OMP_DECLARE_MAPPER_P(NODE) \
+ (LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_mapper_p)
+
/* Nonzero if DECL has been declared threadprivate by
#pragma omp threadprivate. */
#define CP_DECL_THREADPRIVATE_P(DECL) \
@@ -7653,10 +7659,13 @@ extern tree finish_qualified_id_expr (tree, tree, bool, bool,
extern void simplify_aggr_init_expr (tree *);
extern void finalize_nrv (tree *, tree, tree);
extern tree omp_reduction_id (enum tree_code, tree, tree);
+extern tree omp_mapper_id (tree, tree);
extern tree cp_remove_omp_priv_cleanup_stmt (tree *, int *, void *);
extern bool cp_check_omp_declare_reduction (tree);
+extern bool cp_check_omp_declare_mapper (tree);
extern void finish_omp_declare_simd_methods (tree);
extern tree finish_omp_clauses (tree, enum c_omp_region_type);
+extern tree omp_instantiate_mappers (tree);
extern tree push_omp_privatization_clauses (bool);
extern void pop_omp_privatization_clauses (tree);
extern void save_omp_privatization_clauses (vec<tree> &);
@@ -8208,6 +8217,7 @@ 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 bool cxx_omp_privatize_by_reference (const_tree);
extern bool cxx_omp_disregard_value_expr (tree, bool);
extern void cp_fold_function (tree);
@@ -1859,6 +1859,18 @@ duplicate_decls (tree newdecl, tree olddecl, bool hiding, bool was_hidden)
"previous %<pragma omp declare reduction%> declaration");
return error_mark_node;
}
+ else if (TREE_CODE (newdecl) == FUNCTION_DECL
+ && DECL_OMP_DECLARE_MAPPER_P (newdecl))
+ {
+ /* OMP UDMs are never duplicates either. */
+ gcc_assert (DECL_OMP_DECLARE_MAPPER_P (olddecl));
+ error_at (newdecl_loc,
+ "redeclaration of %<pragma omp declare mapper%>");
+ inform (olddecl_loc,
+ "previous %<pragma omp declare mapper%> declaration");
+ return error_mark_node;
+
+ }
else if (TREE_CODE (newdecl) == FUNCTION_DECL
&& ((DECL_TEMPLATE_SPECIALIZATION (olddecl)
&& (!DECL_TEMPLATE_INFO (newdecl)
@@ -17821,7 +17833,8 @@ finish_function (bool inline_p)
/* Perform delayed folding before NRV transformation. */
if (!processing_template_decl
&& !DECL_IMMEDIATE_FUNCTION_P (fndecl)
- && !DECL_OMP_DECLARE_REDUCTION_P (fndecl))
+ && !DECL_OMP_DECLARE_REDUCTION_P (fndecl)
+ && !DECL_OMP_DECLARE_MAPPER_P (fndecl))
cp_fold_function (fndecl);
/* Set up the named return value optimization, if we can. Candidate
@@ -17898,7 +17911,8 @@ finish_function (bool inline_p)
/* Genericize before inlining. */
if (!processing_template_decl
&& !DECL_IMMEDIATE_FUNCTION_P (fndecl)
- && !DECL_OMP_DECLARE_REDUCTION_P (fndecl))
+ && !DECL_OMP_DECLARE_REDUCTION_P (fndecl)
+ && !DECL_OMP_DECLARE_MAPPER_P (fndecl))
cp_genericize (fndecl);
/* Emit the resumer and destroyer functions now, providing that we have
@@ -931,10 +931,11 @@ decl_mangling_context (tree decl)
tcontext = CP_DECL_CONTEXT (decl);
- /* Ignore the artificial declare reduction functions. */
+ /* Ignore the artificial declare reduction and declare mapper functions. */
if (tcontext
&& TREE_CODE (tcontext) == FUNCTION_DECL
- && DECL_OMP_DECLARE_REDUCTION_P (tcontext))
+ && (DECL_OMP_DECLARE_REDUCTION_P (tcontext)
+ || DECL_OMP_DECLARE_MAPPER_P (tcontext)))
return decl_mangling_context (tcontext);
return tcontext;
@@ -3354,7 +3354,8 @@ set_decl_context_in_fn (tree ctx, tree decl)
gcc_checking_assert (DECL_LOCAL_DECL_P (decl)
&& (DECL_NAMESPACE_SCOPE_P (decl)
|| (TREE_CODE (decl) == FUNCTION_DECL
- && DECL_OMP_DECLARE_REDUCTION_P (decl))));
+ && (DECL_OMP_DECLARE_REDUCTION_P (decl)
+ || DECL_OMP_DECLARE_MAPPER_P (decl)))));
if (!DECL_CONTEXT (decl)
/* When parsing the parameter list of a function declarator,
@@ -26203,10 +26203,12 @@ cp_parser_class_specifier_1 (cp_parser* parser)
{
/* OpenMP UDRs need to be parsed before all other functions. */
FOR_EACH_VEC_SAFE_ELT (unparsed_funs_with_definitions, ix, decl)
- if (DECL_OMP_DECLARE_REDUCTION_P (decl))
+ if (DECL_OMP_DECLARE_REDUCTION_P (decl)
+ || DECL_OMP_DECLARE_MAPPER_P (decl))
cp_parser_late_parsing_for_member (parser, decl);
FOR_EACH_VEC_SAFE_ELT (unparsed_funs_with_definitions, ix, decl)
- if (!DECL_OMP_DECLARE_REDUCTION_P (decl))
+ if (!DECL_OMP_DECLARE_REDUCTION_P (decl)
+ && !DECL_OMP_DECLARE_MAPPER_P (decl))
cp_parser_late_parsing_for_member (parser, decl);
}
else
@@ -32133,6 +32135,8 @@ cp_parser_enclosed_template_argument_list (cp_parser* parser)
return arguments;
}
+static bool cp_parser_omp_declare_mapper_maplist (tree, cp_parser *);
+
/* MEMBER_FUNCTION is a member function, or a friend. If default
arguments, or the body of the function have not yet been parsed,
parse them now. */
@@ -32193,6 +32197,16 @@ cp_parser_late_parsing_for_member (cp_parser* parser, tree member_function)
finish_function (/*inline_p=*/true);
cp_check_omp_declare_reduction (member_function);
}
+ else if (DECL_OMP_DECLARE_MAPPER_P (member_function))
+ {
+ parser->lexer->in_pragma = true;
+ cp_parser_omp_declare_mapper_maplist (member_function, parser);
+ finish_function (/*inline_p=*/true);
+ cp_check_omp_declare_mapper (member_function);
+ /* If this is a template class, this forces the body of the mapper
+ to be instantiated. */
+ DECL_PRESERVE_P (member_function) = 1;
+ }
else
/* Now, parse the body of the function. */
cp_parser_function_definition_after_declarator (parser,
@@ -39488,10 +39502,9 @@ cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc)
always | close */
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;
@@ -39509,11 +39522,24 @@ 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)
+ && ((cp_lexer_peek_nth_token (parser->lexer, pos + 2)->type
+ == CPP_NAME)
+ || (cp_lexer_peek_nth_token (parser->lexer, pos + 2)->keyword
+ == RID_DEFAULT))
+ && (cp_lexer_peek_nth_token (parser->lexer, pos + 3)->type
+ == CPP_CLOSE_PAREN)
+ && (cp_lexer_peek_nth_token (parser->lexer, pos + 4)->type
+ == CPP_COMMA))
+ pos += 4;
pos++;
}
bool always_modifier = false;
bool close_modifier = false;
+ bool mapper_modifier = false;
+ tree mapper_name = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -39536,6 +39562,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)
{
@@ -39549,20 +39576,83 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
return list;
}
close_modifier = true;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else if (strcmp ("mapper", p) == 0)
+ {
+ cp_lexer_consume_token (parser->lexer);
+
+ matching_parens parens;
+ if (parens.require_open (parser))
+ {
+ if (mapper_modifier)
+ {
+ cp_parser_error (parser, "too many %<mapper%> modifiers");
+ /* Assume it's a well-formed mapper modifier, even if it
+ seems to be in the wrong place. */
+ cp_lexer_consume_token (parser->lexer);
+ parens.require_close (parser);
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/
+ true);
+ return list;
+ }
+
+ tok = cp_lexer_peek_token (parser->lexer);
+ switch (tok->type)
+ {
+ case CPP_NAME:
+ {
+ cp_expr e = cp_parser_identifier (parser);
+ if (e != error_mark_node)
+ mapper_name = e;
+ else
+ goto err;
+ }
+ break;
+
+ case CPP_KEYWORD:
+ if (tok->keyword == RID_DEFAULT)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ break;
+ }
+ /* Fallthrough. */
+
+ default:
+ err:
+ cp_parser_error (parser,
+ "expected identifier or %<default>%");
+ return list;
+ }
+
+ if (!parens.require_close (parser))
+ {
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/
+ true);
+ return list;
+ }
+
+ mapper_modifier = true;
+ pos += 3;
+ }
}
else
{
cp_parser_error (parser, "%<#pragma omp target%> with "
- "modifier other than %<always%> or %<close%>"
- "on %<map%> clause");
+ "modifier other than %<always%>, %<close%> "
+ "or %<mapper%> on %<map%> clause");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
/*or_comma=*/false,
/*consume_paren=*/true);
return list;
}
-
- cp_lexer_consume_token (parser->lexer);
}
if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
@@ -39603,8 +39693,30 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
NULL, true);
+ 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;
}
@@ -40405,7 +40517,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:
@@ -44549,6 +44661,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = nc;
}
+ clauses = omp_instantiate_mappers (clauses);
clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
c_omp_adjust_map_clauses (clauses, true);
@@ -46491,6 +46604,235 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok,
obstack_free (&declarator_obstack, p);
}
+/* OpenMP 5.0:
+ Parse a variable name and a list of map clauses for "omp declare mapper"
+ directives:
+
+ ... var) [clause[[,] clause] ... ] new-line */
+
+static bool
+cp_parser_omp_declare_mapper_maplist (tree fndecl, cp_parser *parser)
+{
+ pragma_omp_clause c_kind;
+ tree maplist = NULL_TREE, stmt = NULL_TREE;
+ tree mapper_name = NULL_TREE;
+ tree type = TREE_VALUE (TYPE_ARG_TYPES (TREE_TYPE (fndecl)));
+ tree id = cp_parser_declarator_id (parser, false);
+
+ if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ return false;
+
+ gcc_assert (TYPE_REF_P (type));
+ type = TREE_TYPE (type);
+
+ keep_next_level (true);
+ tree block = begin_omp_structured_block ();
+
+ tree var = build_lang_decl (VAR_DECL, id, type);
+ pushdecl (var);
+ cp_finish_decl (var, NULL_TREE, 0, NULL_TREE, 0);
+ DECL_ARTIFICIAL (var) = 1;
+ TREE_USED (var) = 1;
+
+ const char *fnname = IDENTIFIER_POINTER (DECL_NAME (fndecl));
+ if (startswith (fnname, "omp declare mapper "))
+ fnname += sizeof "omp declare mapper " - 1;
+ const char *mapname_end = strchr (fnname, '~');
+ if (mapname_end && mapname_end != fnname)
+ {
+ char *tmp = XALLOCAVEC (char, mapname_end - fnname + 1);
+ strncpy (tmp, fnname, mapname_end - fnname);
+ tmp[mapname_end - fnname] = '\0';
+ mapper_name = get_identifier (tmp);
+ }
+
+ while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+ {
+ c_kind = cp_parser_omp_clause_name (parser);
+ if (c_kind != PRAGMA_OMP_CLAUSE_MAP)
+ return false;
+ maplist = cp_parser_omp_clause_map (parser, maplist, GOMP_MAP_UNSET);
+ }
+
+ if (maplist == NULL_TREE)
+ {
+ cp_parser_error (parser, "missing %<map%> clause");
+ return false;
+ }
+
+ 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);
+
+ block = finish_omp_structured_block (block);
+
+ add_stmt (block);
+
+ return true;
+}
+
+/* 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;
+ cp_token *first_token = NULL;
+ cp_token_cache *cp = NULL;
+ tree type = NULL_TREE, fndecl = 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, fntype;
+
+ 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);
+
+ fntype = build_function_type_list (void_type_node,
+ cp_build_reference_type (type, false),
+ NULL_TREE);
+ fndecl = build_lang_decl (FUNCTION_DECL, mapper_id, fntype);
+ DECL_SOURCE_LOCATION (fndecl) = pragma_tok->location;
+ DECL_ARTIFICIAL (fndecl) = 1;
+ DECL_EXTERNAL (fndecl) = 1;
+ DECL_DECLARED_INLINE_P (fndecl) = 1;
+ DECL_IGNORED_P (fndecl) = 1;
+ DECL_OMP_DECLARE_MAPPER_P (fndecl) = 1;
+ SET_DECL_ASSEMBLER_NAME (fndecl, get_identifier ("<udm>"));
+ DECL_ATTRIBUTES (fndecl)
+ = tree_cons (get_identifier ("gnu_inline"), NULL_TREE,
+ DECL_ATTRIBUTES (fndecl));
+
+ if (block_scope)
+ block = begin_omp_structured_block ();
+
+ first_token = cp_lexer_peek_token (parser->lexer);
+
+ if (processing_template_decl)
+ fndecl = push_template_decl (fndecl);
+
+ if (block_scope)
+ {
+ DECL_CONTEXT (fndecl) = current_function_decl;
+ DECL_LOCAL_DECL_P (fndecl) = 1;
+ }
+ else if (current_class_type)
+ {
+ while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
+ cp_lexer_consume_token (parser->lexer);
+ cp = cp_token_cache_new (first_token,
+ cp_lexer_peek_nth_token (parser->lexer, 2));
+ DECL_STATIC_FUNCTION_P (fndecl) = 1;
+ finish_member_declaration (fndecl);
+ DECL_PENDING_INLINE_INFO (fndecl) = cp;
+ DECL_PENDING_INLINE_P (fndecl) = 1;
+ vec_safe_push (unparsed_funs_with_definitions, fndecl);
+ cp_parser_require_pragma_eol (parser, pragma_tok);
+ return;
+ }
+ else
+ {
+ DECL_CONTEXT (fndecl) = current_namespace;
+ tree d = pushdecl (fndecl);
+ gcc_checking_assert (d == error_mark_node || d == fndecl);
+
+ start_preparsed_function (fndecl, NULL_TREE, SF_PRE_PARSED);
+ }
+
+ if (!cp_parser_omp_declare_mapper_maplist (fndecl, parser))
+ goto fail;
+
+ if (!block_scope)
+ {
+ tree fn = finish_function (/*inline_p=*/false);
+ expand_or_defer_fn (fn);
+ }
+ else
+ {
+ DECL_CONTEXT (fndecl) = current_function_decl;
+ if (DECL_TEMPLATE_INFO (fndecl))
+ DECL_CONTEXT (DECL_TI_TEMPLATE (fndecl)) = current_function_decl;
+
+ block = finish_omp_structured_block (block);
+ if (TREE_CODE (block) == BIND_EXPR)
+ DECL_SAVED_TREE (fndecl) = BIND_EXPR_BODY (block);
+ else if (TREE_CODE (block) == STATEMENT_LIST)
+ DECL_SAVED_TREE (fndecl) = block;
+ if (processing_template_decl)
+ add_decl_expr (fndecl);
+ else
+ pushdecl (fndecl);
+ }
+
+ cp_check_omp_declare_mapper (fndecl);
+
+ 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) \
@@ -46531,6 +46873,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);
@@ -46544,7 +46892,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;
}
@@ -12093,9 +12093,13 @@ instantiate_class_template_1 (tree type)
/* Instantiate members marked with attribute used. */
if (r != error_mark_node && DECL_PRESERVE_P (r))
used.safe_push (r);
- if (TREE_CODE (r) == FUNCTION_DECL
- && DECL_OMP_DECLARE_REDUCTION_P (r))
- cp_check_omp_declare_reduction (r);
+ if (TREE_CODE (r) == FUNCTION_DECL)
+ {
+ if (DECL_OMP_DECLARE_REDUCTION_P (r))
+ cp_check_omp_declare_reduction (r);
+ else if (DECL_OMP_DECLARE_MAPPER_P (r))
+ cp_check_omp_declare_mapper (r);
+ }
}
else if ((DECL_CLASS_TEMPLATE_P (t) || DECL_IMPLICIT_TYPEDEF_P (t))
&& LAMBDA_TYPE_P (TREE_TYPE (t)))
@@ -14129,6 +14133,14 @@ tsubst_function_decl (tree t, tree args, tsubst_flags_t complain,
DECL_NAME (r) = omp_reduction_id (ERROR_MARK, DECL_NAME (t),
argtype);
}
+ else if (DECL_OMP_DECLARE_MAPPER_P (t))
+ {
+ tree argtype
+ = TREE_TYPE (TREE_VALUE (TYPE_ARG_TYPES (TREE_TYPE (t))));
+ argtype = tsubst (argtype, args, complain, in_decl);
+ if (strchr (IDENTIFIER_POINTER (DECL_NAME (t)), '~') == NULL)
+ DECL_NAME (r) = omp_mapper_id (DECL_NAME (t), argtype);
+ }
if (member && DECL_CONV_FN_P (r))
/* Type-conversion operator. Reconstruct the name, in
@@ -17852,6 +17864,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
new_clauses = nreverse (new_clauses);
if (ort != C_ORT_OMP_DECLARE_SIMD)
{
+ if (ort == C_ORT_OMP_TARGET)
+ new_clauses = 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))
@@ -18568,6 +18582,10 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
&& DECL_OMP_DECLARE_REDUCTION_P (decl)
&& cp_check_omp_declare_reduction (decl))
instantiate_body (pattern_decl, args, decl, true);
+ else if (TREE_CODE (decl) == FUNCTION_DECL
+ && DECL_OMP_DECLARE_MAPPER_P (decl)
+ && cp_check_omp_declare_mapper (decl))
+ instantiate_body (pattern_decl, args, decl, true);
}
else
{
@@ -19513,6 +19531,66 @@ tsubst_omp_udr (tree t, tree args, tsubst_flags_t complain, tree in_decl)
}
}
+static void
+tsubst_omp_udm (tree t, tree args, tsubst_flags_t complain, tree in_decl)
+{
+ if (t == NULL_TREE || t == error_mark_node)
+ return;
+
+ gcc_assert ((TREE_CODE (t) == STATEMENT_LIST
+ || TREE_CODE (t) == OMP_DECLARE_MAPPER)
+ && current_function_decl);
+
+ tree decl = NULL_TREE, mapper;
+
+ /* The function body is:
+
+ statement-list:
+ TYPE t;
+ #pragma omp declare mapper (TYPE t) map(...)
+ */
+
+ if (TREE_CODE (t) == STATEMENT_LIST)
+ {
+ tree_stmt_iterator tsi = tsi_start (t);
+ decl = tsi_stmt (tsi);
+ tsi_next (&tsi);
+ mapper = tsi_stmt (tsi);
+ tsi_next (&tsi);
+ gcc_assert (tsi_end_p (tsi));
+
+ gcc_assert (TREE_CODE (decl) == DECL_EXPR);
+
+ decl = tsubst (DECL_EXPR_DECL (decl), args, complain, in_decl);
+ pushdecl (decl);
+ }
+ else
+ fatal_error (input_location, "malformed OpenMP user-defined mapper");
+
+ if (TREE_CODE (mapper) == DECL_EXPR)
+ mapper = DECL_EXPR_DECL (mapper);
+
+ gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+ tree id = OMP_DECLARE_MAPPER_ID (mapper);
+ tree type = OMP_DECLARE_MAPPER_TYPE (mapper);
+ tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+
+ type = tsubst (type, args, complain, in_decl);
+ /* The _DECLARE_SIMD variant prevents calling finish_omp_clauses on the
+ substituted OMP clauses just yet. */
+ clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD, args,
+ complain, in_decl);
+
+ mapper = make_node (OMP_DECLARE_MAPPER);
+ OMP_DECLARE_MAPPER_ID (mapper) = id;
+ OMP_DECLARE_MAPPER_TYPE (mapper) = type;
+ OMP_DECLARE_MAPPER_DECL (mapper) = decl;
+ OMP_DECLARE_MAPPER_CLAUSES (mapper) = clauses;
+ SET_EXPR_LOCATION (mapper, EXPR_LOCATION (t));
+ add_stmt (mapper);
+}
+
/* T is a postfix-expression that is not being used in a function
call. Return the substituted version of T. */
@@ -26274,7 +26352,8 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p)
}
else
/* Only OMP reductions are nested. */
- gcc_checking_assert (DECL_OMP_DECLARE_REDUCTION_P (code_pattern));
+ gcc_checking_assert (DECL_OMP_DECLARE_REDUCTION_P (code_pattern)
+ || DECL_OMP_DECLARE_MAPPER_P (code_pattern));
vec<tree> omp_privatization_save;
if (current_function_decl)
@@ -26373,6 +26452,9 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p)
if (DECL_OMP_DECLARE_REDUCTION_P (code_pattern))
tsubst_omp_udr (DECL_SAVED_TREE (code_pattern), args,
tf_warning_or_error, d);
+ else if (DECL_OMP_DECLARE_MAPPER_P (code_pattern))
+ tsubst_omp_udm (DECL_SAVED_TREE (code_pattern), args,
+ tf_warning_or_error, d);
else
{
tsubst_expr (DECL_SAVED_TREE (code_pattern), args,
@@ -26400,6 +26482,8 @@ instantiate_body (tree pattern, tree args, tree d, bool nested_p)
if (DECL_OMP_DECLARE_REDUCTION_P (code_pattern))
cp_check_omp_declare_reduction (d);
+ else if (DECL_OMP_DECLARE_MAPPER_P (code_pattern))
+ cp_check_omp_declare_mapper (d);
}
/* We're not deferring instantiation any more. */
@@ -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
@@ -4729,7 +4730,8 @@ expand_or_defer_fn_1 (tree fn)
be handled. */;
else if (!at_eof
|| DECL_IMMEDIATE_FUNCTION_P (fn)
- || DECL_OMP_DECLARE_REDUCTION_P (fn))
+ || DECL_OMP_DECLARE_REDUCTION_P (fn)
+ || DECL_OMP_DECLARE_MAPPER_P (fn))
tentative_decl_linkage (fn);
else
import_export_decl (fn);
@@ -4742,6 +4744,7 @@ expand_or_defer_fn_1 (tree fn)
&& !DECL_REALLY_EXTERN (fn)
&& !DECL_IMMEDIATE_FUNCTION_P (fn)
&& !DECL_OMP_DECLARE_REDUCTION_P (fn)
+ && !DECL_OMP_DECLARE_MAPPER_P (fn)
&& (flag_keep_inline_functions
|| (flag_keep_inline_dllexport
&& lookup_attribute ("dllexport", DECL_ATTRIBUTES (fn)))))
@@ -4775,7 +4778,8 @@ expand_or_defer_fn_1 (tree fn)
return false;
}
- if (DECL_OMP_DECLARE_REDUCTION_P (fn))
+ if (DECL_OMP_DECLARE_REDUCTION_P (fn)
+ || DECL_OMP_DECLARE_MAPPER_P (fn))
return false;
return true;
@@ -5959,6 +5963,76 @@ 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);
+}
+
+static tree
+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);
+}
+
+static tree
+omp_extract_mapper_directive (tree fndecl)
+{
+ if (BASELINK_P (fndecl))
+ /* See through BASELINK nodes to the underlying function. */
+ fndecl = BASELINK_FUNCTIONS (fndecl);
+
+ tree body = DECL_SAVED_TREE (fndecl);
+
+ if (TREE_CODE (body) == BIND_EXPR)
+ body = BIND_EXPR_BODY (body);
+
+ 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;
+}
+
/* Helper function for cp_parser_omp_declare_reduction_exprs
and tsubst_omp_udr.
Remove CLEANUP_STMT for data (omp_priv variable).
@@ -6440,6 +6514,31 @@ 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_VALUE (TYPE_ARG_TYPES (TREE_TYPE (udm)));
+ location_t loc = DECL_SOURCE_LOCATION (udm);
+ gcc_assert (TYPE_REF_P (type));
+ type = TREE_TYPE (type);
+
+ 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. */
@@ -6700,6 +6799,242 @@ 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. */
@@ -7942,6 +8277,13 @@ 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))
@@ -9286,6 +9628,108 @@ 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
@@ -9311,6 +9755,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;
+
+ mapper_list *mappers;
};
/* Helper function of finish_omp_target_clauses, called via
@@ -9324,6 +9770,8 @@ 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;
+ tree aggr_type = NULL_TREE;
/* References inside of these expression codes shouldn't incur any
form of mapping, so return early. */
@@ -9337,6 +9785,22 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
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 = 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);
@@ -9440,10 +9904,38 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
else
data.current_closure = NULL_TREE;
+ hash_set<omp_name_type> seen_types;
+ auto_vec<tree> mapper_fns;
+ mapper_list 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)
+ 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);
+ 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);
+ }
+
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> *implicit_mappers;
hash_set<tree> *privatized_types;
tree clauses;
/* Iteration variables in an OMP_FOR. */
@@ -451,6 +452,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>;
c->privatized_types = new hash_set<tree>;
c->location = input_location;
c->region_type = region_type;
@@ -474,6 +476,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);
}
@@ -10350,23 +10353,176 @@ error_out:
return success;
}
-/* Scan the OMP clauses in *LIST_P, installing mappings into a new
- and previous omp contexts. */
-
-static void
-gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
- enum omp_region_type region_type,
- enum tree_code code)
+struct instantiate_mapper_info
{
- struct gimplify_omp_ctx *ctx, *outer_ctx;
- tree c;
- tree *prev_list_p = NULL, *orig_list_p = list_p;
- int handled_depend_iterators = -1;
- int nowait = -1;
+ vec<omp_mapping_group> *groups;
+ hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
+ tree *mapper_clauses_p;
+ struct gimplify_omp_ctx *omp_ctx;
+};
- ctx = new_omp_context (region_type);
+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 = unshare_expr (map_info->expr);
+ *walk_subtrees = 0;
+ }
+
+ return NULL_TREE;
+}
+
+static tree *
+omp_instantiate_mapper (hash_map<omp_name_type, tree> *implicit_mappers,
+ tree mapper, tree expr, enum gomp_map_kind outer_kind,
+ tree *mapper_clauses_p)
+{
+ tree mapper_name = NULL_TREE;
+ gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+
+ tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
+ tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
+
+ remap_mapper_decl_info map_info;
+
+ map_info.dummy_var = dummy_var;
+ map_info.expr = expr;
+
+ for (; clause; clause = OMP_CLAUSE_CHAIN (clause))
+ {
+ enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause);
+
+ if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
+ {
+ mapper_name = OMP_CLAUSE_DECL (clause);
+ continue;
+ }
+ else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
+ {
+ mapper_name = NULL_TREE;
+ continue;
+ }
+
+ tree decl = OMP_CLAUSE_DECL (clause), unshared;
+
+ if (TREE_CODE (decl) == OMP_ARRAY_SECTION
+ && TREE_OPERAND (decl, 2)
+ && integer_onep (TREE_OPERAND (decl, 2)))
+ {
+ 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));
+ else
+ OMP_CLAUSE_DECL (unshared) = decl;
+ OMP_CLAUSE_SIZE (unshared) = OMP_CLAUSE_SIZE (clause);
+ }
+ else
+ unshared = unshare_expr (clause);
+
+ 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);
+
+ decl = OMP_CLAUSE_DECL (unshared);
+ tree type = TYPE_MAIN_VARIANT (TREE_TYPE (decl));
+
+ tree *nested_mapper_p = implicit_mappers->get ({ mapper_name, type });
+
+ if (nested_mapper_p && *nested_mapper_p != mapper)
+ {
+ if (clause_kind == GOMP_MAP_UNSET)
+ clause_kind = outer_kind;
+
+ mapper_clauses_p
+ = omp_instantiate_mapper (implicit_mappers, *nested_mapper_p,
+ decl, clause_kind, mapper_clauses_p);
+ continue;
+ }
+
+ *mapper_clauses_p = unshared;
+ mapper_clauses_p = &OMP_CLAUSE_CHAIN (unshared);
+ }
+
+ return mapper_clauses_p;
+}
+
+static bool
+handled_struct_p (vec<omp_mapping_group> *groups, tree s)
+{
+ if (!groups)
+ return false;
+
+ for (auto &i : *groups)
+ {
+ tree node = *i.grp_start;
+ if (OMP_CLAUSE_CODE (node) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_STRUCT
+ && OMP_CLAUSE_DECL (node) == s)
+ return true;
+ }
+ return false;
+}
+
+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;
+
+ 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 });
+ }
+
+ bool handled_p = handled_struct_p (im_info->groups, decl);
+
+ if (mapper_p && !handled_p)
+ {
+ /* If we have a reference, map the pointed-to object rather than the
+ reference itself. */
+ if (ref_p)
+ decl = build_fold_indirect_ref (decl);
+
+ im_info->mapper_clauses_p
+ = omp_instantiate_mapper (ctx->implicit_mappers, *mapper_p, decl,
+ GOMP_MAP_TOFROM, im_info->mapper_clauses_p);
+ }
+
+ return 0;
+}
+
+static struct gimplify_omp_ctx *
+new_omp_context_for_scan (enum omp_region_type region_type,
+ enum tree_code code)
+{
+ struct gimplify_omp_ctx *ctx = new_omp_context (region_type);
ctx->code = code;
- outer_ctx = ctx->outer_context;
if (code == OMP_TARGET)
{
if (!lang_GNU_Fortran ())
@@ -10390,6 +10546,34 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
default:
break;
}
+ gimplify_omp_ctxp = ctx;
+ return ctx;
+}
+
+/* Scan the OMP clauses in *LIST_P, installing mappings into a new
+ and previous omp contexts. */
+
+static void
+gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
+ enum omp_region_type region_type,
+ enum tree_code code,
+ struct gimplify_omp_ctx *existing_ctx = NULL,
+ bool use_mappers = false)
+{
+ struct gimplify_omp_ctx *ctx, *outer_ctx;
+ tree c;
+ tree *prev_list_p = NULL, *orig_list_p = list_p;
+ int handled_depend_iterators = -1;
+ int nowait = -1;
+ tree mapper_clauses = NULL_TREE;
+
+ if (existing_ctx)
+ ctx = existing_ctx;
+ else
+ ctx = new_omp_context_for_scan (region_type, code);
+
+ outer_ctx = ctx->outer_context;
+ gimplify_omp_ctxp = outer_ctx;
if (code == OMP_TARGET
|| code == OMP_TARGET_DATA
@@ -10398,11 +10582,50 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
{
vec<omp_mapping_group> *groups;
groups = omp_gather_mapping_groups (list_p);
+ hash_map<tree_operand_hash, omp_mapping_group *> *grpmap = NULL;
+
+ if (groups)
+ grpmap = omp_index_mapping_groups (groups);
+
+ if (use_mappers)
+ {
+ instantiate_mapper_info im_info;
+
+ im_info.groups = groups;
+ im_info.grpmap = grpmap;
+ im_info.mapper_clauses_p = &mapper_clauses;
+ im_info.omp_ctx = ctx;
+
+ 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);
+
+ if (groups)
+ {
+ delete grpmap;
+ delete groups;
+ grpmap = NULL;
+ }
+
+ if (!mapper_clauses)
+ {
+ gimplify_omp_ctxp = ctx;
+ return;
+ }
+
+ list_p = &mapper_clauses;
+ groups = omp_gather_mapping_groups (list_p);
+
+ if (groups)
+ grpmap = omp_index_mapping_groups (groups);
+ }
+
if (groups)
{
- hash_map<tree_operand_hash, omp_mapping_group *> *grpmap;
- grpmap = omp_index_mapping_groups (groups);
-
omp_build_struct_sibling_lists (code, region_type, groups, &grpmap);
omp_mapping_group *outlist = NULL;
@@ -10445,6 +10668,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
}
+ gcc_assert (!use_mappers || mapper_clauses);
+
while ((c = *list_p) != NULL)
{
bool remove = false;
@@ -11200,6 +11425,29 @@ 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);
+ 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);
+ }
+ gcc_assert (mapper != NULL_TREE
+ && TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
+ ctx->implicit_mappers->put ({ name, type }, mapper);
+ remove = true;
+ break;
+ }
+
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
flags = GOVD_EXPLICIT;
@@ -11739,7 +11987,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
list_p = &OMP_CLAUSE_CHAIN (c);
}
- ctx->clauses = *orig_list_p;
+ if (mapper_clauses)
+ {
+ /* Put clauses implicitly created from "declare mapper" mappers
+ at the front of the clause list. */
+ *list_p = *orig_list_p;
+ *orig_list_p = ctx->clauses = mapper_clauses;
+ }
+ else
+ ctx->clauses = *orig_list_p;
gimplify_omp_ctxp = ctx;
}
@@ -14913,8 +15169,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
bool save_in_omp_construct = in_omp_construct;
if ((ort & ORT_ACC) == 0)
in_omp_construct = false;
+ struct gimplify_omp_ctx *ctx
+ = new_omp_context_for_scan (ort, TREE_CODE (expr));
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
- TREE_CODE (expr));
+ TREE_CODE (expr), ctx, false);
if (TREE_CODE (expr) == OMP_TARGET)
optimize_target_teams (expr, pre_p);
if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -14952,6 +15210,12 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
}
else
gimplify_and_add (OMP_BODY (expr), &body);
+ if (TREE_CODE (expr) == OMP_TARGET
+ || TREE_CODE (expr) == OMP_TARGET_DATA
+ || TREE_CODE (expr) == OMP_TARGET_ENTER_DATA
+ || TREE_CODE (expr) == OMP_TARGET_EXIT_DATA)
+ gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
+ TREE_CODE (expr), ctx, true);
gimplify_adjust_omp_clauses (pre_p, body, &OMP_CLAUSES (expr),
TREE_CODE (expr));
in_omp_construct = save_in_omp_construct;
@@ -15611,6 +15875,16 @@ gimplify_omp_ordered (tree expr, gimple_seq body)
return gimple_build_omp_ordered (body, OMP_ORDERED_CLAUSES (expr));
}
+static enum gimplify_status
+gimplify_omp_declare_mapper (tree *expr_p)
+{
+ /* We don't want assembler output -- this inhibits it. */
+ DECL_DECLARED_INLINE_P (current_function_decl) = 1;
+
+ *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
@@ -16520,6 +16794,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;
@@ -84,6 +84,7 @@ extern enum omp_clause_default_kind lhd_omp_predetermined_sharing (tree);
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);
struct gimplify_omp_ctx;
extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
tree);
@@ -270,6 +271,7 @@ 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_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
@@ -303,6 +305,7 @@ 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_ALLOCATABLE_P, \
LANG_HOOKS_OMP_SCALAR_P, \
LANG_HOOKS_OMP_SCALAR_TARGET_P, \
@@ -634,6 +634,15 @@ lhd_omp_finish_clause (tree, gimple_seq *, bool)
{
}
+/* Finalize clause list C after expanding custom mappers for implicitly-mapped
+ variables. */
+
+tree
+lhd_omp_finish_mapper_clauses (tree c)
+{
+ return c;
+}
+
/* 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. */
@@ -306,6 +306,10 @@ 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);
+
/* Return true if DECL is an allocatable variable (for the purpose of
implicit mapping). */
bool (*omp_allocatable_p) (tree decl);
@@ -149,4 +149,56 @@ get_openacc_privatization_dump_flags ()
return l_dump_flags;
}
+struct omp_name_type
+{
+ tree name;
+ tree type;
+};
+
+template <>
+struct default_hash_traits <omp_name_type>
+ : typed_noop_remove <omp_name_type>
+{
+ GTY((skip)) typedef omp_name_type value_type;
+ GTY((skip)) typedef omp_name_type compare_type;
+
+ static hashval_t
+ hash (omp_name_type 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 p)
+ {
+ return p.type == NULL;
+ }
+
+ static bool
+ is_deleted (omp_name_type)
+ {
+ return false;
+ }
+
+ static bool
+ equal (const omp_name_type &a, const omp_name_type &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 &e)
+ {
+ e.type = NULL;
+ }
+};
+
#endif /* GCC_OMP_GENERAL_H */
@@ -13,10 +13,12 @@ foo (void)
#pragma omp target map (to:a)
;
- #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+ #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */
+/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */
;
- #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+ #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always' or 'close'" "" { target c } } */
+/* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'mapper'" "" { target c++ } .-1 } */
;
#pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */
new file mode 100644
@@ -0,0 +1,58 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+// "omp declare mapper" support -- check expansion in gimple.
+
+struct S {
+ int *ptr;
+ int size;
+};
+
+#define N 64
+
+#pragma omp declare mapper (S w) map(w.size, w.ptr, w.ptr[:w.size])
+#pragma omp declare mapper (foo:S w) map(to:w.size, w.ptr) map(w.ptr[:w.size])
+
+int main (int argc, char *argv[])
+{
+ S s;
+ s.ptr = new int[N];
+ s.size = N;
+
+#pragma omp declare mapper (bar:S w) map(w.size, w.ptr, w.ptr[:w.size])
+
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+#pragma omp target map(tofrom: s)
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+#pragma omp target map(mapper(default), tofrom: s)
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+#pragma omp target map(mapper(foo), alloc: s)
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+#pragma omp target map(mapper(bar), tofrom: s)
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+ return 0;
+}
+
+// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(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" } }
+// { 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" } }
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])
+
+ /* 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'" }
+
+ /* ...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;
+}
new file mode 100644
@@ -0,0 +1,27 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+// Test named mapper invocation.
+
+struct S {
+ int *ptr;
+ int size;
+};
+
+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])
+
+ S s;
+ s.ptr = new 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" } }
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,74 @@
+// { 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" } }
@@ -342,6 +342,10 @@ enum omp_clause_code {
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
+ /* 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_,
@@ -940,6 +940,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, "attach_zero_length_array_section");
break;
+ case GOMP_MAP_UNSET:
+ pp_string (pp, "unset");
+ break;
+ case GOMP_MAP_PUSH_MAPPER_NAME:
+ pp_string (pp, "push_mapper");
+ break;
+ case GOMP_MAP_POP_MAPPER_NAME:
+ pp_string (pp, "pop_mapper");
+ break;
default:
gcc_unreachable ();
}
@@ -1003,6 +1012,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))
@@ -3769,6 +3795,22 @@ 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, OMP_DECLARE_MAPPER_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);
+ is_expr = false;
+ break;
+
case TRANSACTION_EXPR:
if (TRANSACTION_EXPR_OUTER (node))
pp_string (pp, "__transaction_atomic [[outer]]");
@@ -289,6 +289,7 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
+ 3, /* OMP_CLAUSE__MAPPER_BINDING_ */
2, /* OMP_CLAUSE__CACHE_ */
2, /* OMP_CLAUSE_GANG */
1, /* OMP_CLAUSE_ASYNC */
@@ -378,6 +379,7 @@ const char * const omp_clause_code_name[] =
"from",
"to",
"map",
+ "_mapper_binding_",
"_cache_",
"gang",
"async",
@@ -1247,6 +1247,13 @@ DEFTREECODE (OMP_SECTION, "omp_section", tcc_statement, 1)
Operand 0: OMP_MASTER_BODY: Master section body. */
DEFTREECODE (OMP_MASTER, "omp_master", tcc_statement, 1)
+/* OpenMP - #pragma omp declare mapper ([id:] type var) [clause1 ... clauseN]
+ Operand 0: Identifier.
+ Operand 1: Type.
+ Operand 2: Variable decl.
+ Operand 3: List of clauses. */
+DEFTREECODE (OMP_DECLARE_MAPPER, "omp_declare_mapper", tcc_statement, 4)
+
/* OpenACC - #pragma acc cache (variable1 ... variableN)
Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into
OMP_CLAUSE__CACHE_ clauses). */
@@ -1475,6 +1475,15 @@ 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_TYPE(NODE) \
+ TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 1)
+#define OMP_DECLARE_MAPPER_DECL(NODE) \
+ TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 2)
+#define OMP_DECLARE_MAPPER_CLAUSES(NODE) \
+ TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 3)
+
#define OMP_SCAN_BODY(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0)
#define OMP_SCAN_CLAUSES(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1)
@@ -1885,6 +1894,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
@@ -182,7 +182,13 @@ enum gomp_map_kind
/* An attach or detach operation. Rewritten to the appropriate type during
gimplification, depending on directive (i.e. "enter data" or
parallel/kernels region vs. "exit data"). */
- GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3)
+ GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3),
+ /* Unset, used for "declare mapper" maps with no explicit data movement
+ specified. These use the movement specified at the invocation site. */
+ GOMP_MAP_UNSET = (GOMP_MAP_LAST | 4),
+ /* Used to record the name of a named mapper. */
+ GOMP_MAP_PUSH_MAPPER_NAME = (GOMP_MAP_LAST | 5),
+ GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 6)
};
#define GOMP_MAP_COPY_TO_P(X) \
new file mode 100644
@@ -0,0 +1,87 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+#define N 64
+
+struct points
+{
+ double *x;
+ double *y;
+ double *z;
+ size_t len;
+};
+
+#pragma omp declare mapper(points p) map(to:p.x, p.y, p.z) \
+ map(p.x[0:p.len]) \
+ map(p.y[0:p.len]) \
+ map(p.z[0:p.len])
+
+struct shape
+{
+ points tmp;
+ points *pts;
+ int metadata[128];
+};
+
+#pragma omp declare mapper(shape s) map(tofrom:s.pts, *s.pts) map(alloc:s.tmp)
+
+void
+alloc_points (points *pts, size_t sz)
+{
+ pts->x = new double[sz];
+ pts->y = new double[sz];
+ pts->z = new double[sz];
+ pts->len = sz;
+ for (int i = 0; i < sz; i++)
+ pts->x[i] = pts->y[i] = pts->z[i] = 0;
+}
+
+int main (int argc, char *argv[])
+{
+ shape myshape;
+ points mypts;
+
+ myshape.pts = &mypts;
+
+ alloc_points (&myshape.tmp, N);
+ myshape.pts = new points;
+ alloc_points (myshape.pts, N);
+
+ #pragma omp target map(myshape)
+ {
+ for (int i = 0; i < N; i++)
+ {
+ myshape.pts->x[i]++;
+ myshape.pts->y[i]++;
+ myshape.pts->z[i]++;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (myshape.pts->x[i] == 1);
+ assert (myshape.pts->y[i] == 1);
+ assert (myshape.pts->z[i] == 1);
+ }
+
+ #pragma omp target
+ {
+ for (int i = 0; i < N; i++)
+ {
+ myshape.pts->x[i]++;
+ myshape.pts->y[i]++;
+ myshape.pts->z[i]++;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ {
+ assert (myshape.pts->x[i] == 2);
+ assert (myshape.pts->y[i] == 2);
+ assert (myshape.pts->z[i] == 2);
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,55 @@
+// { dg-do run }
+
+#include <cassert>
+
+#define N 256
+
+struct doublebuf
+{
+ int buf_a[N][N];
+ int buf_b[N][N];
+};
+
+#pragma omp declare mapper(lo:doublebuf b) map(b.buf_a[0:N/2][0:N]) \
+ map(b.buf_b[0:N/2][0:N])
+
+#pragma omp declare mapper(hi:doublebuf b) map(b.buf_a[N/2:N/2][0:N]) \
+ map(b.buf_b[N/2:N/2][0:N])
+
+int main (int argc, char *argv[])
+{
+ doublebuf db;
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ db.buf_a[i][j] = db.buf_b[i][j] = 0;
+
+ #pragma omp target map(mapper(lo), tofrom:db)
+ {
+ for (int i = 0; i < N / 2; i++)
+ for (int j = 0; j < N; j++)
+ {
+ db.buf_a[i][j]++;
+ db.buf_b[i][j]++;
+ }
+ }
+
+ #pragma omp target map(mapper(hi), tofrom:db)
+ {
+ for (int i = N / 2; i < N; i++)
+ for (int j = 0; j < N; j++)
+ {
+ db.buf_a[i][j]++;
+ db.buf_b[i][j]++;
+ }
+ }
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ {
+ assert (db.buf_a[i][j] == 1);
+ assert (db.buf_b[i][j] == 1);
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,63 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+struct S {
+ int *myarr;
+};
+
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20])
+
+namespace A {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100])
+}
+
+namespace B {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100])
+}
+
+namespace A
+{
+ void incr_a (S my_s)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 100; i++)
+ my_s.myarr[i]++;
+ }
+ }
+}
+
+namespace B
+{
+ void incr_b (S my_s)
+ {
+#pragma omp target
+ {
+ for (int i = 100; i < 200; i++)
+ my_s.myarr[i]++;
+ }
+ }
+}
+
+int main (int argc, char *argv[])
+{
+ S my_s;
+
+ my_s.myarr = (int *) calloc (200, sizeof (int));
+
+#pragma omp target
+ {
+ for (int i = 0; i < 20; i++)
+ my_s.myarr[i]++;
+ }
+
+ A::incr_a (my_s);
+ B::incr_b (my_s);
+
+ for (int i = 0; i < 200; i++)
+ assert (my_s.myarr[i] == (i < 20) ? 2 : 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,63 @@
+// { dg-do run }
+
+#include <cstdlib>
+#include <cassert>
+
+struct S {
+ int *myarr;
+};
+
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20])
+
+namespace A {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100])
+}
+
+namespace B {
+#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100])
+}
+
+namespace A
+{
+ void incr_a (S &my_s)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 100; i++)
+ my_s.myarr[i]++;
+ }
+ }
+}
+
+namespace B
+{
+ void incr_b (S &my_s)
+ {
+#pragma omp target
+ {
+ for (int i = 100; i < 200; i++)
+ my_s.myarr[i]++;
+ }
+ }
+}
+
+int main (int argc, char *argv[])
+{
+ S my_s;
+
+ my_s.myarr = (int *) calloc (200, sizeof (int));
+
+#pragma omp target
+ {
+ for (int i = 0; i < 20; i++)
+ my_s.myarr[i]++;
+ }
+
+ A::incr_a (my_s);
+ B::incr_b (my_s);
+
+ for (int i = 0; i < 200; i++)
+ assert (my_s.myarr[i] == (i < 20) ? 2 : 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,52 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+ int *myarr;
+ int len;
+};
+
+class C
+{
+ S smemb;
+#pragma omp declare mapper (custom:S s) map(to:s.myarr) \
+ map(tofrom:s.myarr[0:s.len])
+
+public:
+ C(int l)
+ {
+ smemb.myarr = new int[l];
+ smemb.len = l;
+ for (int i = 0; i < l; i++)
+ smemb.myarr[i] = 0;
+ }
+ void bump();
+ void check();
+};
+
+void
+C::bump ()
+{
+#pragma omp target map(mapper(custom), tofrom: smemb)
+ {
+ for (int i = 0; i < smemb.len; i++)
+ smemb.myarr[i]++;
+ }
+}
+
+void
+C::check ()
+{
+ for (int i = 0; i < smemb.len; i++)
+ assert (smemb.myarr[i] == 1);
+}
+
+int main (int argc, char *argv[])
+{
+ C test (100);
+ test.bump ();
+ test.check ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+// { dg-do run }
+
+#include <cassert>
+
+template <typename T>
+void adjust (T param)
+{
+#pragma omp declare mapper (T x) map(to:x.len, x.base) \
+ map(tofrom:x.base[0:x.len])
+
+#pragma omp target
+ for (int i = 0; i < param.len; i++)
+ param.base[i]++;
+}
+
+struct S {
+ int len;
+ int *base;
+};
+
+int main (int argc, char *argv[])
+{
+ S a;
+
+ a.len = 100;
+ a.base = new int[a.len];
+
+ for (int i = 0; i < a.len; i++)
+ a.base[i] = 0;
+
+ adjust (a);
+
+ for (int i = 0; i < a.len; i++)
+ assert (a.base[i] == 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,48 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+ int *myarr;
+};
+
+struct T
+{
+ S *s;
+};
+
+#pragma omp declare mapper (s100: S x) map(to: x.myarr) \
+ map(tofrom: x.myarr[0:100])
+
+void
+bump (T t)
+{
+ /* Here we have an implicit/default mapper invoking a named mapper. We
+ need to make sure that can be located properly at gimplification
+ time. */
+#pragma omp declare mapper (T t) map(to:t.s) map(mapper(s100), tofrom: t.s[0])
+
+#pragma omp target
+ for (int i = 0; i < 100; i++)
+ t.s->myarr[i]++;
+}
+
+int main (int argc, char *argv[])
+{
+ S my_s;
+ T my_t;
+
+ my_s.myarr = new int[100];
+ my_t.s = &my_s;
+
+ for (int i = 0; i < 100; i++)
+ my_s.myarr[i] = 0;
+
+ bump (my_t);
+
+ for (int i = 0; i < 100; i++)
+ assert (my_s.myarr[i] == 1);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,61 @@
+// { dg-do run }
+
+#include <cassert>
+
+struct S
+{
+ int *myarr;
+ int len;
+};
+
+template<typename T>
+class C
+{
+ T memb;
+#pragma omp declare mapper (T t) map(to:t.len, t.myarr) \
+ map(tofrom:t.myarr[0:t.len])
+
+public:
+ C(int sz);
+ ~C();
+ void bump();
+ void check();
+};
+
+template<typename T>
+C<T>::C(int sz)
+{
+ memb.myarr = new int[sz];
+ for (int i = 0; i < sz; i++)
+ memb.myarr[i] = 0;
+ memb.len = sz;
+}
+
+template<typename T>
+C<T>::~C()
+{
+ delete[] memb.myarr;
+}
+
+template<typename T>
+void C<T>::bump()
+{
+#pragma omp target map(memb)
+ for (int i = 0; i < memb.len; i++)
+ memb.myarr[i]++;
+}
+
+template<typename T>
+void C<T>::check()
+{
+ for (int i = 0; i < memb.len; i++)
+ assert (memb.myarr[i] == 1);
+}
+
+int main(int argc, char *argv[])
+{
+ C<S> c_int(100);
+ c_int.bump();
+ c_int.check();
+ return 0;
+}