@@ -521,6 +521,8 @@ const struct attribute_spec c_common_gnu_attributes[] =
handle_omp_declare_variant_attribute, NULL },
{ "omp declare variant variant", 0, -1, true, false, false, false,
handle_omp_declare_variant_attribute, NULL },
+ { "omp declare variant overload", 0, -1, true, false, false, false,
+ handle_omp_declare_variant_attribute, NULL },
{ "simd", 0, 1, true, false, false, false,
handle_simd_attribute, NULL },
{ "omp declare target", 0, -1, true, false, false, false,
@@ -4020,8 +4022,8 @@ handle_omp_declare_simd_attribute (tree *, tree, tree, int, bool *)
return NULL_TREE;
}
-/* Handle an "omp declare variant {base,variant}" attribute; arguments as in
- struct attribute_spec.handler. */
+/* Handle an "omp declare variant {base,variant,overload}" attribute;
+ arguments as in struct attribute_spec.handler. */
static tree
handle_omp_declare_variant_attribute (tree *, tree, tree, int, bool *)
@@ -646,6 +646,7 @@ coverage_begin_function (unsigned lineno_checksum, unsigned cfg_checksum)
(DECL_ASSEMBLER_NAME (current_function_decl)));
gcov_write_unsigned (DECL_ARTIFICIAL (current_function_decl)
&& !DECL_FUNCTION_VERSIONED (current_function_decl)
+ && !DECL_FUNCTION_OMP_VARIANT (current_function_decl)
&& !DECL_LAMBDA_FUNCTION_P (current_function_decl));
gcov_write_filename (remap_profile_filename (startloc.file));
gcov_write_unsigned (startloc.line);
@@ -13093,6 +13093,29 @@ joust (struct z_candidate *cand1, struct z_candidate *cand2, bool warn,
}
}
+ if (TREE_CODE (cand1->fn) == FUNCTION_DECL
+ && DECL_FUNCTION_OMP_VARIANT (cand1->fn)
+ && TREE_CODE (cand2->fn) == FUNCTION_DECL
+ && DECL_FUNCTION_OMP_VARIANT (cand2->fn))
+ {
+ tree f1 = TREE_TYPE (cand1->fn);
+ tree f2 = TREE_TYPE (cand2->fn);
+ tree p1 = TYPE_ARG_TYPES (f1);
+ tree p2 = TYPE_ARG_TYPES (f2);
+ if (compparms (p1, p2)
+ && same_type_p (TREE_TYPE (f1), TREE_TYPE (f2)))
+ {
+ tree a1 = lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (cand1->fn));
+ tree a2 = lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (cand2->fn));
+ if (!a1 && a2)
+ return 1;
+ else if (a1 && !a2)
+ return -1;
+ }
+ }
+
/* If the two function declarations represent the same function (this can
happen with declarations in multiple scopes and arg-dependent lookup),
arbitrarily choose one. But first make sure the default args we're
@@ -1180,6 +1180,18 @@ add_method (tree type, tree method, bool via_using)
&& maybe_version_functions (method, fn, true))
continue;
+ if (flag_openmp
+ && (lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (method))
+ || lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (fn))))
+ {
+ bool record = (!DECL_FUNCTION_OMP_VARIANT (method)
+ || !DECL_FUNCTION_OMP_VARIANT (fn));
+ maybe_omp_variant_functions (method, fn, record);
+ continue;
+ }
+
if (DECL_INHERITED_CTOR (method))
{
if (!DECL_INHERITED_CTOR (fn))
@@ -8759,6 +8771,9 @@ resolve_address_of_overloaded_function (tree target_type,
mark_versions_used (fn);
}
+ // TBD.
+ gcc_assert (!DECL_FUNCTION_OMP_VARIANT (fn));
+
/* If we're doing overload resolution purely for the purpose of
determining conversion sequences, we should not consider the
function used. If this conversion sequence is selected, the
@@ -1845,6 +1845,11 @@ struct GTY(()) cp_omp_begin_assumes_data {
bool attr_syntax;
};
+struct GTY(()) cp_omp_declare_variant_attr {
+ bool attr_syntax;
+ tree selector;
+};
+
/* Global state. */
struct GTY(()) saved_scope {
@@ -1894,6 +1899,7 @@ struct GTY(()) saved_scope {
hash_map<tree, tree> *GTY((skip)) x_local_specializations;
vec<cp_omp_declare_target_attr, va_gc> *omp_declare_target_attribute;
vec<cp_omp_begin_assumes_data, va_gc> *omp_begin_assumes;
+ vec<cp_omp_declare_variant_attr, va_gc> *omp_declare_variant_attribute;
struct saved_scope *prev;
};
@@ -6937,6 +6943,7 @@ extern bool member_like_constrained_friend_p (tree);
extern bool fns_correspond (tree, tree);
extern int decls_match (tree, tree, bool = true);
extern bool maybe_version_functions (tree, tree, bool);
+extern void maybe_omp_variant_functions (tree, tree, bool);
extern bool merge_default_template_args (tree, tree, bool);
extern tree duplicate_decls (tree, tree,
bool hiding = false,
@@ -1180,6 +1180,21 @@ decls_match (tree newdecl, tree olddecl, bool record_versions /* = true */)
|| !DECL_FUNCTION_VERSIONED (olddecl)));
return 0;
}
+ if (flag_openmp
+ && types_match
+ && (lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (newdecl))
+ || lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (olddecl))))
+ {
+ if (record_versions)
+ {
+ bool record = (!DECL_FUNCTION_OMP_VARIANT (newdecl)
+ || !DECL_FUNCTION_OMP_VARIANT (olddecl));
+ maybe_omp_variant_functions (newdecl, olddecl, record);
+ }
+ return 0;
+ }
}
else if (TREE_CODE (newdecl) == TEMPLATE_DECL)
{
@@ -1243,6 +1258,17 @@ maybe_mark_function_versioned (tree decl)
}
}
+static void
+maybe_mark_omp_variant_function (tree decl)
+{
+ if (!DECL_FUNCTION_OMP_VARIANT (decl))
+ {
+ DECL_FUNCTION_OMP_VARIANT (decl) = 1;
+ if (DECL_ASSEMBLER_NAME_SET_P (decl))
+ mangle_decl (decl);
+ }
+}
+
/* NEWDECL and OLDDECL have identical signatures. If they are
different versions adjust them and return true.
If RECORD is set to true, record function versions. */
@@ -1277,6 +1303,116 @@ maybe_version_functions (tree newdecl, tree olddecl, bool record)
return true;
}
+void
+maybe_omp_variant_functions (tree newdecl, tree olddecl, bool record)
+{
+ maybe_mark_omp_variant_function (olddecl);
+ if (DECL_LOCAL_DECL_P (olddecl))
+ {
+ /* Is this meaningful for these functions? */
+ olddecl = DECL_LOCAL_DECL_ALIAS (olddecl);
+ maybe_mark_omp_variant_function (olddecl);
+ }
+
+ maybe_mark_omp_variant_function (newdecl);
+ if (DECL_LOCAL_DECL_P (newdecl))
+ {
+ /* again? */
+ if (!DECL_LOCAL_DECL_ALIAS (newdecl))
+ {
+ if (!DECL_LOCAL_DECL_ALIAS (newdecl))
+ push_local_extern_decl_alias (newdecl);
+ newdecl = DECL_LOCAL_DECL_ALIAS (newdecl);
+ maybe_mark_omp_variant_function (newdecl);
+ }
+ }
+
+ if (record)
+ cgraph_node::record_function_versions (olddecl, newdecl);
+
+ if (record)
+ {
+ cgraph_node *oldn = cgraph_node::get (olddecl);
+ cgraph_node *newn = cgraph_node::get (newdecl);
+ cgraph_function_version_info *oldv = oldn->function_version ();
+ cgraph_function_version_info *first_ver, *ver;
+
+ for (ver = oldv; ver; ver = ver->prev)
+ first_ver = ver;
+
+ tree have_base = NULL_TREE;
+ hash_set<tree> linked_variants;
+
+ const char *base_attr_name = "omp declare variant base";
+ size_t base_attr_len = strlen (base_attr_name);
+
+ for (ver = first_ver; ver; ver = ver->next)
+ {
+ tree fn = ver->this_node->decl;
+ tree variant = lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (fn));
+ if (!variant)
+ {
+ /* There should only be one base function (non-variant) in the
+ list. FIXME: Might not be true for versioned functions. */
+ gcc_assert (!have_base);
+ have_base = fn;
+ tree attrs = DECL_ATTRIBUTES (fn);
+ while (attrs)
+ {
+ tree attr = get_attribute_name (attrs);
+ size_t ident_len = IDENTIFIER_LENGTH (attr);
+ if (cmp_attribs (base_attr_name, base_attr_len,
+ IDENTIFIER_POINTER (attr), ident_len))
+ {
+ /* This is a TREE_LIST node used as a tuple of:
+ (variant decl, context selector, location). */
+ tree info = TREE_VALUE (attr);
+ tree linked_variant = TREE_PURPOSE (info);
+ /* Linked variant should appear only once. */
+ gcc_assert (!linked_variants.contains (linked_variant));
+ linked_variants.add (linked_variant);
+ }
+ attrs = TREE_CHAIN (attrs);
+ }
+ }
+ }
+
+ /* We haven't seen the base function yet; do nothing. */
+ if (!have_base)
+ return;
+
+ for (ver = first_ver; ver; ver = ver->next)
+ {
+ tree fn = ver->this_node->decl;
+ if (fn == have_base)
+ continue;
+ if (linked_variants.contains (fn))
+ continue;
+ /* The base function doesn't have a link to this variant yet. Add
+ one now. */
+ tree attrib = lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (fn));
+ gcc_assert (attrib);
+ tree ctx_selector = TREE_VALUE (attrib);
+ /* This location stuff is nonsense. FIXME. */
+ location_t fn_loc = DECL_SOURCE_LOCATION (fn);
+ tree wrapped_loc = maybe_wrap_with_location (integer_zero_node,
+ fn_loc);
+ cp_id_kind idk = CP_ID_KIND_NONE;
+ wrapped_loc
+ = tree_cons (wrapped_loc, build_int_cst (integer_type_node, idk),
+ build_tree_list (wrapped_loc, integer_zero_node));
+ tree info = tree_cons (fn, ctx_selector, wrapped_loc);
+ DECL_ATTRIBUTES (have_base)
+ = tree_cons (get_identifier (base_attr_name), info,
+ DECL_ATTRIBUTES (have_base));
+ if (processing_template_decl)
+ ATTR_IS_DEPENDENT (DECL_ATTRIBUTES (have_base)) = 1;
+ }
+ }
+}
+
/* If NEWDECL is `static' and an `extern' was seen previously,
warn about it. OLDDECL is the previous declaration.
@@ -1547,10 +1683,14 @@ duplicate_function_template_decls (tree newdecl, tree olddecl)
tree oldres = DECL_TEMPLATE_RESULT (olddecl);
/* Function template declarations can be differentiated by parameter
and return type. */
- if (compparms (TYPE_ARG_TYPES (TREE_TYPE (oldres)),
- TYPE_ARG_TYPES (TREE_TYPE (newres)))
- && same_type_p (TREE_TYPE (TREE_TYPE (newdecl)),
- TREE_TYPE (TREE_TYPE (olddecl))))
+ if ((!lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (newres))
+ && !lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (oldres)))
+ && compparms (TYPE_ARG_TYPES (TREE_TYPE (oldres)),
+ TYPE_ARG_TYPES (TREE_TYPE (newres)))
+ && same_type_p (TREE_TYPE (TREE_TYPE (newdecl)),
+ TREE_TYPE (TREE_TYPE (olddecl))))
{
/* ... and also by their template-heads and requires-clauses. */
if (template_heads_equivalent_p (newdecl, olddecl)
@@ -1970,6 +2110,8 @@ duplicate_decls (tree newdecl, tree olddecl, bool hiding, bool was_hidden)
are not ambiguous. */
else if ((!DECL_FUNCTION_VERSIONED (newdecl)
&& !DECL_FUNCTION_VERSIONED (olddecl))
+ && (!DECL_FUNCTION_OMP_VARIANT (newdecl)
+ && !DECL_FUNCTION_OMP_VARIANT (olddecl))
/* Let constrained hidden friends coexist for now, we'll
check satisfaction later. */
&& !member_like_constrained_friend_p (newdecl)
@@ -3143,6 +3285,15 @@ duplicate_decls (tree newdecl, tree olddecl, bool hiding, bool was_hidden)
cgraph_node::delete_function_version_by_decl (newdecl);
}
+ if (TREE_CODE (newdecl) == FUNCTION_DECL
+ && DECL_FUNCTION_OMP_VARIANT (olddecl))
+ {
+ /* Set the flag for newdecl so that it gets copied to olddecl. */
+ DECL_FUNCTION_OMP_VARIANT (newdecl) = 1;
+ /* FIXME: Do we need this here? */
+ //cgraph_node::delete_function_version_by_decl (newdecl);
+ }
+
if (TREE_CODE (newdecl) == FUNCTION_DECL)
{
int function_size;
@@ -8168,6 +8319,20 @@ omp_declare_variant_finalize_one (tree decl, tree attr)
if (idk == CP_ID_KIND_QUALIFIED)
variant = finish_call_expr (variant, &args, /*disallow_virtual=*/true,
koenig_p, tf_warning_or_error);
+ else if (idk == CP_ID_KIND_NONE
+ && DECL_NONSTATIC_MEMBER_FUNCTION_P (variant)
+ && CLASS_TYPE_P (DECL_CONTEXT (decl)))
+ {
+ tree saved_ccp = current_class_ptr;
+ tree saved_ccr = current_class_ref;
+ current_class_ptr = NULL_TREE;
+ current_class_ref = NULL_TREE;
+ inject_this_parameter (DECL_CONTEXT (decl), TYPE_UNQUALIFIED);
+ variant = finish_call_expr (variant, &args, /*disallow_virtual=*/false,
+ koenig_p, tf_warning_or_error);
+ current_class_ptr = saved_ccp;
+ current_class_ref = saved_ccr;
+ }
else
variant = finish_call_expr (variant, &args, /*disallow_virtual=*/false,
koenig_p, tf_warning_or_error);
@@ -8194,6 +8359,9 @@ omp_declare_variant_finalize_one (tree decl, tree attr)
error_at (varid_loc, "variant %qD is a built-in", variant);
return true;
}
+ else if (lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (variant)))
+ TREE_PURPOSE (TREE_VALUE (attr)) = variant;
else
{
tree construct
@@ -1774,6 +1774,17 @@ cplus_decl_attributes (tree *decl, tree attributes, int flags)
}
}
+ if (vec_safe_length (scope_chain->omp_declare_variant_attribute)
+ && TREE_CODE (*decl) == FUNCTION_DECL)
+ {
+ int length = scope_chain->omp_declare_variant_attribute->length ();
+ cp_omp_declare_variant_attr a
+ = (*scope_chain->omp_declare_variant_attribute)[length - 1];
+ tree ctx = copy_list (a.selector);
+ attributes = tree_cons (get_identifier ("omp declare variant overload"),
+ ctx, attributes);
+ }
+
tree late_attrs = NULL_TREE;
if (processing_template_decl)
{
@@ -33,6 +33,7 @@ along with GCC; see the file COPYING3. If not see
#include "gcc-rich-location.h"
#include "cp-name-hint.h"
#include "langhooks.h"
+#include "omp-general.h"
static int interface_strcmp (const char *);
static void init_cp_pragma (void);
@@ -328,6 +329,8 @@ cxx_init (void)
init_cp_semantics ();
init_operators ();
init_method ();
+ if (flag_openmp)
+ omp_init_mangle ();
current_function_decl = NULL;
@@ -55,6 +55,7 @@ along with GCC; see the file COPYING3. If not see
#include "stor-layout.h"
#include "flags.h"
#include "attribs.h"
+#include "omp-general.h"
/* Debugging support. */
@@ -828,6 +829,21 @@ write_mangled_name (const tree decl, bool top_level)
else
gcc_unreachable ();
}
+
+ if (flag_openmp
+ && TREE_CODE (decl) == FUNCTION_DECL
+ && DECL_FUNCTION_OMP_VARIANT (decl))
+ {
+ tree ctx;
+ if ((ctx = lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (decl))))
+ {
+ write_string (JOIN_STR "ompvariant");
+ const char *append
+ = omp_mangle_context_selector (JOIN_STR, TREE_VALUE (ctx));
+ write_string (append);
+ }
+ }
}
/* Returns true if the return type of DECL is part of its signature, and
@@ -48440,6 +48440,62 @@ cp_parser_omp_begin (cp_parser *parser, cp_token *pragma_tok)
= { in_omp_attribute_pragma, device_type, indirect };
vec_safe_push (scope_chain->omp_declare_target_attribute, a);
}
+ else if (strcmp (p, "variant") == 0)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ const char *clause = "";
+ matching_parens parens;
+ location_t match_loc = cp_lexer_peek_token (parser->lexer)->location;
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+ clause = IDENTIFIER_POINTER (id);
+ }
+ if (strcmp (clause, "match") != 0)
+ {
+ cp_parser_error (parser, "expected %<match%>");
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return;
+ }
+
+ cp_lexer_consume_token (parser->lexer);
+
+ if (!parens.require_open (parser))
+ {
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return;
+ }
+
+ tree ctx = cp_parser_omp_context_selector_specification (parser,
+ true);
+ if (ctx == error_mark_node)
+ {
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return;
+ }
+
+ if (vec_safe_length (scope_chain->omp_declare_variant_attribute) > 0)
+ {
+ int length
+ = scope_chain->omp_declare_variant_attribute->length ();
+ cp_omp_declare_variant_attr a
+ = (*scope_chain->omp_declare_variant_attribute)[length - 1];
+ tree outer_ctx = a.selector;
+ ctx = omp_merge_context_selectors (outer_ctx, ctx);
+ }
+
+ ctx = omp_check_context_selector (match_loc, ctx);
+
+ if (ctx != error_mark_node)
+ {
+ cp_omp_declare_variant_attr a
+ = { parser->lexer->in_omp_attribute_pragma, ctx };
+ vec_safe_push (scope_chain->omp_declare_variant_attribute, a);
+ }
+
+ parens.require_close (parser);
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ }
else
{
cp_parser_error (parser, "expected %<target%>");
@@ -48486,41 +48542,70 @@ cp_parser_omp_end (cp_parser *parser, cp_token *pragma_tok)
p = IDENTIFIER_POINTER (id);
}
if (strcmp (p, "target") == 0)
- cp_lexer_consume_token (parser->lexer);
+ {
+ cp_lexer_consume_token (parser->lexer);
+ cp_parser_require_pragma_eol (parser, pragma_tok);
+ if (!vec_safe_length (scope_chain->omp_declare_target_attribute))
+ error_at (pragma_tok->location,
+ "%<#pragma omp end declare target%> without "
+ "corresponding %<#pragma omp declare target%> or "
+ "%<#pragma omp begin declare target%>");
+ else
+ {
+ cp_omp_declare_target_attr
+ a = scope_chain->omp_declare_target_attribute->pop ();
+ if (a.attr_syntax != in_omp_attribute_pragma)
+ {
+ if (a.attr_syntax)
+ error_at (pragma_tok->location,
+ "%qs in attribute syntax terminated "
+ "with %qs in pragma syntax",
+ a.device_type >= 0 ? "begin declare target"
+ : "declare target",
+ "end declare target");
+ else
+ error_at (pragma_tok->location,
+ "%qs in pragma syntax terminated "
+ "with %qs in attribute syntax",
+ a.device_type >= 0 ? "begin declare target"
+ : "declare target",
+ "end declare target");
+ }
+ }
+ }
+ else if (strcmp (p, "variant") == 0)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ cp_parser_require_pragma_eol (parser, pragma_tok);
+ if (!vec_safe_length (scope_chain->omp_declare_variant_attribute))
+ error_at (pragma_tok->location,
+ "%<#pragma omp end declare variant%> without "
+ "corresponding %<#pragma omp begin declare variant%>");
+ else
+ {
+ cp_omp_declare_variant_attr
+ a = scope_chain->omp_declare_variant_attribute->pop ();
+ if (a.attr_syntax != in_omp_attribute_pragma)
+ {
+ if (a.attr_syntax)
+ error_at (pragma_tok->location,
+ "%<begin declare variant%> in attribute syntax "
+ "terminated with %<end declare variant%> in "
+ "pragma syntax");
+ else
+ error_at (pragma_tok->location,
+ "%<begin declare variant%> in pragma syntax "
+ "terminated with %<end declare variant%> in "
+ "attribute syntax");
+ }
+ }
+ }
else
{
cp_parser_error (parser, "expected %<target%>");
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
return;
}
- cp_parser_require_pragma_eol (parser, pragma_tok);
- if (!vec_safe_length (scope_chain->omp_declare_target_attribute))
- error_at (pragma_tok->location,
- "%<#pragma omp end declare target%> without corresponding "
- "%<#pragma omp declare target%> or "
- "%<#pragma omp begin declare target%>");
- else
- {
- cp_omp_declare_target_attr
- a = scope_chain->omp_declare_target_attribute->pop ();
- if (a.attr_syntax != in_omp_attribute_pragma)
- {
- if (a.attr_syntax)
- error_at (pragma_tok->location,
- "%qs in attribute syntax terminated "
- "with %qs in pragma syntax",
- a.device_type >= 0 ? "begin declare target"
- : "declare target",
- "end declare target");
- else
- error_at (pragma_tok->location,
- "%qs in pragma syntax terminated "
- "with %qs in attribute syntax",
- a.device_type >= 0 ? "begin declare target"
- : "declare target",
- "end declare target");
- }
- }
}
else if (strcmp (p, "assumes") == 0)
{
@@ -1669,6 +1669,18 @@ spec_hasher::equal (spec_entry *e1, spec_entry *e2)
tree c2 = e2->spec ? get_constraints (e2->spec) : NULL_TREE;
equal = equivalent_constraints (c1, c2);
}
+ if (equal && flag_openmp)
+ {
+ tree r1 = DECL_TEMPLATE_RESULT (e1->tmpl);
+ tree r2 = DECL_TEMPLATE_RESULT (e2->tmpl);
+ if (TREE_CODE (r1) == FUNCTION_DECL
+ && TREE_CODE (r2) == FUNCTION_DECL
+ && (lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (r1))
+ || lookup_attribute ("omp declare variant overload",
+ DECL_ATTRIBUTES (r2))))
+ equal = false;
+ }
--processing_template_decl;
--comparing_dependent_aliases;
--comparing_specializations;
@@ -11775,6 +11787,10 @@ tsubst_contract_attributes (tree decl, tree args, tsubst_flags_t complain, tree
DECL_ATTRIBUTES (decl) = list;
}
+static tree
+tsubst_function_decl (tree t, tree args, tsubst_flags_t complain,
+ tree lambda_fntype, bool use_spec_table = true);
+
/* Instantiate a single dependent attribute T (a TREE_LIST), and return either
T or a new TREE_LIST, possibly a chain in the case of a pack expansion. */
@@ -11812,9 +11828,16 @@ tsubst_attribute (tree t, tree *decl_p, tree args,
&& is_attribute_p ("omp declare variant base",
get_attribute_name (t)))
{
- ++cp_unevaluated_operand;
- tree varid = tsubst_expr (TREE_PURPOSE (val), args, complain, in_decl);
- --cp_unevaluated_operand;
+ tree varid;
+ if (TREE_CODE (TREE_PURPOSE (val)) == FUNCTION_DECL)
+ varid = tsubst_function_decl (TREE_PURPOSE (val), args, complain,
+ NULL_TREE);
+ else
+ {
+ ++cp_unevaluated_operand;
+ varid = tsubst_expr (TREE_PURPOSE (val), args, complain, in_decl);
+ --cp_unevaluated_operand;
+ }
tree chain = TREE_CHAIN (val);
location_t match_loc = cp_expr_loc_or_input_loc (TREE_PURPOSE (chain));
tree ctx = copy_list (TREE_VALUE (val));
@@ -11908,6 +11931,10 @@ tsubst_attribute (tree t, tree *decl_p, tree args,
}
val = tree_cons (varid, ctx, chain);
}
+ else if (flag_openmp
+ && is_attribute_p ("omp declare variant overload",
+ get_attribute_name (t)))
+ ;
/* If the first attribute argument is an identifier, don't
pass it through tsubst. Attributes like mode, format,
cleanup and several target specific attributes expect it
@@ -14329,7 +14356,7 @@ maybe_rebuild_function_decl_type (tree decl)
static tree
tsubst_function_decl (tree t, tree args, tsubst_flags_t complain,
- tree lambda_fntype, bool use_spec_table = true)
+ tree lambda_fntype, bool use_spec_table /*= true*/)
{
tree gen_tmpl = NULL_TREE, argvec = NULL_TREE;
hashval_t hash = 0;
@@ -14630,7 +14657,11 @@ tsubst_function_decl (tree t, tree args, tsubst_flags_t complain,
if (flag_openmp)
if (tree attr = lookup_attribute ("omp declare variant base",
DECL_ATTRIBUTES (r)))
- omp_declare_variant_finalize (r, attr);
+ {
+ DECL_ATTRIBUTES (r) = tsubst_attributes (DECL_ATTRIBUTES (r), args,
+ complain, r);
+ omp_declare_variant_finalize (r, attr);
+ }
return r;
}
@@ -3467,6 +3467,13 @@ finish_translation_unit (void)
"#pragma omp end declare target");
vec_safe_truncate (scope_chain->omp_declare_target_attribute, 0);
}
+ if (vec_safe_length (scope_chain->omp_declare_variant_attribute))
+ {
+ if (!errorcount)
+ error ("%<omp begin declare variant%> without corresponding "
+ "%<omp end declare variant%>");
+ vec_safe_truncate (scope_chain->omp_declare_variant_attribute, 0);
+ }
if (vec_safe_length (scope_chain->omp_begin_assumes))
{
if (!errorcount)
@@ -10099,6 +10099,8 @@ expand_omp_target (struct omp_region *region)
= DECL_FUNCTION_SPECIFIC_TARGET (current_function_decl);
DECL_FUNCTION_VERSIONED (child_fn2)
= DECL_FUNCTION_VERSIONED (current_function_decl);
+ DECL_FUNCTION_OMP_VARIANT (child_fn2)
+ = DECL_FUNCTION_OMP_VARIANT (current_function_decl);
fn2_node = cgraph_node::get_create (child_fn2);
fn2_node->offloadable = 1;
@@ -45,6 +45,7 @@ along with GCC; see the file COPYING3. If not see
#include "data-streamer.h"
#include "streamer-hooks.h"
#include "opts.h"
+#include "tree-hash-traits.h"
enum omp_requires omp_requires_mask;
@@ -1125,6 +1126,8 @@ const char *omp_tss_map[] =
/* Arrays of property candidates must be null-terminated. */
static const char *const kind_properties[] =
{ "host", "nohost", "cpu", "gpu", "fpga", "any", NULL };
+static const char *const kind_abbrevs[] =
+ { "h", "n", "c", "g", "f", "" };
static const char *const vendor_properties[] =
{ "amd", "arm", "bsc", "cray", "fujitsu", "gnu", "hpe", "ibm", "intel",
"llvm", "nvidia", "pgi", "ti", "unknown", NULL };
@@ -1132,111 +1135,112 @@ static const char *const extension_properties[] =
{ NULL };
static const char *const atomic_default_mem_order_properties[] =
{ "seq_cst", "relaxed", "acq_rel", "acquire", "release", NULL };
+static const char *const atomic_default_mem_order_abbrevs[] =
+ { "sc", "rx", "ar", "ac", "re" };
struct omp_ts_info omp_ts_map[] =
{
- { "kind",
+ { "kind", "k",
(1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
OMP_TRAIT_PROPERTY_NAME_LIST, false,
- kind_properties
+ kind_properties, kind_abbrevs
},
- { "isa",
+ { "isa", "i",
(1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
OMP_TRAIT_PROPERTY_NAME_LIST, false,
- NULL
+ NULL, NULL
},
- { "arch",
+ { "arch", "a",
(1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
OMP_TRAIT_PROPERTY_NAME_LIST, false,
- NULL
+ NULL, NULL
},
- { "device_num",
+ { "device_num", NULL,
(1 << OMP_TRAIT_SET_TARGET_DEVICE),
OMP_TRAIT_PROPERTY_EXPR, false,
- NULL
+ NULL, NULL
},
- { "vendor",
+ { "vendor", "v",
(1 << OMP_TRAIT_SET_IMPLEMENTATION),
OMP_TRAIT_PROPERTY_NAME_LIST, true,
- vendor_properties,
+ vendor_properties, NULL
},
- { "extension",
+ { "extension", "e",
(1 << OMP_TRAIT_SET_IMPLEMENTATION),
OMP_TRAIT_PROPERTY_NAME_LIST, true,
- extension_properties,
+ extension_properties, NULL
},
- { "atomic_default_mem_order",
+ { "atomic_default_mem_order", "a",
(1 << OMP_TRAIT_SET_IMPLEMENTATION),
OMP_TRAIT_PROPERTY_ID, true,
atomic_default_mem_order_properties,
+ atomic_default_mem_order_abbrevs
},
- { "requires",
+ { "requires", "r",
(1 << OMP_TRAIT_SET_IMPLEMENTATION),
OMP_TRAIT_PROPERTY_CLAUSE_LIST, true,
- NULL
+ NULL, NULL
},
- { "unified_address",
+ { "unified_address", "una",
(1 << OMP_TRAIT_SET_IMPLEMENTATION),
OMP_TRAIT_PROPERTY_NONE, true,
- NULL
+ NULL, NULL
},
- { "unified_shared_memory",
+ { "unified_shared_memory", "usm",
(1 << OMP_TRAIT_SET_IMPLEMENTATION),
OMP_TRAIT_PROPERTY_NONE, true,
- NULL
+ NULL, NULL
},
- { "dynamic_allocators",
+ { "dynamic_allocators", "dna",
(1 << OMP_TRAIT_SET_IMPLEMENTATION),
OMP_TRAIT_PROPERTY_NONE, true,
- NULL
+ NULL, NULL
},
- { "reverse_offload",
+ { "reverse_offload", "rvo",
(1 << OMP_TRAIT_SET_IMPLEMENTATION),
OMP_TRAIT_PROPERTY_NONE, true,
- NULL
+ NULL, NULL
},
- { "condition",
+ { "condition", NULL,
(1 << OMP_TRAIT_SET_USER),
OMP_TRAIT_PROPERTY_EXPR, true,
- NULL
+ NULL, NULL
},
- { "target",
+ { "target", "ta",
(1 << OMP_TRAIT_SET_CONSTRUCT),
OMP_TRAIT_PROPERTY_NONE, false,
- NULL
+ NULL, NULL
},
- { "teams",
+ { "teams", "te",
(1 << OMP_TRAIT_SET_CONSTRUCT),
OMP_TRAIT_PROPERTY_NONE, false,
- NULL
+ NULL, NULL
},
- { "parallel",
+ { "parallel", "pa",
(1 << OMP_TRAIT_SET_CONSTRUCT),
OMP_TRAIT_PROPERTY_NONE, false,
- NULL
+ NULL, NULL
},
- { "for",
+ { "for", "fo",
(1 << OMP_TRAIT_SET_CONSTRUCT),
OMP_TRAIT_PROPERTY_NONE, false,
- NULL
+ NULL, NULL
},
- { "simd",
+ { "simd", "si",
(1 << OMP_TRAIT_SET_CONSTRUCT),
OMP_TRAIT_PROPERTY_CLAUSE_LIST, false,
- NULL
+ NULL, NULL
},
- { NULL, 0, OMP_TRAIT_PROPERTY_NONE, false, NULL } /* OMP_TRAIT_LAST */
+ /* OMP_TRAIT_LAST */
+ { NULL, NULL, 0, OMP_TRAIT_PROPERTY_NONE, false, NULL, NULL }
};
+/* Helper function for omp_context_name_list_prop. Return string from either
+ an identifier or a string constant. */
-/* Return a name from PROP, a property in selectors accepting
- name lists. */
-
-const char *
-omp_context_name_list_prop (tree prop)
+static const char *
+omp_context_name_list_prop_1 (tree val)
{
- gcc_assert (OMP_TP_NAME (prop) == OMP_TP_NAMELIST_NODE);
- tree val = OMP_TP_VALUE (prop);
switch (TREE_CODE (val))
{
case IDENTIFIER_NODE:
@@ -1254,6 +1258,16 @@ omp_context_name_list_prop (tree prop)
}
}
+/* Return a name from PROP, a property in selectors accepting
+ name lists. */
+
+const char *
+omp_context_name_list_prop (tree prop)
+{
+ gcc_assert (OMP_TP_NAME (prop) == OMP_TP_NAMELIST_NODE);
+ return omp_context_name_list_prop_1 (OMP_TP_VALUE (prop));
+}
+
/* Diagnose errors in an OpenMP context selector, return CTX if
it is correct or error_mark_node otherwise. */
@@ -1797,6 +1811,611 @@ omp_context_selector_matches (tree ctx)
return ret;
}
+static void
+omp_copy_trait_set (tree from_ts, tree *to_ts)
+{
+ for (tree ts = from_ts; ts; ts = TREE_CHAIN (ts))
+ *to_ts = make_trait_selector (OMP_TS_CODE (ts), OMP_TS_SCORE (ts),
+ OMP_TS_PROPERTIES (ts), *to_ts);
+}
+
+/* Trait properties are stored in a tree list, and for namelist properties,
+ may use either an IDENTIFIER_POINTER or STRING_CST for their value. We
+ want these to be considered equal if they represent the same string (e.g.
+ quoted or not in the source program). Arrange a hash table to allow us to
+ do this. */
+
+typedef std::pair<tree, tree> omp_trait_prop;
+
+struct omp_trait_prop_hash : typed_noop_remove <omp_trait_prop>
+{
+ typedef omp_trait_prop value_type;
+ typedef omp_trait_prop compare_type;
+
+ static inline hashval_t hash (omp_trait_prop);
+ static const bool empty_zero_p = true;
+ static inline bool is_empty (omp_trait_prop);
+ static inline bool is_deleted (omp_trait_prop);
+ static inline bool equal (const omp_trait_prop &, const omp_trait_prop &);
+ static inline void mark_empty (omp_trait_prop &);
+};
+
+inline hashval_t
+omp_trait_prop_hash::hash (omp_trait_prop t)
+{
+ if (t.first == OMP_TP_NAMELIST_NODE)
+ return htab_hash_string (omp_context_name_list_prop_1 (t.second));
+ if (t.first && t.second)
+ return iterative_hash_expr (t.second, iterative_hash_expr (t.first, 0));
+ else if (t.first)
+ return iterative_hash_expr (t.first, 0);
+ else if (t.second)
+ return iterative_hash_expr (t.second, 0);
+ else
+ gcc_unreachable ();
+}
+
+inline bool
+omp_trait_prop_hash::is_empty (omp_trait_prop t)
+{
+ return !t.first && !t.second;
+}
+
+inline bool
+omp_trait_prop_hash::is_deleted (omp_trait_prop)
+{
+ return false;
+}
+
+inline bool
+omp_trait_prop_hash::equal (const omp_trait_prop &a, const omp_trait_prop &b)
+{
+ if (a.first == OMP_TP_NAMELIST_NODE
+ && b.first == OMP_TP_NAMELIST_NODE)
+ {
+ const char *a_name = omp_context_name_list_prop_1 (a.second);
+ const char *b_name = omp_context_name_list_prop_1 (b.second);
+ return strcmp (a_name, b_name) == 0;
+ }
+
+ if (a.first && a.second && b.first && b.second)
+ return operand_equal_p (a.first, b.first)
+ && operand_equal_p (a.second, b.second);
+ else if (a.first && b.first)
+ return !a.second && !b.second && a.first == b.first;
+ else if (a.second && b.second)
+ return !a.first && !b.first && operand_equal_p (a.second, b.second);
+ else
+ return false;
+}
+
+inline void
+omp_trait_prop_hash::mark_empty (omp_trait_prop &t)
+{
+ t.first = NULL_TREE;
+ t.second = NULL_TREE;
+}
+
+typedef hash_set<omp_trait_prop_hash> omp_trait_prop_set;
+
+static void
+omp_gather_trait_sets (omp_trait_prop_set sets[], tree scores[], tree inlist)
+{
+ for (tree ts = inlist; ts; ts = TREE_CHAIN (ts))
+ {
+ enum omp_ts_code ts_code = OMP_TS_CODE (ts);
+
+ if (scores && omp_ts_map[ts_code].allow_score)
+ scores[ts_code] = OMP_TS_SCORE (ts);
+
+ for (tree tp = OMP_TS_PROPERTIES (ts); tp; tp = TREE_CHAIN (tp))
+ {
+ tree name = OMP_TP_NAME (tp);
+ tree value = OMP_TP_VALUE (tp);
+ sets[ts_code].add (std::make_pair (name, value));
+ }
+ }
+}
+
+/* We might have things like:
+
+ outer: arch(avx2,sse4.1,3dnow)
+ inner: arch(sse4.1)
+
+ outer: kind(any) kind(gpu,cpu) kind(gpu,cpu)
+ inner: kind(gpu) kind(gpu) kind(any)
+
+ We want to add the intersection of the inner and outer sets -- most of the
+ time that just means adding the inner set, but there are some special cases
+ that need handling (e.g. "any" on inner, but some more restrictive kind on
+ outer). */
+
+static void
+omp_combine_trait_properties (omp_trait_prop_set &outer_props,
+ omp_trait_prop_set &inner_props,
+ omp_tss_code tss_code, omp_ts_code ts_code,
+ tree *to_list, tree score)
+{
+ tree anyt = get_identifier ("any");
+ tree combined_traits = NULL_TREE;
+
+ std::pair<tree, tree> any = std::make_pair (OMP_TP_NAMELIST_NODE, anyt);
+
+ if (outer_props.is_empty ())
+ for (std::pair<tree, tree> e : inner_props)
+ combined_traits = make_trait_property (e.first, e.second,
+ combined_traits);
+ else if (inner_props.is_empty ()
+ || (tss_code == OMP_TRAIT_SET_DEVICE
+ && ts_code == OMP_TRAIT_DEVICE_KIND
+ && inner_props.contains (any)))
+ for (std::pair<tree, tree> e : outer_props)
+ combined_traits = make_trait_property (e.first, e.second,
+ combined_traits);
+ else
+ for (std::pair<tree, tree> e : inner_props)
+ if (tss_code == OMP_TRAIT_SET_IMPLEMENTATION
+ && ts_code == OMP_TRAIT_IMPLEMENTATION_ADMO
+ && !outer_props.contains (e))
+ error ("cannot combine nested %<atomic_default_mem_order%> properties");
+ else if (outer_props.contains (e)
+ || (tss_code == OMP_TRAIT_SET_DEVICE
+ && ts_code == OMP_TRAIT_DEVICE_KIND
+ && outer_props.contains (any)))
+ combined_traits = make_trait_property (e.first, e.second,
+ combined_traits);
+
+ if (combined_traits)
+ *to_list = make_trait_selector (ts_code, score,
+ nreverse (combined_traits), *to_list);
+}
+
+static void
+omp_combine_trait_sets (tree outer_ts, tree inner_ts, omp_tss_code tss_code,
+ tree *to_list)
+{
+ unsigned HOST_WIDE_INT used_traits = 0;
+
+ for (tree t = outer_ts; t; t = TREE_CHAIN (t))
+ used_traits |= 1 << OMP_TS_CODE (t);
+ for (tree t = inner_ts; t; t = TREE_CHAIN (t))
+ used_traits |= 1 << OMP_TS_CODE (t);
+
+ omp_trait_prop_set outer_sets[OMP_TRAIT_LAST];
+ omp_trait_prop_set inner_sets[OMP_TRAIT_LAST];
+
+ tree outer_scores[OMP_TRAIT_LAST] = { };
+ tree inner_scores[OMP_TRAIT_LAST] = { };
+
+ omp_gather_trait_sets (outer_sets, outer_scores, outer_ts);
+ omp_gather_trait_sets (inner_sets, inner_scores, inner_ts);
+
+ for (unsigned i = 0; i < OMP_TRAIT_LAST; i++)
+ if (used_traits & (1 << i))
+ {
+ omp_ts_code ts_code = static_cast<omp_ts_code>(i);
+ /* We can technically have a score on both the inner and outer context
+ selectors (e.g. for the "implementation" trait set selector). Just
+ ignore the outer one. */
+ tree use_score = inner_scores[i] ? inner_scores[i] : outer_scores[i];
+
+ switch (omp_ts_map[ts_code].tp_type)
+ {
+ case OMP_TRAIT_PROPERTY_ID:
+ case OMP_TRAIT_PROPERTY_NAME_LIST:
+ omp_combine_trait_properties (outer_sets[i], inner_sets[i],
+ tss_code, ts_code, to_list,
+ use_score);
+ break;
+
+ case OMP_TRAIT_PROPERTY_EXPR:
+ /* The user trait set selector just contains the 'condition' trait.
+ No attempt is made to combine conditions in outer/inner context
+ selectors here -- instead, the outer condition is just
+ dropped. */
+ if (ts_code == OMP_TRAIT_USER_CONDITION && inner_ts && outer_ts)
+ warning (OPT_Wopenmp, "ignoring condition on enclosing context "
+ "selector");
+
+ if (inner_ts)
+ omp_copy_trait_set (inner_ts, to_list);
+ else if (outer_ts)
+ omp_copy_trait_set (outer_ts, to_list);
+ break;
+
+ case OMP_TRAIT_PROPERTY_NONE:
+ switch (ts_code)
+ {
+ case OMP_TRAIT_IMPLEMENTATION_UNIFIED_ADDRESS:
+ case OMP_TRAIT_IMPLEMENTATION_UNIFIED_SHARED_MEMORY:
+ case OMP_TRAIT_IMPLEMENTATION_DYNAMIC_ALLOCATORS:
+ case OMP_TRAIT_IMPLEMENTATION_REVERSE_OFFLOAD:
+ /* For these ones, we want the union of the features listed on
+ the outer and inner context selectors. We already checked
+ the trait has been mentioned in one or the other selector
+ via 'used_traits', so add it here. */
+ *to_list = make_trait_selector (ts_code, NULL_TREE,
+ NULL_TREE, *to_list);
+ break;
+
+ default:
+ sorry ("don't know how to merge nested context selectors");
+ }
+ break;
+
+ default:
+ gcc_unreachable ();
+ }
+ }
+}
+
+extern tree omp_get_context_selector_list (tree ctx, enum omp_tss_code set);
+
+tree
+omp_merge_context_selectors (tree outer_ctx, tree inner_ctx)
+{
+ tree merged_ctx = NULL_TREE;
+
+ for (unsigned i = OMP_TRAIT_SET_CONSTRUCT; i != OMP_TRAIT_SET_LAST; i++)
+ {
+ omp_tss_code tss_code = static_cast<omp_tss_code>(i);
+ tree outer_ts = omp_get_context_selector_list (outer_ctx, tss_code);
+ tree inner_ts = omp_get_context_selector_list (inner_ctx, tss_code);
+ tree merged_ts = NULL_TREE;
+
+ switch (tss_code)
+ {
+ case OMP_TRAIT_SET_CONSTRUCT:
+ if (outer_ts)
+ omp_copy_trait_set (outer_ts, &merged_ts);
+ if (inner_ts)
+ omp_copy_trait_set (inner_ts, &merged_ts);
+ break;
+
+ case OMP_TRAIT_SET_DEVICE:
+ case OMP_TRAIT_SET_TARGET_DEVICE:
+ case OMP_TRAIT_SET_IMPLEMENTATION:
+ case OMP_TRAIT_SET_USER:
+ if (outer_ts && inner_ts)
+ omp_combine_trait_sets (outer_ts, inner_ts, tss_code, &merged_ts);
+ else if (outer_ts)
+ omp_copy_trait_set (outer_ts, &merged_ts);
+ else if (inner_ts)
+ omp_copy_trait_set (inner_ts, &merged_ts);
+ break;
+
+ default:
+ gcc_unreachable ();
+ }
+
+ if (merged_ts)
+ merged_ctx = make_trait_set_selector (tss_code, nreverse (merged_ts),
+ merged_ctx);
+ }
+
+ return nreverse (merged_ctx);
+}
+
+/* OpenMP "declare variant" mangling obstack bits. These are essentially
+ copied from cp/mangle.cc, but we'll need them for C also, so they need to go
+ here. */
+
+static struct obstack *omp_mangle_obstack;
+static struct obstack omp_name_obstack;
+static void *omp_name_base;
+
+/* Append a single character to the end of the mangled
+ representation. */
+#define write_char(CHAR) \
+ obstack_1grow (omp_mangle_obstack, (CHAR))
+
+/* Append a NUL-terminated string to the end of the mangled
+ representation. */
+#define write_string(STRING) \
+ obstack_grow (omp_mangle_obstack, (STRING), strlen (STRING))
+
+static void
+omp_start_mangling (void)
+{
+ obstack_free (&omp_name_obstack, omp_name_base);
+ omp_mangle_obstack = &omp_name_obstack;
+ omp_name_base = obstack_alloc (&omp_name_obstack, 0);
+}
+
+static const char *
+omp_finish_mangling (void)
+{
+ write_char ('\0');
+ return (const char *) obstack_finish (omp_mangle_obstack);
+}
+
+void
+omp_init_mangle (void)
+{
+ gcc_obstack_init (&omp_name_obstack);
+ omp_name_base = obstack_alloc (&omp_name_obstack, 0);
+}
+
+static int
+omp_string_compare (const void *a, const void *b)
+{
+ const char *as = *(char * const *) a;
+ const char *bs = *(char * const *) b;
+ return strcmp (as, bs);
+}
+
+static void
+omp_stringify_sorted_property_set (const char *sep, omp_ts_code ts_code,
+ omp_trait_prop_set &props)
+{
+ vec<const char *> names = vNULL;
+ hash_map<nofree_string_hash, const char *> abbrevs;
+ omp_tp_type tp_type = omp_ts_map[ts_code].tp_type;
+
+ /* This function only really works for name lists. */
+ gcc_assert (tp_type == OMP_TRAIT_PROPERTY_NAME_LIST
+ || tp_type == OMP_TRAIT_PROPERTY_ID);
+
+ if (omp_ts_map[ts_code].prop_abbrevs)
+ for (int i = 0; omp_ts_map[ts_code].valid_properties[i]; i++)
+ abbrevs.put (omp_ts_map[ts_code].valid_properties[i],
+ omp_ts_map[ts_code].prop_abbrevs[i]);
+
+ for (auto prop : props)
+ {
+ if (prop.first == OMP_TP_NAMELIST_NODE)
+ {
+ const char *name = omp_context_name_list_prop_1 (prop.second);
+ if (omp_ts_map[ts_code].prop_abbrevs)
+ {
+ const char **repl = abbrevs.get (name);
+ if (repl)
+ names.safe_push (*repl);
+ else
+ error ("missing abbreviation for %qs", name);
+ }
+ else
+ names.safe_push (name);
+ }
+ else if (tp_type == OMP_TRAIT_PROPERTY_ID
+ && TREE_CODE (prop.first) == IDENTIFIER_NODE)
+ {
+ const char *name = IDENTIFIER_POINTER (prop.first);
+ if (omp_ts_map[ts_code].prop_abbrevs)
+ {
+ const char **repl = abbrevs.get (name);
+ if (repl)
+ names.safe_push (*repl);
+ else
+ error ("missing abbreviation for %qs", name);
+ }
+ else
+ names.safe_push (name);
+ }
+ else
+ gcc_unreachable ();
+ }
+
+ names.qsort (omp_string_compare);
+
+ const char *use_sep = "";
+
+ for (auto &n : names)
+ {
+ bool needs_escaping = false;
+ for (int i = 0; n[i]; i++)
+ if (!ISALPHA (n[i]) && !ISDIGIT (n[i]) && n[i] != '_')
+ needs_escaping = true;
+ if (needs_escaping)
+ {
+ write_string (use_sep);
+ for (int i = 0; n[i]; i++)
+ {
+ /* We may have an arbirary string constants representing e.g. an
+ ISA, which could contain characters not valid for symbols in
+ assembly output. The below isn't really safe for arbitrary
+ character encodings (either in source or assembly output), but
+ should be a good-enough approximation for our purposes. */
+ if (ISALPHA (n[i]) || ISDIGIT (n[i]) || n[i] == '_')
+ write_char (n[i]);
+ else
+ {
+ char tmpbuf[4];
+ sprintf (tmpbuf, "%.2x", (unsigned char) n[i]);
+ write_string (sep);
+ write_string (tmpbuf);
+ }
+ }
+ }
+ else
+ {
+ write_string (use_sep);
+ write_string (n);
+ }
+ /* If we're using abbreviations for these trait properties, we don't need
+ extra separators between them. */
+ if (!omp_ts_map[ts_code].prop_abbrevs)
+ use_sep = sep;
+ }
+}
+
+/* Key:
+
+ C: start construct
+ ta: target
+ te: teams
+ pa: parallel
+ fo: for/do
+ si: simd
+ D: start device
+ i: isa
+ (ordered) [<string> SEP] list
+ a: arch
+ (ordered) [<string> SEP] list
+ k: kind
+ h: host
+ n: nohost
+ c: cpu
+ g: gpu
+ f: fpga
+ (omitted): any
+ T: start target_device
+ (as above, also)
+ N: (ordered) [<device number> SEP] list
+ I: start implementation
+ v: vendor
+ (ordered) [<string> SEP] list
+ e: extension
+ (nothing here)
+ r: requires
+ (tbd)
+ a: atomic_default_mem_order
+ ar: acq_rel
+ ac: acquire
+ rx: relaxed
+ re: release
+ sc: seq_cst
+ U: start user
+ (no further encoding, just 'U' or not)
+ SEP: separate trait selector sets
+*/
+
+const char *
+omp_mangle_context_selector (const char *sep, tree ctx)
+{
+ gcc_assert (ctx);
+
+ omp_start_mangling ();
+
+ for (int i = OMP_TRAIT_SET_CONSTRUCT; i != OMP_TRAIT_SET_LAST; i++)
+ {
+ omp_tss_code tss_code = static_cast<omp_tss_code>(i);
+ tree ts_list = omp_get_context_selector_list (ctx, tss_code);
+
+ if (!ts_list)
+ continue;
+
+ switch (tss_code)
+ {
+ case OMP_TRAIT_SET_CONSTRUCT:
+ {
+ write_string (sep);
+ write_char ('C');
+ for (tree ts = ts_list; ts; ts = TREE_CHAIN (ts))
+ write_string (omp_ts_map[OMP_TS_CODE (ts)].abbrev_name);
+ }
+ break;
+ case OMP_TRAIT_SET_DEVICE:
+ {
+ omp_trait_prop_set trait_sets[OMP_TRAIT_LAST];
+
+ omp_gather_trait_sets (trait_sets, NULL, ts_list);
+ write_string (sep);
+ write_char ('D');
+
+ const char *use_sep = "";
+
+ for (unsigned j = 0; j < OMP_TRAIT_LAST; j++)
+ {
+ omp_ts_code ts_code = static_cast<omp_ts_code>(j);
+
+ if (!(omp_ts_map[ts_code].tss_mask
+ & (1 << OMP_TRAIT_SET_DEVICE)))
+ continue;
+
+ if (trait_sets[ts_code].is_empty ())
+ continue;
+
+ write_string (use_sep);
+ write_string (omp_ts_map[ts_code].abbrev_name);
+
+ omp_stringify_sorted_property_set (sep, ts_code,
+ trait_sets[ts_code]);
+
+ use_sep = sep;
+ }
+ }
+ break;
+ case OMP_TRAIT_SET_TARGET_DEVICE:
+ {
+ write_string (sep);
+ write_char ('T');
+ /* This should share the 'device' code above, but for now is
+ unimplemented. */
+ gcc_unreachable ();
+ }
+ break;
+ case OMP_TRAIT_SET_IMPLEMENTATION:
+ {
+ omp_trait_prop_set trait_sets[OMP_TRAIT_LAST];
+
+ omp_gather_trait_sets (trait_sets, NULL, ts_list);
+ write_string (sep);
+ write_char ('I');
+ const char *use_sep = "";
+
+ for (unsigned j = 0; j < OMP_TRAIT_LAST; j++)
+ {
+ omp_ts_code ts_code = static_cast<omp_ts_code>(j);
+
+ if (!(omp_ts_map[ts_code].tss_mask
+ & (1 << OMP_TRAIT_SET_IMPLEMENTATION)))
+ continue;
+
+ switch (omp_ts_map[ts_code].tp_type)
+ {
+ case OMP_TRAIT_PROPERTY_NONE:
+ if (omp_get_context_selector (ctx, tss_code, ts_code))
+ {
+ write_string (use_sep);
+ write_string (omp_ts_map[ts_code].abbrev_name);
+ }
+ break;
+
+ case OMP_TRAIT_PROPERTY_ID:
+ case OMP_TRAIT_PROPERTY_NAME_LIST:
+ if (trait_sets[i].is_empty ())
+ continue;
+
+ write_string (use_sep);
+ write_string (omp_ts_map[ts_code].abbrev_name);
+
+ omp_stringify_sorted_property_set (sep, ts_code,
+ trait_sets[ts_code]);
+ break;
+
+ case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
+ if (omp_get_context_selector (ctx, tss_code, ts_code))
+ {
+ write_string (use_sep);
+ write_string ("?????");
+ }
+ break;
+
+ default:
+ gcc_unreachable ();
+ }
+
+ use_sep = sep;
+ }
+ continue;
+ }
+ break;
+ case OMP_TRAIT_SET_USER:
+ {
+ write_string (sep);
+ write_char ('U');
+ continue;
+ }
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ }
+
+ return omp_finish_mangling ();
+}
+
/* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
in omp_context_selector_set_compare. */
@@ -170,6 +170,9 @@ extern tree omp_check_context_selector (location_t loc, tree ctx);
extern void omp_mark_declare_variant (location_t loc, tree variant,
tree construct);
extern int omp_context_selector_matches (tree);
+extern tree omp_merge_context_selectors (tree, tree);
+extern void omp_init_mangle (void);
+extern const char * omp_mangle_context_selector (const char *, tree);
extern int omp_context_selector_set_compare (enum omp_tss_code, tree, tree);
extern tree omp_get_context_selector (tree, enum omp_tss_code,
enum omp_ts_code);
@@ -2089,6 +2089,8 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
= DECL_FUNCTION_SPECIFIC_TARGET (current_function_decl);
DECL_FUNCTION_VERSIONED (decl)
= DECL_FUNCTION_VERSIONED (current_function_decl);
+ DECL_FUNCTION_OMP_VARIANT (decl)
+ = DECL_FUNCTION_OMP_VARIANT (current_function_decl);
if (omp_maybe_offloaded_ctx (ctx))
{
@@ -80,10 +80,12 @@ extern const char *omp_tss_map [];
null-terminated array of strings. */
struct omp_ts_info {
const char *name;
+ const char *abbrev_name;
unsigned int tss_mask;
enum omp_tp_type tp_type;
bool allow_score;
const char * const *valid_properties;
+ const char * const *prop_abbrevs;
};
extern struct omp_ts_info omp_ts_map[];
new file mode 100644
@@ -0,0 +1,82 @@
+#if 1
+#ifdef _OPENMP
+namespace bar {
+#pragma omp begin declare variant match(construct={target})
+#pragma omp begin declare variant match(construct={parallel})
+namespace qux {
+int foo (char *buf, int len)
+{
+ return 5;
+}
+} // namespace qux
+#pragma omp end declare variant
+#pragma omp end declare variant
+} // namespace bar
+#endif
+
+namespace bar {
+namespace qux {
+int foo (char *buf, int len)
+{
+ return 3;
+}
+} // namespace qux
+} // namespace bar
+
+#pragma omp begin declare variant match(implementation={vendor(gnu)})
+#pragma omp begin declare variant match(implementation={vendor(gnu,nvidia)})
+int quux (int c, int d)
+{
+ return c+d;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(device={arch("mips","riscv")})
+#pragma omp begin declare variant match(device={arch("mips"),isa("mips4")})
+int quuux (int c, int d)
+{
+ return c+d;
+}
+#pragma omp end declare variant
+#pragma omp begin declare variant match(device={arch("riscv"),isa("rv3.5")})
+int quuux (int c, int d)
+{
+ return c+d;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+int quux (int c, int d)
+{
+ return c+d;
+}
+#endif
+
+#if 1
+#pragma omp begin declare variant match(device={kind(any)})
+#pragma omp begin declare variant match(device={kind(gpu,nohost,fpga)})
+int warble (int c, int d)
+{
+ return c*d;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+#endif
+
+#if 1
+#pragma omp begin declare variant match(device={kind(gpu,nohost,fpga)})
+#pragma omp begin declare variant match(device={kind(any)})
+int warble2 (int c, int d)
+{
+ return c*d;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+#endif
+
+
+#pragma omp begin declare variant match(device={arch("mips",riscv)})
+#pragma omp begin declare variant match(device={arch(mips,"riscv")})
+#pragma omp end declare variant
+#pragma omp end declare variant
new file mode 100644
@@ -0,0 +1,20 @@
+#pragma omp begin declare variant match(implementation={dynamic_allocators,unified_shared_memory})
+/* This is a "sorry" for now. */
+#pragma omp begin declare variant match(implementation={requires(reverse_offload)})
+int foo (int c)
+{
+ return c+5;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+int foo (int c)
+{
+ return c+3;
+}
+
+int main(void)
+{
+ foo (6);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,32 @@
+#include <stdio.h>
+
+int r_foo (int c, int d)
+{
+ return (c+d) * 3;
+}
+
+#pragma omp declare variant (r_foo) match(implementation={requires(unified_shared_memory)})
+int foo (int c, int d)
+{
+ return c+d;
+}
+
+int main()
+{
+ printf ("host: %d\n", foo (3, 4));
+
+#pragma omp parallel
+ {
+#pragma omp single
+ {
+ printf ("parallel: %d\n", foo (3, 4));
+ }
+ }
+
+#pragma omp target
+ {
+ printf ("target: %d\n", foo (3, 4));
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,26 @@
+#pragma omp begin declare variant match(implementation={vendor(score(5):nvidia)})
+#pragma omp begin declare variant match(implementation={vendor(nvidia)})
+int foo (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(implementation={vendor(nvidia)})
+#pragma omp begin declare variant match(implementation={vendor(score(7):nvidia)})
+int bar (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(implementation={vendor(score(9):nvidia)})
+#pragma omp begin declare variant match(implementation={vendor(score(13):"nvidia")})
+int qux (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+ const bool enable_p = true;
+ const bool enable2_p = false;
+
+#pragma omp begin declare variant match(implementation={vendor(score(9):nvidia,amd)},user={condition(score(17):enable_p)})
+#pragma omp begin declare variant match(implementation={vendor(score(13):"nvidia")},user={condition(enable2_p)})
+int quux (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
new file mode 100644
@@ -0,0 +1,6 @@
+#pragma omp begin declare variant match(implementation={vendor(amd,nvidia),atomic_default_mem_order(relaxed)})
+#pragma omp begin declare variant match(implementation={atomic_default_mem_order(acq_rel)})
+int foo (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
new file mode 100644
@@ -0,0 +1,10 @@
+#pragma omp begin declare variant match(implementation={vendor(amd,nvidia),atomic_default_mem_order(acq_rel)})
+#pragma omp begin declare variant match(implementation={atomic_default_mem_order(acq_rel)})
+int foo (int c) { return c; }
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+
+int foo (int c) { return c; }
+
+int main() { foo (5); return 0; }
new file mode 100644
@@ -0,0 +1,4 @@
+#pragma omp begin declare variant match(device={arch("super power")})
+int foo (int c) { return c; }
+#pragma omp end declare variant
+
@@ -8300,6 +8300,10 @@ dump_function_to_file (tree fndecl, FILE *file, dump_flags_t flags)
print_omp_context_selector (file, TREE_VALUE (a),
dump_flags);
}
+ else if (!strcmp (IDENTIFIER_POINTER (name),
+ "omp declare variant overload"))
+ print_omp_context_selector (file, TREE_VALUE (chain),
+ dump_flags);
else
print_generic_expr (file, TREE_VALUE (chain), dump_flags);
fprintf (file, ")");
@@ -1994,8 +1994,9 @@ struct GTY(()) tree_function_decl {
unsigned has_debug_args_flag : 1;
unsigned versioned_function : 1;
unsigned replaceable_operator : 1;
+ unsigned omp_variant : 1;
- /* 11 bits left for future expansion. */
+ /* 10 bits left for future expansion. */
/* 32 bits on 64-bit HW. */
};
@@ -3505,6 +3505,13 @@ extern vec<tree, va_gc> **decl_debug_args_insert (tree);
#define DECL_FUNCTION_VERSIONED(NODE)\
(FUNCTION_DECL_CHECK (NODE)->function_decl.versioned_function)
+/* In FUNCTION_DECL, a OpenMP 'declare variant' function from a
+ 'begin/end declare variant' block. These have the same name as some base
+ function, so this flag is used to disambiguate them (similar to
+ DECL_FUNCTION_VERSIONED). */
+#define DECL_FUNCTION_OMP_VARIANT(NODE) \
+ (FUNCTION_DECL_CHECK (NODE)->function_decl.omp_variant)
+
/* In FUNCTION_DECL, this is set if this function is a C++ constructor.
Devirtualization machinery uses this knowledge for determing type of the
object constructed. Also we assume that constructor address is not
new file mode 100644
@@ -0,0 +1,19 @@
+#pragma omp begin declare variant match(implementation={dynamic_allocators,unified_shared_memory})
+#pragma omp begin declare variant match(implementation={reverse_offload})
+int foo (int c)
+{
+ return c+5;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+int foo (int c)
+{
+ return c+3;
+}
+
+int main(void)
+{
+ foo (6);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,51 @@
+#include <stdio.h>
+
+int foo(int c)
+{
+ return c+1;
+}
+
+int foo(int c, int d)
+{
+ return c+d;
+}
+
+int foo(int c, int d, int e, int f = 10)
+{
+ return c+d+e+f;
+}
+
+#pragma omp begin declare variant match(construct={target})
+int foo(int c)
+{
+ return (c+1)*2;
+}
+
+int foo(int c, int d)
+{
+ return (c+d)*2;
+}
+
+int foo(int c, int d, int e, int f = 10)
+{
+ return (c+d+e+f)*2;
+}
+#pragma omp end declare variant
+
+int main()
+{
+ printf ("foo(5) = %d\n", foo(5));
+ printf ("foo(5,2) = %d\n", foo(5,2));
+ printf ("foo(5,2,3) = %d\n", foo(5,2,3));
+ printf ("foo(5,2,3,4) = %d\n", foo(5,2,3,4));
+
+#pragma omp target
+ {
+ printf ("foo(5) = %d\n", foo(5));
+ printf ("foo(5,2) = %d\n", foo(5,2));
+ printf ("foo(5,2,3) = %d\n", foo(5,2,3));
+ printf ("foo(5,2,3,4) = %d\n", foo(5,2,3,4));
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,57 @@
+#include <stdio.h>
+
+template<typename T>
+T foo(T c)
+{
+ return c+1;
+}
+
+template<typename T>
+T foo(T c, T d)
+{
+ return c+d;
+}
+
+template<typename T, int U>
+T foo(T c, T d, T e, T f = U)
+{
+ return c+d+e+f;
+}
+
+#pragma omp begin declare variant match(construct={target})
+template<typename T>
+T foo(T c)
+{
+ return (c+1)*2;
+}
+
+template<typename T>
+T foo(T c, T d)
+{
+ return (c+d)*2;
+}
+
+template<typename T, int U>
+T foo(T c, T d, T e, T f = U)
+{
+ return (c+d+e+f)*2;
+}
+#pragma omp end declare variant
+
+int main()
+{
+ printf ("foo(5) = %d\n", foo<int>(5));
+ printf ("foo(5,2) = %d\n", foo<int>(5,2));
+ printf ("foo(5,2,3) = %d\n", foo<int, 5>(5,2,3));
+ printf ("foo(5,2,3,4) = %d\n", foo<int, 10>(5,2,3,4));
+
+#pragma omp target
+ {
+ printf ("foo(5) = %d\n", foo<int>(5));
+ printf ("foo(5,2) = %d\n", foo<int>(5,2));
+ printf ("foo(5,2,3) = %d\n", foo<int, 5>(5,2,3));
+ printf ("foo(5,2,3,4) = %d\n", foo<int, 10>(5,2,3,4));
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,55 @@
+#include <stdio.h>
+
+template<typename T>
+T foo(T c)
+{
+ return c+1;
+}
+
+template<typename T>
+T foo(T c, T d)
+{
+ return c+d;
+}
+
+template<typename T, typename S>
+T foo(T c, S d)
+{
+ return c+d+2;
+}
+
+#pragma omp begin declare variant match(construct={target})
+template<typename T>
+T foo(T c)
+{
+ return (c+1)*2;
+}
+
+template<typename T>
+T foo(T c, T d)
+{
+ return (c+d)*2;
+}
+
+template<typename T, typename S>
+T foo(T c, S d)
+{
+ return (c+d+2)*2;
+}
+#pragma omp end declare variant
+
+int main()
+{
+ printf ("foo<int>(5) = %d\n", foo<int>(5));
+ printf ("foo<int>(5,2) = %d\n", foo<int>(5,2));
+ printf ("foo<int,long>(5,2) = %d\n", foo<int, long>(5,2));
+
+#pragma omp target
+ {
+ printf ("foo<int>(5) = %d\n", foo<int>(5));
+ printf ("foo<int>(5,2) = %d\n", foo<int>(5,2));
+ printf ("foo<int,long>(5,2) = %d\n", foo<int, long>(5,2));
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,30 @@
+#include <cstdio>
+
+class foo {
+public:
+ int bar (int c, int d) { return c + d; }
+ static int qux (int c, int d) { return c + d + 5; }
+
+#if 1
+#pragma omp begin declare variant match(construct={target})
+ int bar (int c, int d) { return c + d + 3; }
+ static int qux (int c, int d) { return c + d + 7; }
+#pragma omp end declare variant
+#endif
+};
+
+int main()
+{
+ foo fv;
+
+ printf ("fv.bar (1, 2)=%d\n", fv.bar (1, 2));
+ printf ("foo::qux (1, 2)=%d\n", foo::qux (1, 2));
+
+#pragma omp target
+ {
+ printf ("fv.bar (1, 2)=%d\n", fv.bar (1, 2));
+ printf ("foo::qux (1, 2)=%d\n", foo::qux (1, 2));
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,29 @@
+#include <cstdio>
+
+class foo {
+public:
+ int t_bar (int c, int d) { return c + d; }
+ static int t_qux (int c, int d) { return c + d + 5; }
+
+#pragma omp declare variant (t_bar) match(construct={target})
+ int bar (int c, int d) { return c + d + 3; }
+
+#pragma omp declare variant (t_qux) match(construct={target})
+ static int qux (int c, int d) { return c + d + 7; }
+};
+
+int main()
+{
+ foo fv;
+
+ printf ("fv.bar (1, 2)=%d\n", fv.bar (1, 2));
+ printf ("foo::qux (1, 2)=%d\n", foo::qux (1, 2));
+
+#pragma omp target
+ {
+ printf ("fv.bar (1, 2)=%d\n", fv.bar (1, 2));
+ printf ("foo::qux (1, 2)=%d\n", foo::qux (1, 2));
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,46 @@
+#include <cassert>
+#include <cstdio>
+
+#if 0
+#pragma omp begin declare variant match(device={arch(x86)})
+template<typename T>
+T foo (T v)
+{
+ return v + 3;
+}
+#pragma omp end declare variant
+#endif
+
+#pragma omp begin declare variant match(device={arch(powerpc)})
+template<typename T>
+T foo (T v)
+{
+ return v + 5;
+}
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(device={arch(aarch64)})
+template<typename T>
+T foo (T v)
+{
+ return v + 7;
+}
+#pragma omp end declare variant
+
+template<typename T>
+T foo (T v)
+{
+ return v + 1;
+}
+
+int main ()
+{
+ int res;
+#pragma omp dispatch novariants(1)
+ {
+ res = foo<int> (3);
+ }
+ printf ("dispatch foo(3)=%d\n", res);
+ printf ("regular foo(3)=%d\n", foo<int> (3));
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,41 @@
+#include <cassert>
+#include <cstdio>
+
+template<typename T>
+T foo_x86 (T v)
+{
+ return v + 3;
+}
+
+template<typename T>
+T foo_powerpc (T v)
+{
+ return v + 5;
+}
+
+template<typename T>
+T foo_aarch64 (T v)
+{
+ return v + 7;
+}
+
+#pragma omp declare variant (foo_x86) match(device={arch(x86)})
+#pragma omp declare variant (foo_powerpc) match(device={arch(powerpc)})
+#pragma omp declare variant (foo_aarch64) match(device={arch(arm)})
+template<typename T>
+T foo (T v)
+{
+ return v + 1;
+}
+
+int main ()
+{
+ int res;
+#pragma omp dispatch novariants(1)
+ {
+ res = foo<int> (3);
+ }
+ printf ("dispatch foo(3)=%d\n", res);
+ printf ("regular foo(3)=%d\n", foo<int> (3));
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,39 @@
+#include <cassert>
+#include <cstdio>
+
+typedef int T;
+
+T foo_x86 (T v)
+{
+ return v + 3;
+}
+
+T foo_powerpc (T v)
+{
+ return v + 5;
+}
+
+T foo_aarch64 (T v)
+{
+ return v + 7;
+}
+
+#pragma omp declare variant (foo_x86) match(device={arch(x86)})
+#pragma omp declare variant (foo_powerpc) match(device={arch(powerpc)})
+#pragma omp declare variant (foo_aarch64) match(device={arch(aarch64)})
+T foo (T v)
+{
+ return v + 1;
+}
+
+int main ()
+{
+ int res;
+#pragma omp dispatch novariants(1)
+ {
+ res = foo (3);
+ }
+ printf ("dispatch foo(3)=%d\n", res);
+ printf ("regular foo(3)=%d\n", foo (3));
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,56 @@
+#include <cstdio>
+
+template<typename T>
+class foo
+{
+ T var;
+public:
+ void set_var(T c) { var = c; }
+ T get_var () { return var + 1; }
+
+#pragma omp begin declare variant match(construct={target})
+ T get_var () { return var + 3; }
+#pragma omp end declare variant
+};
+
+void inty ()
+{
+ foo<int> fv;
+ fv.set_var (6);
+
+ printf ("fv.get_var () = %d\n", fv.get_var ());
+
+ int res;
+
+#pragma omp target map(from: res)
+ {
+ res = fv.get_var ();
+ }
+
+ printf ("fv.get_var () [target] = %d\n", res);
+}
+
+void longy ()
+{
+ foo<long> fvl;
+ fvl.set_var (7);
+
+ printf ("fvl.get_var () = %ld\n", fvl.get_var ());
+
+ long res;
+
+#pragma omp target map(from: res)
+ {
+ res = fvl.get_var ();
+ }
+
+ printf ("fvl.get_var () [target] = %ld\n", res);
+}
+
+int main ()
+{
+ inty ();
+ longy ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+#include <stdio.h>
+
+#pragma omp begin declare variant match(construct={target})
+int foo (void)
+{
+ return 7;
+}
+#pragma omp begin declare variant match(construct={teams})
+int foo (void)
+{
+ return 9;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+
+int foo (void)
+{
+ return 5;
+}
+
+int main (void)
+{
+ int c = foo ();
+ printf ("host c=%d\n", c);
+
+#pragma omp target map(from:c)
+ {
+ c = foo ();
+ }
+ printf ("target c=%d\n", c);
+#pragma omp target teams map(from:c)
+ {
+ c = foo ();
+ }
+ printf ("target teams c=%d\n", c);
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,38 @@
+#include <stdio.h>
+
+int p_foo (int c, int d)
+{
+ return (c+d) * 2;
+}
+
+int t_foo (int c, int d)
+{
+ return (c+d) * 3;
+}
+
+#pragma omp declare variant (p_foo) match(construct={parallel})
+#pragma omp declare variant (t_foo) match(construct={target})
+int foo (int c, int d)
+{
+ return c+d;
+}
+
+int main()
+{
+ printf ("host: %d\n", foo (3, 4));
+
+#pragma omp parallel
+ {
+#pragma omp single
+ {
+ printf ("parallel: %d\n", foo (3, 4));
+ }
+ }
+
+#pragma omp target
+ {
+ printf ("target: %d\n", foo (3, 4));
+ }
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,61 @@
+#include <stdio.h>
+
+#define VARIANTS
+
+#ifdef VARIANTS
+#pragma omp begin declare variant match(construct={target})
+template<typename T>
+T foo (void)
+{
+ return 7;
+}
+#pragma omp begin declare variant match(construct={teams})
+template<typename T>
+T foo (void)
+{
+ return 9;
+}
+#pragma omp end declare variant
+#pragma omp end declare variant
+#endif
+
+template<typename T>
+T foo (void)
+{
+ return 5;
+}
+
+int main (void)
+{
+ int c = foo<int> ();
+ printf ("host c=%d\n", c);
+
+#pragma omp target map(from:c)
+ {
+ c = foo<int> ();
+ }
+ printf ("target c=%d\n", c);
+#pragma omp target teams map(from:c)
+ {
+ c = foo<int> ();
+ }
+ printf ("target teams c=%d\n", c);
+
+ long cl = foo<long> ();
+ printf ("host cl=%ld\n", cl);
+
+#pragma omp target map(from:cl)
+ {
+ cl = foo<long> ();
+ }
+ printf ("target cl=%ld\n", cl);
+#pragma omp target teams map(from:cl)
+ {
+ cl = foo<long> ();
+ }
+ printf ("target teams cl=%ld\n", cl);
+
+ // int q = foo ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,54 @@
+#include <stdio.h>
+
+template<typename T>
+T foo_targ (void)
+{
+ return 7;
+}
+
+template<typename T>
+T foo_targteams (void)
+{
+ return 9;
+}
+
+#pragma omp declare variant(foo_targ<T>) match(construct={target})
+#pragma omp declare variant(foo_targteams<T>) match(construct={target,teams})
+template<typename T>
+T foo (void)
+{
+ return 5;
+}
+
+int main (void)
+{
+ int c = foo<int> ();
+ printf ("host c=%d\n", c);
+
+#pragma omp target map(from:c)
+ {
+ c = foo<int> ();
+ }
+ printf ("target c=%d\n", c);
+#pragma omp target teams map(from:c)
+ {
+ c = foo<int> ();
+ }
+ printf ("target teams c=%d\n", c);
+
+ long cl = foo<long> ();
+ printf ("host cl=%ld\n", cl);
+
+#pragma omp target map(from:cl)
+ {
+ cl = foo<long> ();
+ }
+ printf ("target cl=%ld\n", cl);
+#pragma omp target teams map(from:cl)
+ {
+ cl = foo<long> ();
+ }
+ printf ("target teams cl=%ld\n", cl);
+
+ return 0;
+}