@@ -24635,7 +24635,7 @@ c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context)
static tree
c_parser_omp_context_selector (c_parser *parser, enum omp_tss_code set,
- tree parms, bool metadirective_p)
+ tree parms)
{
tree ret = NULL_TREE;
do
@@ -24782,16 +24782,7 @@ c_parser_omp_context_selector (c_parser *parser, enum omp_tss_code set,
{
mark_exp_read (t);
t = c_fully_fold (t, false, NULL);
- /* FIXME: I believe it is an unimplemented feature rather
- than a user error to have non-constant expressions
- inside "declare variant". */
- if (!metadirective_p
- && (!INTEGRAL_TYPE_P (TREE_TYPE (t))
- || !tree_fits_shwi_p (t)))
- error_at (token->location,
- "property must be constant integer expression");
- else if (metadirective_p
- && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
error_at (token->location,
"property must be integer expression");
else
@@ -24874,8 +24865,7 @@ c_parser_omp_context_selector (c_parser *parser, enum omp_tss_code set,
user */
static tree
-c_parser_omp_context_selector_specification (c_parser *parser, tree parms,
- bool metadirective_p)
+c_parser_omp_context_selector_specification (c_parser *parser, tree parms)
{
tree ret = NULL_TREE;
do
@@ -24900,8 +24890,7 @@ c_parser_omp_context_selector_specification (c_parser *parser, tree parms,
if (!braces.require_open (parser))
return error_mark_node;
- tree selectors = c_parser_omp_context_selector (parser, set, parms,
- metadirective_p);
+ tree selectors = c_parser_omp_context_selector (parser, set, parms);
if (selectors == error_mark_node)
ret = error_mark_node;
else if (ret != error_mark_node)
@@ -24977,8 +24966,7 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
if (parms == NULL_TREE)
parms = error_mark_node;
- tree ctx = c_parser_omp_context_selector_specification (parser,
- parms, false);
+ tree ctx = c_parser_omp_context_selector_specification (parser, parms);
if (ctx == error_mark_node)
goto fail;
ctx = omp_check_context_selector (match_loc, ctx, false);
@@ -25014,7 +25002,7 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms)
tree construct
= omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT);
omp_mark_declare_variant (match_loc, variant, construct);
- if (omp_context_selector_matches (ctx, false, true))
+ if (omp_context_selector_matches (ctx, NULL_TREE, false))
{
tree attr
= tree_cons (get_identifier ("omp declare variant base"),
@@ -26688,7 +26676,7 @@ c_parser_omp_metadirective (c_parser *parser, bool *if_p)
if (!default_p)
{
ctx = c_parser_omp_context_selector_specification (parser,
- NULL_TREE, true);
+ NULL_TREE);
if (ctx == error_mark_node)
goto error;
ctx = omp_check_context_selector (match_loc, ctx, true);
@@ -26697,7 +26685,7 @@ c_parser_omp_metadirective (c_parser *parser, bool *if_p)
/* Remove the selector from further consideration if can be
evaluated as a non-match at this point. */
- skip = omp_context_selector_matches (ctx, true, true) == 0;
+ skip = (omp_context_selector_matches (ctx, NULL_TREE, false) == 0);
if (c_parser_next_token_is_not (parser, CPP_COLON))
{
@@ -1891,6 +1891,7 @@ struct GTY(()) saved_scope {
int suppress_location_wrappers;
BOOL_BITFIELD x_processing_explicit_instantiation : 1;
BOOL_BITFIELD need_pop_function_context : 1;
+ BOOL_BITFIELD x_processing_omp_trait_property_expr : 1;
/* Nonzero if we are parsing the discarded statement of a constexpr
if-statement. */
@@ -1962,6 +1963,7 @@ extern GTY(()) struct saved_scope *scope_chain;
#define processing_template_decl scope_chain->x_processing_template_decl
#define processing_specialization scope_chain->x_processing_specialization
#define processing_explicit_instantiation scope_chain->x_processing_explicit_instantiation
+#define processing_omp_trait_property_expr scope_chain->x_processing_omp_trait_property_expr
/* Nonzero if we are parsing the conditional expression of a contract
condition. These expressions appear outside the paramter list (like a
@@ -8328,7 +8328,7 @@ omp_declare_variant_finalize_one (tree decl, tree attr)
tree construct
= omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT);
omp_mark_declare_variant (match_loc, variant, construct);
- if (!omp_context_selector_matches (ctx, false, true))
+ if (!omp_context_selector_matches (ctx, NULL_TREE, false))
return true;
TREE_PURPOSE (TREE_VALUE (attr)) = variant;
}
@@ -47939,7 +47939,7 @@ cp_parser_omp_declare_simd (cp_parser *parser, cp_token *pragma_tok,
static tree
cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set,
- bool has_parms_p, bool metadirective_p)
+ bool has_parms_p)
{
tree ret = NULL_TREE;
do
@@ -48100,31 +48100,24 @@ cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set,
break;
case OMP_TRAIT_PROPERTY_DEV_NUM_EXPR:
case OMP_TRAIT_PROPERTY_BOOL_EXPR:
- /* FIXME: I believe it is an unimplemented feature rather
- than a user error to have non-constant expressions
- inside "declare variant". */
- t = metadirective_p
- ? cp_parser_expression (parser)
- : cp_parser_constant_expression (parser);
- if (t != error_mark_node)
+ processing_omp_trait_property_expr = true;
+ /* This actually parses a not-necessarily-constant
+ conditional-expression. */
+ t = cp_parser_constant_expression (parser, true, NULL, false);
+ processing_omp_trait_property_expr = false;
+ if (t == error_mark_node)
+ return error_mark_node;
+ if (!type_dependent_expression_p (t))
{
t = fold_non_dependent_expr (t);
- if (!metadirective_p
- && !value_dependent_expression_p (t)
- && (!INTEGRAL_TYPE_P (TREE_TYPE (t))
- || !tree_fits_shwi_p (t)))
- error_at (token->location, "property must be "
- "constant integer expression");
- if (metadirective_p
- && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
- error_at (token->location,
- "property must be integer expression");
- else
- properties = make_trait_property (NULL_TREE, t,
- properties);
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ error_at (token->location,
+ "property must be integer expression");
+ return error_mark_node;
+ }
}
- else
- return error_mark_node;
+ properties = make_trait_property (NULL_TREE, t, properties);
break;
case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
if (sel == OMP_TRAIT_CONSTRUCT_SIMD)
@@ -48197,8 +48190,7 @@ cp_parser_omp_context_selector (cp_parser *parser, enum omp_tss_code set,
static tree
cp_parser_omp_context_selector_specification (cp_parser *parser,
- bool has_parms_p,
- bool metadirective_p)
+ bool has_parms_p)
{
tree ret = NULL_TREE;
do
@@ -48225,8 +48217,7 @@ cp_parser_omp_context_selector_specification (cp_parser *parser,
return error_mark_node;
tree selectors
- = cp_parser_omp_context_selector (parser, set, has_parms_p,
- metadirective_p);
+ = cp_parser_omp_context_selector (parser, set, has_parms_p);
if (selectors == error_mark_node)
{
cp_parser_skip_to_closing_brace (parser);
@@ -48536,8 +48527,7 @@ cp_finish_omp_declare_variant (cp_parser *parser, cp_token *pragma_tok,
if (!parens.require_open (parser))
goto fail;
- tree ctx = cp_parser_omp_context_selector_specification (parser, true,
- false);
+ tree ctx = cp_parser_omp_context_selector_specification (parser, true);
if (ctx == error_mark_node)
goto fail;
ctx = omp_check_context_selector (match_loc, ctx, false);
@@ -49326,8 +49316,7 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok,
if (!default_p)
{
- ctx = cp_parser_omp_context_selector_specification (parser, false,
- true);
+ ctx = cp_parser_omp_context_selector_specification (parser, false);
if (ctx == error_mark_node)
goto fail;
ctx = omp_check_context_selector (match_loc, ctx, true);
@@ -49339,8 +49328,7 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok,
/* FIXME: we could still do this if the context selector
doesn't have any dependent subexpressions. */
skip = (!processing_template_decl
- && omp_context_selector_matches (ctx, true, true) == 0);
-
+ && !omp_context_selector_matches (ctx, NULL_TREE, false));
if (cp_lexer_next_token_is_not (parser->lexer, CPP_COLON))
{
cp_parser_require (parser, CPP_COLON, RT_COLON);
@@ -17906,7 +17906,8 @@ tsubst_omp_context_selector (tree ctx, tree args, tsubst_flags_t complain,
if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
error_at (cp_expr_loc_or_input_loc (t),
"property must be integer expression");
- properties = make_trait_property (NULL_TREE, t, NULL_TREE);
+ else
+ properties = make_trait_property (NULL_TREE, t, NULL_TREE);
break;
case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
if (OMP_TS_CODE (sel) == OMP_TRAIT_CONSTRUCT_SIMD)
@@ -19496,7 +19497,7 @@ tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl)
in_decl);
/* Remove the selector from further consideration if it can be
evaluated as a non-match at this point. */
- if (omp_context_selector_matches (ctx, true, true) == 0)
+ if (omp_context_selector_matches (ctx, NULL_TREE, false) == 0)
continue;
}
s = push_stmt_list ();
@@ -4384,7 +4384,8 @@ finish_id_expression_1 (tree id_expression,
if (TREE_CODE (decl) == PARM_DECL
&& DECL_CONTEXT (decl) == NULL_TREE
&& !cp_unevaluated_operand
- && !processing_contract_condition)
+ && !processing_contract_condition
+ && !processing_omp_trait_property_expr)
{
*error_msg = G_("use of parameter outside function body");
return error_mark_node;
@@ -8555,7 +8555,8 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
omp_mark_declare_variant (gfc_get_location (&odv->where),
gfc_get_symbol_decl (variant_proc_sym),
construct);
- if (omp_context_selector_matches (set_selectors, false, true))
+ if (omp_context_selector_matches (set_selectors,
+ NULL_TREE, false))
{
tree id = get_identifier ("omp declare variant base");
tree variant = gfc_get_symbol_decl (variant_proc_sym);
@@ -8627,7 +8628,7 @@ gfc_trans_omp_metadirective (gfc_code *code)
return error_mark_node;
/* If the selector doesn't match, drop the whole variant. */
- if (!omp_context_selector_matches (ctx, true, true))
+ if (!omp_context_selector_matches (ctx, NULL_TREE, false))
{
variant = variant->next;
continue;
@@ -193,6 +193,9 @@ input_gimple_stmt (class lto_input_block *ib, class data_in *data_in,
= dyn_cast <gomp_metadirective*> (stmt))
{
gimple_alloc_omp_metadirective (metadirective_stmt);
+ gimple_omp_metadirective_set_context (metadirective_stmt,
+ stream_read_tree (ib,
+ data_in));
for (i = 0; i < num_ops; i++)
gimple_omp_metadirective_set_label (metadirective_stmt, i,
stream_read_tree (ib,
@@ -171,9 +171,13 @@ output_gimple_stmt (struct output_block *ob, struct function *fn, gimple *stmt)
stream_write_tree (ob, gimple_call_fntype (stmt), true);
}
if (gimple_code (stmt) == GIMPLE_OMP_METADIRECTIVE)
- for (i = 0; i < gimple_num_ops (stmt); i++)
- stream_write_tree (ob, gimple_omp_metadirective_label (stmt, i),
+ {
+ stream_write_tree (ob, gimple_omp_metadirective_context (stmt),
true);
+ for (i = 0; i < gimple_num_ops (stmt); i++)
+ stream_write_tree (ob, gimple_omp_metadirective_label (stmt, i),
+ true);
+ }
break;
@@ -1331,6 +1331,7 @@ gimple_build_omp_metadirective (int num_variants)
= as_a <gomp_metadirective *> (gimple_alloc (GIMPLE_OMP_METADIRECTIVE,
num_variants));
gimple_alloc_omp_metadirective (p);
+ gimple_omp_metadirective_set_context (p, NULL);
gimple_omp_metadirective_set_variants (p, NULL);
return p;
@@ -398,7 +398,8 @@ DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_PARALLEL_LAYOUT)
CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_omp_ordered", GSS_OMP_SINGLE_LAYOUT)
-/* GIMPLE_OMP_METADIRECTIVE represents #pragma omp metadirective. */
+/* GIMPLE_OMP_METADIRECTIVE represents both #pragma omp metadirective and
+ a call to a function with "declare variant" variants. */
DEFGSCODE(GIMPLE_OMP_METADIRECTIVE, "gimple_omp_metadirective",
GSS_OMP_METADIRECTIVE)
@@ -856,10 +856,14 @@ struct GTY((tag("GSS_OMP_METADIRECTIVE")))
/* [ WORD 8 ] : a list of bodies associated with the directive variants. */
gomp_variant *variants;
- /* [ WORD 9 ] : label vector. */
+ /* [ WORD 9 ] : the cached OpenMP context for this directive, used for
+ post-gimplification resolution. */
+ tree context;
+
+ /* [ WORD 10 ] : label vector. */
tree * GTY((length ("%h.num_ops"))) labels;
- /* [ WORD 10 ] : operand vector. Used to hold the selectors for the
+ /* [ WORD 11 ] : operand vector. Used to hold the selectors for the
directive variants. */
tree GTY((length ("%h.num_ops"))) op[1];
};
@@ -6692,6 +6696,24 @@ gimple_assume_body (const gimple *gs)
}
+static inline tree
+gimple_omp_metadirective_context (const gimple *g)
+{
+ const gomp_metadirective *omp_metadirective
+ = as_a <const gomp_metadirective *> (g);
+ return omp_metadirective->context;
+}
+
+
+static inline void
+gimple_omp_metadirective_set_context (gimple *g, tree context)
+{
+ gomp_metadirective *omp_metadirective
+ = as_a <gomp_metadirective *> (g);
+ omp_metadirective->context = context;
+}
+
+
static inline tree
gimple_omp_metadirective_label (const gimple *g, unsigned i)
{
@@ -96,6 +96,11 @@ reset_cond_uid ()
/* Hash set of poisoned variables in a bind expr. */
static hash_set<tree> *asan_poisoned_variables = NULL;
+/* Hash set of already-resolved calls to OpenMP "declare variant"
+ functions. A call can resolve to the original function and
+ we don't want to repeat the resolution multiple times. */
+static hash_set<tree> *omp_resolved_variant_calls = NULL;
+
enum gimplify_omp_var_data
{
GOVD_SEEN = 0x000001,
@@ -3799,12 +3804,185 @@ maybe_fold_stmt (gimple_stmt_iterator *gsi)
return fold_stmt (gsi);
}
+/* Helper function for gimplify_call_expr: handle "declare variant"
+ resolution and expansion. Arguments are as for gimplify_call_expr.
+ If *EXPR_P is unchanged, the return value should be ignored and the
+ normal gimplify_call_expr handling should be applied. Otherwise GS_OK
+ is returned if the new *EXPR_P is something that needs to be further
+ gimplified, or GS_ALL_DONE if *EXPR_P has been translated into a
+ GIMPLE_OMP_METADIRECTIVE. */
+
+static enum gimplify_status
+gimplify_variant_call_expr (tree *expr_p, gimple_seq *pre_p,
+ fallback_t fallback)
+{
+ /* If we've already processed this call, stop now. This can happen
+ if the variant call resolves to the original function, or to
+ a dynamic conditional that includes the default call to the original
+ function. */
+ gcc_assert (omp_resolved_variant_calls != NULL);
+ if (omp_resolved_variant_calls->contains (*expr_p))
+ return GS_OK;
+
+ tree fndecl = get_callee_fndecl (*expr_p);
+ tree fnptrtype = TREE_TYPE (CALL_EXPR_FN (*expr_p));
+ location_t loc = EXPR_LOCATION (*expr_p);
+ tree construct_context = omp_get_construct_context ();
+ vec<struct omp_variant> all_candidates
+ = omp_declare_variant_candidates (fndecl, construct_context);
+ gcc_assert (!all_candidates.is_empty ());
+ vec<struct omp_variant> candidates
+ = omp_get_dynamic_candidates (all_candidates, construct_context);
+
+ /* If the variant call could be resolved now, build a nest of COND_EXPRs
+ if there are dynamic candidates, and/or a new CALL_EXPR for each
+ candidate call. */
+ if (!candidates.is_empty ())
+ {
+ int n = candidates.length ();
+ tree tail = NULL_TREE;
+
+ for (int i = n - 1; i >= 0; i--)
+ {
+ if (tail)
+ gcc_assert (candidates[i].dynamic_selector);
+ else
+ gcc_assert (!candidates[i].dynamic_selector);
+ if (candidates[i].alternative == fndecl)
+ {
+ /* We should only get the original function back as the
+ default. */
+ gcc_assert (!tail);
+ omp_resolved_variant_calls->add (*expr_p);
+ tail = *expr_p;
+ }
+ else
+ {
+ /* Special case: if there are no adjust_args/append_args for
+ the final static selector, we can re-use the old CALL_EXPR
+ and just replace the function. */
+ /* FIXME: also test for adjust_args/append_args. */
+ tree thiscall = tail ? unshare_expr (*expr_p) : *expr_p;
+ CALL_EXPR_FN (thiscall) = build1 (ADDR_EXPR, fnptrtype,
+ candidates[i].alternative);
+ /* FIXME: Handle adjust_args/append_args here too. */
+ if (!tail)
+ tail = thiscall;
+ else
+ tail = build3 (COND_EXPR, TREE_TYPE (*expr_p),
+ candidates[i].dynamic_selector,
+ thiscall, tail);
+ }
+ }
+ *expr_p = tail;
+ return GS_OK;
+ }
+
+ /* If we couldn't resolve the variant call now, replace the call with a
+ GIMPLE_OMP_METADIRECTIVE that has gimplified calls in each of its
+ alternatives. Yes, this is badly named, but the same logic is used
+ to replace both things in the late resolution from the ompdevlow pass. */
+ else
+ {
+ /* If we need a usable return value, we need a temporary
+ and an assignment in each alternative. This logic was borrowed
+ from gimplify_cond_expr. */
+ tree type = TREE_TYPE (*expr_p);
+ bool want_value = (fallback != fb_none && !VOID_TYPE_P (type));
+ bool pointerize = false;
+ tree tmp = NULL_TREE, result = NULL_TREE;
+
+ if (want_value)
+ {
+ /* If either an rvalue is ok or we do not require an lvalue,
+ create the temporary. But we cannot do that if the type is
+ addressable. */
+ if (((fallback & fb_rvalue) || !(fallback & fb_lvalue))
+ && !TREE_ADDRESSABLE (type))
+ {
+ tmp = create_tmp_var (type, "iftmp");
+ result = tmp;
+ }
+
+ /* Otherwise, only create and copy references to the values. */
+ else
+ {
+ pointerize = true;
+ type = build_pointer_type (type);
+ tmp = create_tmp_var (type, "iftmp");
+ result = build_simple_mem_ref_loc (loc, tmp);
+ }
+ }
+
+ /* The following code was more or less stolen from
+ gimplify_omp_metadirective. FIXME: do we also need to copy
+ the "omp metadirective construct target" part too? */
+ gomp_variant *first_variant = NULL;
+ gomp_variant *prev_variant = NULL;
+ gomp_metadirective *stmt
+ = gimple_build_omp_metadirective (all_candidates.length ());
+ tree end_label = create_artificial_label (UNKNOWN_LOCATION);
+
+ for (unsigned int i = 0; i < all_candidates.length (); i++)
+ {
+ tree decl = all_candidates[i].alternative;
+ gimple_set_op (stmt, i, all_candidates[i].selector);
+ gomp_variant *omp_variant
+ = gimple_build_omp_variant (NULL);
+ gimple_seq *directive_p = gimple_omp_body_ptr (omp_variant);
+ tree thiscall;
+
+ /* We need to turn the decl from the candidate into a function
+ call and possible assignment, gimplify it, and stuff that in
+ the directive seq of the gomp_variant. */
+ if (decl == fndecl)
+ {
+ thiscall = *expr_p;
+ omp_resolved_variant_calls->add (*expr_p);
+ }
+ else
+ {
+ /* FIXME: handle adjust_args/append_args here too. */
+ thiscall = unshare_expr (*expr_p);
+ CALL_EXPR_FN (thiscall) = build1 (ADDR_EXPR, fnptrtype, decl);
+ }
+ if (pointerize)
+ thiscall = build_fold_addr_expr_loc (loc, thiscall);
+ if (want_value)
+ thiscall = build2 (INIT_EXPR, type, tmp, thiscall);
+
+ gimplify_stmt (&thiscall, directive_p);
+ gimplify_seq_add_stmt (directive_p, gimple_build_goto (end_label));
+
+ if (!first_variant)
+ first_variant = omp_variant;
+ if (prev_variant)
+ {
+ prev_variant->next = omp_variant;
+ omp_variant->prev = prev_variant;
+ }
+ prev_variant = omp_variant;
+ }
+
+ gimple_omp_metadirective_set_context (stmt, construct_context);
+ gimple_omp_metadirective_set_variants (stmt, first_variant);
+
+ gimplify_seq_add_stmt (pre_p, stmt);
+ gimplify_seq_add_stmt (pre_p, gimple_build_label (end_label));
+ cgraph_node::get (cfun->decl)->has_metadirectives = 1;
+
+ *expr_p = result;
+ return GS_ALL_DONE;
+ }
+}
+
/* Gimplify the CALL_EXPR node *EXPR_P into the GIMPLE sequence PRE_P.
WANT_VALUE is true if the result of the call is desired. */
static enum gimplify_status
-gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
+gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, fallback_t fallback)
{
+ bool want_value = (fallback != fb_none);
tree fndecl, parms, p, fnptrtype;
enum gimplify_status ret;
int i, nargs;
@@ -3970,14 +4148,23 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
/* Remember the original function pointer type. */
fnptrtype = TREE_TYPE (CALL_EXPR_FN (*expr_p));
+ /* Handle "declare variant" substitution. */
if (flag_openmp
&& fndecl
&& cfun
- && (cfun->curr_properties & PROP_gimple_any) == 0)
+ && (cfun->curr_properties & PROP_gimple_any) == 0
+ && lookup_attribute ("omp declare variant base",
+ DECL_ATTRIBUTES (fndecl)))
{
- tree variant = omp_resolve_declare_variant (fndecl);
- if (variant != fndecl)
- CALL_EXPR_FN (*expr_p) = build1 (ADDR_EXPR, fnptrtype, variant);
+ tree orig = *expr_p;
+ enum gimplify_status ret
+ = gimplify_variant_call_expr (expr_p, pre_p, fallback);
+ /* This may resolve to the same call, or the call expr with just
+ the function replaced, in which case we should just continue to
+ gimplify it normally. Otherwise, if we get something else back,
+ stop here. */
+ if (*expr_p != orig)
+ return ret;
}
/* There is a sequence point before the call, so any side effects in
@@ -14634,57 +14821,59 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
delete_omp_context (ctx);
}
-/* Return 0 if CONSTRUCTS selectors don't match the OpenMP context,
- -1 if unknown yet (simd is involved, won't be known until vectorization)
- and 1 if they do. If SCORES is non-NULL, it should point to an array
- of at least 2*NCONSTRUCTS+2 ints, and will be filled with the positions
- of the CONSTRUCTS (position -1 if it will never match) followed by
- number of constructs in the OpenMP context construct trait. If the
- score depends on whether it will be in a declare simd clone or not,
- the function returns 2 and there will be two sets of the scores, the first
- one for the case that it is not in a declare simd clone, the other
- that it is in a declare simd clone. */
+/* Collect a list of traits for enclosing constructs in the current
+ OpenMP context. The list is in the same format as the trait selector
+ list of construct trait sets built by the front ends.
-int
-omp_construct_selector_matches (enum tree_code *constructs, int nconstructs,
- int *scores)
+ Per the OpenMP specification, the construct trait set includes constructs
+ up to an enclosing "target" construct. If there is no "target" construct,
+ then additional things may be added to the construct trait set (simd for
+ simd clones, additional constructs associated with "declare variant",
+ the target trait for "declare target"); those are not handled here.
+ In particular simd clones are not known during gimplification so
+ matching/scoring of context selectors that might involve them needs
+ to be deferred to the omp_device_lower pass. */
+
+tree
+omp_get_construct_context (void)
{
- int matched = 0, cnt = 0;
- bool simd_seen = false;
- bool target_seen = false;
- int declare_simd_cnt = -1;
- auto_vec<enum tree_code, 16> codes;
+ tree result = NULL_TREE;
for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx;)
{
- if (((ctx->region_type & ORT_PARALLEL) && ctx->code == OMP_PARALLEL)
- || ((ctx->region_type & (ORT_TARGET | ORT_IMPLICIT_TARGET | ORT_ACC))
- == ORT_TARGET && ctx->code == OMP_TARGET)
- || ((ctx->region_type & ORT_TEAMS) && ctx->code == OMP_TEAMS)
- || (ctx->region_type == ORT_WORKSHARE && ctx->code == OMP_FOR)
- || (ctx->region_type == ORT_SIMD
- && ctx->code == OMP_SIMD
- && !omp_find_clause (ctx->clauses, OMP_CLAUSE_BIND)))
+ if (((ctx->region_type & (ORT_TARGET | ORT_IMPLICIT_TARGET | ORT_ACC))
+ == ORT_TARGET)
+ && ctx->code == OMP_TARGET)
{
- ++cnt;
- if (scores)
- codes.safe_push (ctx->code);
- else if (matched < nconstructs && ctx->code == constructs[matched])
- {
- if (ctx->code == OMP_SIMD)
- {
- if (matched)
- return 0;
- simd_seen = true;
- }
- ++matched;
- }
- if (ctx->code == OMP_TARGET)
- {
- if (scores == NULL)
- return matched < nconstructs ? 0 : simd_seen ? -1 : 1;
- target_seen = true;
- break;
- }
+ result = make_trait_selector (OMP_TRAIT_CONSTRUCT_TARGET,
+ NULL_TREE, NULL_TREE, result);
+ /* We're not interested in any outer constructs. */
+ break;
+ }
+ else if ((ctx->region_type & ORT_PARALLEL) && ctx->code == OMP_PARALLEL)
+ result = make_trait_selector (OMP_TRAIT_CONSTRUCT_PARALLEL,
+ NULL_TREE, NULL_TREE, result);
+ else if ((ctx->region_type & ORT_TEAMS) && ctx->code == OMP_TEAMS)
+ result = make_trait_selector (OMP_TRAIT_CONSTRUCT_TEAMS,
+ NULL_TREE, NULL_TREE, result);
+ else if (ctx->region_type == ORT_WORKSHARE && ctx->code == OMP_FOR)
+ result = make_trait_selector (OMP_TRAIT_CONSTRUCT_FOR,
+ NULL_TREE, NULL_TREE, result);
+ else if (ctx->region_type == ORT_SIMD
+ && ctx->code == OMP_SIMD
+ && !omp_find_clause (ctx->clauses, OMP_CLAUSE_BIND))
+ {
+ tree props = NULL_TREE;
+ tree *last = &props;
+ for (tree c = ctx->clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SIMDLEN
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_INBRANCH
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NOTINBRANCH)
+ {
+ *last = unshare_expr (c);
+ last = &(OMP_CLAUSE_CHAIN (c));
+ }
+ result = make_trait_selector (OMP_TRAIT_CONSTRUCT_SIMD,
+ NULL_TREE, props, result);
}
else if (ctx->region_type == ORT_WORKSHARE
&& ctx->code == OMP_LOOP
@@ -14696,89 +14885,8 @@ omp_construct_selector_matches (enum tree_code *constructs, int nconstructs,
ctx = ctx->outer_context->outer_context;
ctx = ctx->outer_context;
}
- if (!target_seen
- && lookup_attribute ("omp declare simd",
- DECL_ATTRIBUTES (current_function_decl)))
- {
- /* Declare simd is a maybe case, it is supposed to be added only to the
- omp-simd-clone.cc added clones and not to the base function. */
- declare_simd_cnt = cnt++;
- if (scores)
- codes.safe_push (OMP_SIMD);
- else if (cnt == 0
- && constructs[0] == OMP_SIMD)
- {
- gcc_assert (matched == 0);
- simd_seen = true;
- if (++matched == nconstructs)
- return -1;
- }
- }
- if (tree attr = lookup_attribute ("omp declare variant variant",
- DECL_ATTRIBUTES (current_function_decl)))
- {
- tree selectors = TREE_VALUE (attr);
- int variant_nconstructs = list_length (selectors);
- enum tree_code *variant_constructs = NULL;
- if (!target_seen && variant_nconstructs)
- {
- variant_constructs
- = (enum tree_code *) alloca (variant_nconstructs
- * sizeof (enum tree_code));
- omp_construct_traits_to_codes (selectors, variant_nconstructs,
- variant_constructs);
- }
- for (int i = 0; i < variant_nconstructs; i++)
- {
- ++cnt;
- if (scores)
- codes.safe_push (variant_constructs[i]);
- else if (matched < nconstructs
- && variant_constructs[i] == constructs[matched])
- {
- if (variant_constructs[i] == OMP_SIMD)
- {
- if (matched)
- return 0;
- simd_seen = true;
- }
- ++matched;
- }
- }
- }
- if (!target_seen
- && lookup_attribute ("omp declare target block",
- DECL_ATTRIBUTES (current_function_decl)))
- {
- if (scores)
- codes.safe_push (OMP_TARGET);
- else if (matched < nconstructs && constructs[matched] == OMP_TARGET)
- ++matched;
- }
- if (scores)
- {
- for (int pass = 0; pass < (declare_simd_cnt == -1 ? 1 : 2); pass++)
- {
- int j = codes.length () - 1;
- for (int i = nconstructs - 1; i >= 0; i--)
- {
- while (j >= 0
- && (pass != 0 || declare_simd_cnt != j)
- && constructs[i] != codes[j])
- --j;
- if (pass == 0 && declare_simd_cnt != -1 && j > declare_simd_cnt)
- *scores++ = j - 1;
- else
- *scores++ = j;
- }
- *scores++ = ((pass == 0 && declare_simd_cnt != -1)
- ? codes.length () - 1 : codes.length ());
- }
- return declare_simd_cnt == -1 ? 1 : 2;
- }
- if (matched == nconstructs)
- return simd_seen ? -1 : 1;
- return 0;
+
+ return result;
}
/* Gimplify OACC_CACHE. */
@@ -17720,8 +17828,11 @@ gimplify_omp_metadirective (tree *expr_p, gimple_seq *pre_p, gimple_seq *,
}
/* Try to resolve the metadirective. */
+ tree construct_context = omp_get_construct_context ();
+ vec<struct omp_variant> all_candidates
+ = omp_metadirective_candidates (*expr_p, construct_context);
vec<struct omp_variant> candidates
- = omp_early_resolve_metadirective (*expr_p);
+ = omp_get_dynamic_candidates (all_candidates, construct_context);
if (!candidates.is_empty ())
return expand_omp_metadirective (candidates, pre_p);
@@ -17733,12 +17844,11 @@ gimplify_omp_metadirective (tree *expr_p, gimple_seq *pre_p, gimple_seq *,
tree body_label = NULL;
tree end_label = create_artificial_label (UNKNOWN_LOCATION);
- for (tree variant = OMP_METADIRECTIVE_VARIANTS (*expr_p); variant != NULL_TREE;
- variant = TREE_CHAIN (variant))
+ for (unsigned int i = 0; i < all_candidates.length (); i++)
{
- tree selector = OMP_METADIRECTIVE_VARIANT_SELECTOR (variant);
- tree directive = OMP_METADIRECTIVE_VARIANT_DIRECTIVE (variant);
- tree body = OMP_METADIRECTIVE_VARIANT_BODY (variant);
+ tree selector = all_candidates[i].selector;
+ tree directive = all_candidates[i].alternative;
+ tree body = all_candidates[i].body;
selectors.safe_push (selector);
gomp_variant *omp_variant
@@ -17770,6 +17880,7 @@ gimplify_omp_metadirective (tree *expr_p, gimple_seq *pre_p, gimple_seq *,
gomp_metadirective *stmt
= gimple_build_omp_metadirective (selectors.length ());
+ gimple_omp_metadirective_set_context (stmt, construct_context);
gimple_omp_metadirective_set_variants (stmt, first_variant);
tree selector;
@@ -18029,7 +18140,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
break;
case CALL_EXPR:
- ret = gimplify_call_expr (expr_p, pre_p, fallback != fb_none);
+ ret = gimplify_call_expr (expr_p, pre_p, fallback);
/* C99 code may assign to an array in a structure returned
from a function, and this has undefined behavior only on
@@ -19533,7 +19644,16 @@ gimplify_function_tree (tree fndecl)
if (asan_sanitize_use_after_scope ())
asan_poisoned_variables = new hash_set<tree> ();
+ if (flag_openmp)
+ omp_resolved_variant_calls = new hash_set<tree> ();
+
bind = gimplify_body (fndecl, true);
+
+ if (omp_resolved_variant_calls)
+ {
+ delete omp_resolved_variant_calls;
+ omp_resolved_variant_calls = NULL;
+ }
if (asan_poisoned_variables)
{
delete asan_poisoned_variables;
@@ -76,7 +76,7 @@ extern void omp_firstprivatize_variable (struct gimplify_omp_ctx *, tree);
extern enum gimplify_status gimplify_expr (tree *, gimple_seq *, gimple_seq *,
bool (*) (tree), fallback_t);
-int omp_construct_selector_matches (enum tree_code *, int, int *);
+extern tree omp_get_construct_context (void);
extern void gimplify_type_sizes (tree, gimple_seq *);
extern void gimplify_one_sizepos (tree *, gimple_seq *);
@@ -1024,31 +1024,6 @@ omp_max_simt_vf (void)
return 0;
}
-/* Store the construct selectors as tree codes from last to first.
- CTX is a list of trait selectors, nconstructs must be equal to its
- length, and the array CONSTRUCTS holds the output. */
-
-void
-omp_construct_traits_to_codes (tree ctx, int nconstructs,
- enum tree_code *constructs)
-{
- int i = nconstructs - 1;
-
- /* Order must match the OMP_TRAIT_CONSTRUCT_* enumerators in
- enum omp_ts_code. */
- static enum tree_code code_map[]
- = { OMP_TARGET, OMP_TEAMS, OMP_PARALLEL, OMP_FOR, OMP_SIMD };
-
- for (tree ts = ctx; ts; ts = TREE_CHAIN (ts), i--)
- {
- enum omp_ts_code sel = OMP_TS_CODE (ts);
- int j = (int)sel - (int)OMP_TRAIT_CONSTRUCT_TARGET;
- gcc_assert (j >= 0 && (unsigned int) j < ARRAY_SIZE (code_map));
- constructs[i] = code_map[j];
- }
- gcc_assert (i == -1);
-}
-
/* Return true if PROP is possibly present in one of the offloading target's
OpenMP contexts. The format of PROPS string is always offloading target's
name terminated by '\0', followed by properties for that offloading
@@ -1096,29 +1071,36 @@ omp_offload_device_kind_arch_isa (const char *props, const char *prop)
region or when unsure, return false otherwise. */
static bool
-omp_maybe_offloaded (void)
+omp_maybe_offloaded (tree construct_context)
{
+ /* No offload targets available? */
if (!ENABLE_OFFLOADING)
return false;
const char *names = getenv ("OFFLOAD_TARGET_NAMES");
if (names == NULL || *names == '\0')
return false;
+ /* Parsing is too early to tell. */
if (symtab->state == PARSING)
/* Maybe. */
return true;
+
+ /* Late resolution of offloaded code happens in the offload compiler,
+ where it's treated as native code instead. So return false here. */
if (cfun && cfun->after_inlining)
return false;
+
+ /* The function is explicitly offloaded? */
if (current_function_decl
&& lookup_attribute ("omp declare target",
DECL_ATTRIBUTES (current_function_decl)))
return true;
- if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0)
- {
- enum tree_code construct = OMP_TARGET;
- if (omp_construct_selector_matches (&construct, 1, NULL))
- return true;
- }
+
+ /* Check for nesting inside a target directive. */
+ for (tree ts = construct_context; ts; ts = TREE_CHAIN (ts))
+ if (OMP_TS_CODE (ts) == OMP_TRAIT_CONSTRUCT_TARGET)
+ return true;
+
return false;
}
@@ -1265,6 +1247,18 @@ omp_context_name_list_prop (tree prop)
}
}
+
+/* Helper function called via walk_tree, to determine if *TP is a
+ PARM_DECL. */
+static tree
+expr_uses_parm_decl (tree *tp, int *walk_subtrees ATTRIBUTE_UNUSED,
+ void *data ATTRIBUTE_UNUSED)
+{
+ if (TREE_CODE (*tp) == PARM_DECL)
+ return *tp;
+ return NULL_TREE;
+}
+
/* Diagnose errors in an OpenMP context selector, return CTX if
it is correct or error_mark_node otherwise. */
@@ -1280,11 +1274,6 @@ omp_check_context_selector (location_t loc, tree ctx, bool metadirective_p)
bool saw_any_prop = false;
bool saw_other_prop = false;
- /* FIXME: not implemented yet. */
- if (!metadirective_p && tss_code == OMP_TRAIT_SET_TARGET_DEVICE)
- sorry_at (loc, "%<target_device%> selector set is not supported "
- "yet for %<declare variant%>");
-
/* Each trait-set-selector-name can only be specified once. */
if (tss_seen[tss_code])
{
@@ -1337,6 +1326,33 @@ omp_check_context_selector (location_t loc, tree ctx, bool metadirective_p)
|| ts_code == OMP_TRAIT_DEVICE_NUM)
saw_other_prop = true;
+ /* This restriction is documented in the spec in the section
+ for the metadirective "when" clause (7.4.1 in the 5.2 spec). */
+ if (metadirective_p && ts_code == OMP_TRAIT_CONSTRUCT_SIMD)
+ {
+ error_at (loc,
+ "properties must not be specified for the %<simd%> "
+ "selector in a %<metadirective%> context-selector");
+ return error_mark_node;
+ }
+
+ /* Reject expressions that reference parameter variables in
+ "declare variant", as this is not yet implemented. FIXME;
+ see PR middle-end/113094. */
+ if (!metadirective_p
+ && (ts_code == OMP_TRAIT_DEVICE_NUM
+ || ts_code == OMP_TRAIT_USER_CONDITION))
+ {
+ tree exp = OMP_TS_PROPERTIES (ts);
+ if (walk_tree (&exp, expr_uses_parm_decl, NULL, NULL))
+ {
+ sorry_at (loc,
+ "reference to function parameter in "
+ "%<declare variant%> dynamic selector expression");
+ return error_mark_node;
+ }
+ }
+
if (omp_ts_map[ts_code].valid_properties == NULL)
continue;
@@ -1401,6 +1417,9 @@ omp_check_context_selector (location_t loc, tree ctx, bool metadirective_p)
return ctx;
}
+/* Forward declarations. */
+static int omp_context_selector_set_compare (enum omp_tss_code, tree, tree);
+static int omp_construct_simd_compare (tree, tree, bool);
/* Register VARIANT as variant of some base function marked with
#pragma omp declare variant. CONSTRUCT is corresponding list of
@@ -1471,24 +1490,118 @@ make_omp_metadirective_variant (tree selector, tree directive, tree body)
return build_tree_list (selector, build_tree_list (directive, body));
}
+/* If the construct selector traits SELECTOR_TRAITS match the corresponding
+ OpenMP context traits CONTEXT_TRAITS, return true and set *SCORE to the
+ corresponding score if it is non-null. */
+static bool
+omp_construct_traits_match (tree selector_traits, tree context_traits,
+ score_wide_int *score)
+{
+ int slength = list_length (selector_traits);
+ int clength = list_length (context_traits);
-/* Return 1 if context selector matches the current OpenMP context, 0
+ /* Trivial failure: the selector has more traits than the OpenMP context. */
+ if (slength > clength)
+ return false;
+
+ /* There's only one trait in the selector and it doesn't have any properties
+ to match. */
+ if (slength == 1 && !OMP_TS_PROPERTIES (selector_traits))
+ {
+ int p = 0, i = 1;
+ enum omp_ts_code code = OMP_TS_CODE (selector_traits);
+ for (tree t = context_traits; t; t = TREE_CHAIN (t), i++)
+ if (OMP_TS_CODE (t) == code)
+ p = i;
+ if (p != 0)
+ {
+ if (score)
+ *score = wi::shifted_mask <score_wide_int> (p - 1, 1, false);
+ return true;
+ }
+ else
+ return false;
+ }
+
+ /* Now handle the more general cases.
+ Both lists of traits are ordered from outside in, corresponding to
+ the c1, ..., cN numbering for the OpenMP context specified in
+ in section 7.1 of the OpenMP 5.2 spec. Section 7.3 of the spec says
+ "if the traits that correspond to the construct selector set appear
+ multiple times in the OpenMP context, the highest valued subset of
+ context traits that contains all trait selectors in the same order
+ are used". This means that we want to start the search for a match
+ from the end of the list, rather than the beginning. To facilitate
+ that, transfer the lists to temporary arrays to allow random access
+ to the elements (their order remains outside in). */
+ int i, j;
+ tree s, c;
+
+ tree *sarray = (tree *) alloca (slength * sizeof (tree));
+ for (s = selector_traits, i = 0; s; s = TREE_CHAIN (s), i++)
+ sarray[i] = s;
+
+ tree *carray = (tree *) alloca (clength * sizeof (tree));
+ for (c = context_traits, j = 0; c; c = TREE_CHAIN (c), j++)
+ carray[j] = c;
+
+ /* The variable "i" indexes the selector, "j" indexes the OpenMP context.
+ Find the "j" corresponding to each sarray[i]. Note that the spec uses
+ "p" as the 1-based position, but "j" is zero-based, e.g. equal to
+ p - 1. */
+ score_wide_int result = 0;
+ j = clength - 1;
+ for (i = slength - 1; i >= 0; i--)
+ {
+ enum omp_ts_code code = OMP_TS_CODE (sarray[i]);
+ tree props = OMP_TS_PROPERTIES (sarray[i]);
+ for (; j >= 0; j--)
+ {
+ if (OMP_TS_CODE (carray[j]) != code)
+ continue;
+ if (code == OMP_TRAIT_CONSTRUCT_SIMD
+ && props
+ && omp_construct_simd_compare (props,
+ OMP_TS_PROPERTIES (carray[j]),
+ true) > 0)
+ continue;
+ break;
+ }
+ /* If j >= 0, we have a match for this trait at position j. */
+ if (j < 0)
+ return false;
+ result += wi::shifted_mask <score_wide_int> (j, 1, false);
+ j--;
+ }
+ if (score)
+ *score = result;
+ return true;
+}
+
+/* Return 1 if context selector CTX matches the current OpenMP context, 0
if it does not and -1 if it is unknown and need to be determined later.
- Some properties can be checked right away during parsing (this routine),
- others need to wait until the whole TU is parsed, others need to wait until
+ Some properties can be checked right away during parsing, others need
+ to wait until the whole TU is parsed, others need to wait until
IPA, others until vectorization.
- METADIRECTIVE_P is true if this is a metadirective context, and DELAY_P
- is true if it's too early in compilation to determine whether some
- properties match.
+ CONSTRUCT_CONTEXT is a list of construct traits from the OpenMP context,
+ which must be collected by omp_get_construct_context during
+ gimplification. It is ignored (and may be null) if this function is
+ called during parsing. Otherwise COMPLETE_P should indicate whether
+ CONSTRUCT_CONTEXT is known to be complete and not missing constructs
+ filled in later during compilation.
Dynamic properties (which are evaluated at run-time) should always
return 1. */
int
-omp_context_selector_matches (tree ctx, bool metadirective_p, bool delay_p)
+omp_context_selector_matches (tree ctx,
+ tree construct_context,
+ bool complete_p)
{
int ret = 1;
+ bool maybe_offloaded = omp_maybe_offloaded (construct_context);
+
for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
{
enum omp_tss_code set = OMP_TSS_CODE (tss);
@@ -1502,61 +1615,45 @@ omp_context_selector_matches (tree ctx, bool metadirective_p, bool delay_p)
if (set == OMP_TRAIT_SET_CONSTRUCT)
{
- /* For now, ignore the construct set. While something can be
- determined already during parsing, we don't know until end of TU
- whether additional constructs aren't added through declare variant
- unless "omp declare variant variant" attribute exists already
- (so in most of the cases), and we'd need to maintain set of
- surrounding OpenMP constructs, which is better handled during
- gimplification. */
+ /* We cannot resolve the construct selector during parsing because
+ the OpenMP context (and CONSTRUCT_CONTEXT) isn't available
+ until gimplification. */
if (symtab->state == PARSING)
{
ret = -1;
continue;
}
- int nconstructs = list_length (selectors);
- enum tree_code *constructs = NULL;
- if (nconstructs)
- {
- /* Even though this alloca appears in a loop over selector
- sets, it does not repeatedly grow the stack, because
- there can be only one construct selector set specified.
- This is enforced by omp_check_context_selector. */
- constructs
- = (enum tree_code *) alloca (nconstructs
- * sizeof (enum tree_code));
- omp_construct_traits_to_codes (selectors, nconstructs,
- constructs);
- }
+ gcc_assert (selectors);
- if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
+ /* During gimplification, CONSTRUCT_CONTEXT is partial, and doesn't
+ include a construct for "declare simd" that may be added
+ when there is not an enclosing "target" construct. We might
+ be able to find a positive match against the partial context
+ (although we cannot yet score it accurately), but if we can't,
+ treat it as unknown instead of no match. */
+ if (!omp_construct_traits_match (selectors, construct_context, NULL))
{
- if (!cfun->after_inlining)
- {
- ret = -1;
- continue;
- }
- int i;
- for (i = 0; i < nconstructs; ++i)
- if (constructs[i] == OMP_SIMD)
- break;
- if (i < nconstructs)
- {
- ret = -1;
- continue;
- }
- /* If there is no simd, assume it is ok after IPA,
- constructs should have been checked before. */
- continue;
- }
+ /* If we've got a complete context, it's definitely a failed
+ match. */
+ if (complete_p)
+ return 0;
- int r = omp_construct_selector_matches (constructs, nconstructs,
- NULL);
- if (r == 0)
- return 0;
- if (r == -1)
- ret = -1;
+ /* If the selector doesn't include simd, then we don't have
+ to worry about whether "declare simd" would cause it to
+ match; so this is also a definite failure. */
+ bool have_simd = false;
+ for (tree ts = construct_context; ts; ts = TREE_CHAIN (ts))
+ if (OMP_TS_CODE (ts) == OMP_TRAIT_CONSTRUCT_SIMD)
+ {
+ have_simd = true;
+ break;
+ }
+ if (!have_simd)
+ return 0;
+ else
+ ret = -1;
+ }
continue;
}
else if (set == OMP_TRAIT_SET_TARGET_DEVICE)
@@ -1570,302 +1667,293 @@ omp_context_selector_matches (tree ctx, bool metadirective_p, bool delay_p)
switch (sel)
{
case OMP_TRAIT_IMPLEMENTATION_VENDOR:
- if (set == OMP_TRAIT_SET_IMPLEMENTATION)
- for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
- {
- const char *prop = omp_context_name_list_prop (p);
- if (prop == NULL)
- return 0;
- if (!strcmp (prop, "gnu"))
- continue;
- return 0;
- }
- break;
- case OMP_TRAIT_IMPLEMENTATION_EXTENSION:
- if (set == OMP_TRAIT_SET_IMPLEMENTATION)
- /* We don't support any extensions right now. */
- return 0;
- break;
- case OMP_TRAIT_IMPLEMENTATION_ADMO:
- if (set == OMP_TRAIT_SET_IMPLEMENTATION)
+ gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
+ for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
{
- if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
- break;
-
- enum omp_memory_order omo
- = ((enum omp_memory_order)
- (omp_requires_mask
- & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER));
- if (omo == OMP_MEMORY_ORDER_UNSPECIFIED)
- {
- /* We don't know yet, until end of TU. */
- if (symtab->state == PARSING)
- {
- ret = -1;
- break;
- }
- else
- omo = OMP_MEMORY_ORDER_RELAXED;
- }
- tree p = OMP_TS_PROPERTIES (ts);
- const char *prop = IDENTIFIER_POINTER (OMP_TP_NAME (p));
- if (!strcmp (prop, "relaxed")
- && omo != OMP_MEMORY_ORDER_RELAXED)
- return 0;
- else if (!strcmp (prop, "seq_cst")
- && omo != OMP_MEMORY_ORDER_SEQ_CST)
- return 0;
- else if (!strcmp (prop, "acq_rel")
- && omo != OMP_MEMORY_ORDER_ACQ_REL)
- return 0;
- else if (!strcmp (prop, "acquire")
- && omo != OMP_MEMORY_ORDER_ACQUIRE)
- return 0;
- else if (!strcmp (prop, "release")
- && omo != OMP_MEMORY_ORDER_RELEASE)
+ const char *prop = omp_context_name_list_prop (p);
+ if (prop == NULL)
return 0;
+ if (!strcmp (prop, "gnu"))
+ continue;
+ return 0;
}
break;
- case OMP_TRAIT_DEVICE_ARCH:
- if (set == OMP_TRAIT_SET_DEVICE)
- for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
- {
- const char *arch = omp_context_name_list_prop (p);
- if (arch == NULL)
- return 0;
- if (metadirective_p && delay_p)
- return -1;
+ case OMP_TRAIT_IMPLEMENTATION_EXTENSION:
+ gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
+ /* We don't support any extensions right now. */
+ return 0;
+ break;
+ case OMP_TRAIT_IMPLEMENTATION_ADMO:
+ gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
+ if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
+ break;
- int r = 0;
- if (targetm.omp.device_kind_arch_isa != NULL)
- r = targetm.omp.device_kind_arch_isa (omp_device_arch,
- arch);
- if (r == 0 || (r == -1 && symtab->state != PARSING))
+ {
+ enum omp_memory_order omo
+ = ((enum omp_memory_order)
+ (omp_requires_mask
+ & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER));
+ if (omo == OMP_MEMORY_ORDER_UNSPECIFIED)
+ {
+ /* We don't know yet, until end of TU. */
+ if (symtab->state == PARSING)
{
- /* If we are or might be in a target region or
- declare target function, need to take into account
- also offloading values. */
- if (!omp_maybe_offloaded ())
- return 0;
- if (ENABLE_OFFLOADING)
- {
- const char *arches = omp_offload_device_arch;
- if (omp_offload_device_kind_arch_isa (arches,
- arch))
- {
- ret = -1;
- continue;
- }
- }
- return 0;
+ ret = -1;
+ break;
}
- else if (r == -1)
- ret = -1;
- /* If arch matches on the host, it still might not match
- in the offloading region. */
- else if (omp_maybe_offloaded ())
- ret = -1;
+ else
+ omo = OMP_MEMORY_ORDER_RELAXED;
}
+ tree p = OMP_TS_PROPERTIES (ts);
+ const char *prop = IDENTIFIER_POINTER (OMP_TP_NAME (p));
+ if (!strcmp (prop, "relaxed")
+ && omo != OMP_MEMORY_ORDER_RELAXED)
+ return 0;
+ else if (!strcmp (prop, "seq_cst")
+ && omo != OMP_MEMORY_ORDER_SEQ_CST)
+ return 0;
+ else if (!strcmp (prop, "acq_rel")
+ && omo != OMP_MEMORY_ORDER_ACQ_REL)
+ return 0;
+ else if (!strcmp (prop, "acquire")
+ && omo != OMP_MEMORY_ORDER_ACQUIRE)
+ return 0;
+ else if (!strcmp (prop, "release")
+ && omo != OMP_MEMORY_ORDER_RELEASE)
+ return 0;
+ }
+ break;
+ case OMP_TRAIT_DEVICE_ARCH:
+ gcc_assert (set == OMP_TRAIT_SET_DEVICE);
+ for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
+ {
+ const char *arch = omp_context_name_list_prop (p);
+ if (arch == NULL)
+ return 0;
+ int r = 0;
+ if (targetm.omp.device_kind_arch_isa != NULL)
+ r = targetm.omp.device_kind_arch_isa (omp_device_arch,
+ arch);
+ if (r == 0 || (r == -1 && symtab->state != PARSING))
+ {
+ /* If we are or might be in a target region or
+ declare target function, need to take into account
+ also offloading values.
+ Note that maybe_offloaded is always false in late
+ resolution; that's handled as native code (the
+ above case) in the offload compiler instead. */
+ if (!maybe_offloaded)
+ return 0;
+ if (ENABLE_OFFLOADING)
+ {
+ const char *arches = omp_offload_device_arch;
+ if (omp_offload_device_kind_arch_isa (arches, arch))
+ {
+ ret = -1;
+ continue;
+ }
+ }
+ return 0;
+ }
+ else if (r == -1)
+ ret = -1;
+ /* If arch matches on the host, it still might not match
+ in the offloading region. */
+ else if (maybe_offloaded)
+ ret = -1;
+ }
break;
case OMP_TRAIT_IMPLEMENTATION_UNIFIED_ADDRESS:
- if (set == OMP_TRAIT_SET_IMPLEMENTATION)
- {
- if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
- break;
+ gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
+ if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
+ break;
- if ((omp_requires_mask & OMP_REQUIRES_UNIFIED_ADDRESS) == 0)
- {
- if (symtab->state == PARSING)
- ret = -1;
- else
- return 0;
- }
+ if ((omp_requires_mask & OMP_REQUIRES_UNIFIED_ADDRESS) == 0)
+ {
+ if (symtab->state == PARSING)
+ ret = -1;
+ else
+ return 0;
}
break;
case OMP_TRAIT_IMPLEMENTATION_UNIFIED_SHARED_MEMORY:
- if (set == OMP_TRAIT_SET_IMPLEMENTATION)
- {
- if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
- break;
+ gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
+ if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
+ break;
- if ((omp_requires_mask
- & OMP_REQUIRES_UNIFIED_SHARED_MEMORY) == 0)
- {
- if (symtab->state == PARSING)
- ret = -1;
- else
- return 0;
- }
+ if ((omp_requires_mask
+ & OMP_REQUIRES_UNIFIED_SHARED_MEMORY) == 0)
+ {
+ if (symtab->state == PARSING)
+ ret = -1;
+ else
+ return 0;
}
break;
case OMP_TRAIT_IMPLEMENTATION_DYNAMIC_ALLOCATORS:
- if (set == OMP_TRAIT_SET_IMPLEMENTATION)
- {
- if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
- break;
+ gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
+ if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
+ break;
- if ((omp_requires_mask
- & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
- {
- if (symtab->state == PARSING)
- ret = -1;
- else
- return 0;
- }
+ if ((omp_requires_mask
+ & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
+ {
+ if (symtab->state == PARSING)
+ ret = -1;
+ else
+ return 0;
}
break;
case OMP_TRAIT_IMPLEMENTATION_REVERSE_OFFLOAD:
- if (set == OMP_TRAIT_SET_IMPLEMENTATION)
- {
- if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
- break;
+ gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
+ if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
+ break;
- if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
- {
- if (symtab->state == PARSING)
- ret = -1;
- else
- return 0;
- }
+ if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
+ {
+ if (symtab->state == PARSING)
+ ret = -1;
+ else
+ return 0;
}
break;
case OMP_TRAIT_DEVICE_KIND:
- if (set == OMP_TRAIT_SET_DEVICE)
- for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
- {
- const char *prop = omp_context_name_list_prop (p);
- if (prop == NULL)
- return 0;
- if (!strcmp (prop, "any"))
- continue;
- if (!strcmp (prop, "host"))
- {
+ gcc_assert (set == OMP_TRAIT_SET_DEVICE);
+ for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
+ {
+ const char *prop = omp_context_name_list_prop (p);
+ if (prop == NULL)
+ return 0;
+ if (!strcmp (prop, "any"))
+ continue;
+ if (!strcmp (prop, "host"))
+ {
#ifdef ACCEL_COMPILER
- return 0;
+ return 0;
#else
- if (omp_maybe_offloaded ())
- ret = -1;
- continue;
+ if (maybe_offloaded)
+ ret = -1;
+ continue;
#endif
- }
- if (!strcmp (prop, "nohost"))
- {
+ }
+ if (!strcmp (prop, "nohost"))
+ {
#ifndef ACCEL_COMPILER
- if (omp_maybe_offloaded ())
- ret = -1;
- else
- return 0;
-#endif
- continue;
- }
- if (metadirective_p && delay_p)
- return -1;
-
- int r = 0;
- if (targetm.omp.device_kind_arch_isa != NULL)
- r = targetm.omp.device_kind_arch_isa (omp_device_kind,
- prop);
- else
- r = strcmp (prop, "cpu") == 0;
- if (r == 0 || (r == -1 && symtab->state != PARSING))
- {
- /* If we are or might be in a target region or
- declare target function, need to take into account
- also offloading values. */
- if (!omp_maybe_offloaded ())
- return 0;
- if (ENABLE_OFFLOADING)
- {
- const char *kinds = omp_offload_device_kind;
- if (omp_offload_device_kind_arch_isa (kinds, prop))
- {
- ret = -1;
- continue;
- }
- }
+ if (maybe_offloaded)
+ ret = -1;
+ else
return 0;
- }
- else if (r == -1)
- ret = -1;
- /* If kind matches on the host, it still might not match
- in the offloading region. */
- else if (omp_maybe_offloaded ())
- ret = -1;
- }
+#endif
+ continue;
+ }
+
+ int r = 0;
+ if (targetm.omp.device_kind_arch_isa != NULL)
+ r = targetm.omp.device_kind_arch_isa (omp_device_kind,
+ prop);
+ else
+ r = strcmp (prop, "cpu") == 0;
+ if (r == 0 || (r == -1 && symtab->state != PARSING))
+ {
+ /* If we are or might be in a target region or
+ declare target function, need to take into account
+ also offloading values.
+ Note that maybe_offloaded is always false in late
+ resolution; that's handled as native code (the
+ above case) in the offload compiler instead. */
+ if (!maybe_offloaded)
+ return 0;
+ if (ENABLE_OFFLOADING)
+ {
+ const char *kinds = omp_offload_device_kind;
+ if (omp_offload_device_kind_arch_isa (kinds, prop))
+ {
+ ret = -1;
+ continue;
+ }
+ }
+ return 0;
+ }
+ else if (r == -1)
+ ret = -1;
+ /* If kind matches on the host, it still might not match
+ in the offloading region. */
+ else if (maybe_offloaded)
+ ret = -1;
+ }
break;
case OMP_TRAIT_DEVICE_ISA:
- if (set == OMP_TRAIT_SET_DEVICE)
- for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
- {
- const char *isa = omp_context_name_list_prop (p);
- if (isa == NULL)
- return 0;
- if (metadirective_p && delay_p)
- return -1;
-
- int r = 0;
- if (targetm.omp.device_kind_arch_isa != NULL)
- r = targetm.omp.device_kind_arch_isa (omp_device_isa,
- isa);
- if (r == 0 || (r == -1 && symtab->state != PARSING))
- {
- /* If isa is valid on the target, but not in the
- current function and current function has
- #pragma omp declare simd on it, some simd clones
- might have the isa added later on. */
- if (r == -1
- && targetm.simd_clone.compute_vecsize_and_simdlen
- && (cfun == NULL || !cfun->after_inlining))
- {
- tree attrs
- = DECL_ATTRIBUTES (current_function_decl);
- if (lookup_attribute ("omp declare simd", attrs))
- {
- ret = -1;
- continue;
- }
- }
- /* If we are or might be in a target region or
- declare target function, need to take into account
- also offloading values. */
- if (!omp_maybe_offloaded ())
- return 0;
- if (ENABLE_OFFLOADING)
- {
- const char *isas = omp_offload_device_isa;
- if (omp_offload_device_kind_arch_isa (isas, isa))
- {
- ret = -1;
- continue;
- }
- }
+ gcc_assert (set == OMP_TRAIT_SET_DEVICE);
+ for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
+ {
+ const char *isa = omp_context_name_list_prop (p);
+ if (isa == NULL)
+ return 0;
+ int r = 0;
+ if (targetm.omp.device_kind_arch_isa != NULL)
+ r = targetm.omp.device_kind_arch_isa (omp_device_isa,
+ isa);
+ if (r == 0 || (r == -1 && symtab->state != PARSING))
+ {
+ /* If isa is valid on the target, but not in the
+ current function and current function has
+ #pragma omp declare simd on it, some simd clones
+ might have the isa added later on. */
+ if (r == -1
+ && targetm.simd_clone.compute_vecsize_and_simdlen
+ && (cfun == NULL || !cfun->after_inlining))
+ {
+ tree attrs
+ = DECL_ATTRIBUTES (current_function_decl);
+ if (lookup_attribute ("omp declare simd", attrs))
+ {
+ ret = -1;
+ continue;
+ }
+ }
+ /* If we are or might be in a target region or
+ declare target function, need to take into account
+ also offloading values.
+ Note that maybe_offloaded is always false in late
+ resolution; that's handled as native code (the
+ above case) in the offload compiler instead. */
+ if (!maybe_offloaded)
return 0;
- }
- else if (r == -1)
- ret = -1;
- /* If isa matches on the host, it still might not match
- in the offloading region. */
- else if (omp_maybe_offloaded ())
- ret = -1;
- }
+ if (ENABLE_OFFLOADING)
+ {
+ const char *isas = omp_offload_device_isa;
+ if (omp_offload_device_kind_arch_isa (isas, isa))
+ {
+ ret = -1;
+ continue;
+ }
+ }
+ return 0;
+ }
+ else if (r == -1)
+ ret = -1;
+ /* If isa matches on the host, it still might not match
+ in the offloading region. */
+ else if (maybe_offloaded)
+ ret = -1;
+ }
break;
case OMP_TRAIT_USER_CONDITION:
- if (set == OMP_TRAIT_SET_USER)
- for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
- if (OMP_TP_NAME (p) == NULL_TREE)
- {
- /* OpenMP 5.1 allows non-constant conditions for
- metadirectives. */
- if (metadirective_p
- && !tree_fits_shwi_p (OMP_TP_VALUE (p)))
- break;
+ gcc_assert (set == OMP_TRAIT_SET_USER);
+ for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
+ if (OMP_TP_NAME (p) == NULL_TREE)
+ {
+ /* If the expression is not a constant, the selector
+ is dynamic. */
+ if (!tree_fits_shwi_p (OMP_TP_VALUE (p)))
+ break;
- if (integer_zerop (OMP_TP_VALUE (p)))
- return 0;
- if (integer_nonzerop (OMP_TP_VALUE (p)))
- break;
- ret = -1;
- }
+ if (integer_zerop (OMP_TP_VALUE (p)))
+ return 0;
+ if (integer_nonzerop (OMP_TP_VALUE (p)))
+ break;
+ ret = -1;
+ }
break;
default:
break;
@@ -1876,10 +1964,13 @@ omp_context_selector_matches (tree ctx, bool metadirective_p, bool delay_p)
}
/* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
- in omp_context_selector_set_compare. */
+ in omp_context_selector_set_compare. If MATCH_P is true, additionally
+ apply the special matching rules for the "simdlen" and "aligned" clauses
+ used to determine whether the selector CLAUSES1 is part of matches
+ the OpenMP context containing CLAUSES2. */
static int
-omp_construct_simd_compare (tree clauses1, tree clauses2)
+omp_construct_simd_compare (tree clauses1, tree clauses2, bool match_p)
{
if (clauses1 == NULL_TREE)
return clauses2 == NULL_TREE ? 0 : -1;
@@ -1896,6 +1987,7 @@ omp_construct_simd_compare (tree clauses1, tree clauses2)
: inbranch(false), notinbranch(false), simdlen(NULL_TREE) {}
} data[2];
unsigned int i;
+ tree e0, e1;
for (i = 0; i < 2; i++)
for (tree c = i ? clauses2 : clauses1; c; c = OMP_CLAUSE_CHAIN (c))
{
@@ -1934,10 +2026,23 @@ omp_construct_simd_compare (tree clauses1, tree clauses2)
r |= data[0].inbranch ? 2 : 1;
if (data[0].notinbranch != data[1].notinbranch)
r |= data[0].notinbranch ? 2 : 1;
- if (!simple_cst_equal (data[0].simdlen, data[1].simdlen))
+ e0 = data[0].simdlen;
+ e1 = data[1].simdlen;
+ if (!simple_cst_equal (e0, e1))
{
- if (data[0].simdlen && data[1].simdlen)
- return 2;
+ if (e0 && e1)
+ {
+ if (match_p && tree_fits_uhwi_p (e0) && tree_fits_uhwi_p (e1))
+ {
+ /* The two simdlen clauses match if m is a multiple of n. */
+ unsigned HOST_WIDE_INT n = tree_to_uhwi (e0);
+ unsigned HOST_WIDE_INT m = tree_to_uhwi (e1);
+ if (m % n != 0)
+ return 2;
+ }
+ else
+ return 2;
+ }
r |= data[0].simdlen ? 2 : 1;
}
if (data[0].data_sharing.length () < data[1].data_sharing.length ()
@@ -1978,9 +2083,22 @@ omp_construct_simd_compare (tree clauses1, tree clauses2)
}
if (c1 == NULL_TREE)
continue;
- if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1),
- OMP_CLAUSE_ALIGNED_ALIGNMENT (c2)))
- return 2;
+ e0 = OMP_CLAUSE_ALIGNED_ALIGNMENT (c1);
+ e1 = OMP_CLAUSE_ALIGNED_ALIGNMENT (c2);
+ if (!simple_cst_equal (e0, e1))
+ {
+ if (e0 && e1
+ && match_p && tree_fits_uhwi_p (e0) && tree_fits_uhwi_p (e1))
+ {
+ /* The two aligned clauses match if n is a multiple of m. */
+ unsigned HOST_WIDE_INT n = tree_to_uhwi (e0);
+ unsigned HOST_WIDE_INT m = tree_to_uhwi (e1);
+ if (n % m != 0)
+ return 2;
+ }
+ else
+ return 2;
+ }
}
switch (r)
{
@@ -2059,7 +2177,7 @@ omp_context_selector_props_compare (enum omp_tss_code set,
1 if CTX2 is a strict subset of CTX1, or
2 if neither context is a subset of another one. */
-int
+static int
omp_context_selector_set_compare (enum omp_tss_code set, tree ctx1, tree ctx2)
{
@@ -2096,7 +2214,8 @@ omp_context_selector_set_compare (enum omp_tss_code set, tree ctx1, tree ctx2)
int r = 0;
if (OMP_TS_CODE (ts1) == OMP_TRAIT_CONSTRUCT_SIMD)
r = omp_construct_simd_compare (OMP_TS_PROPERTIES (ts1),
- OMP_TS_PROPERTIES (ts2));
+ OMP_TS_PROPERTIES (ts2),
+ false);
if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
return 2;
if (ret == 0)
@@ -2296,11 +2415,9 @@ omp_dynamic_cond (tree ctx)
{
tree expr_list = OMP_TS_PROPERTIES (user);
- gcc_assert (OMP_TP_NAME (expr_list) == NULL_TREE);
-
/* The user condition is not dynamic if it is constant. */
- if (!tree_fits_shwi_p (TREE_VALUE (expr_list)))
- expr = TREE_VALUE (expr_list);
+ if (!tree_fits_shwi_p (OMP_TP_VALUE (expr_list)))
+ expr = OMP_TP_VALUE (expr_list);
}
tree target_device
@@ -2359,86 +2476,145 @@ omp_dynamic_cond (tree ctx)
return expr;
}
+/* Compute *SCORE for context selector CTX, which is already known to match.
+ CONSTRUCT_CONTEXT is the OpenMP construct context; if this is null or
+ incomplete (e.g., during parsing or gimplification) then it may not be
+ possible to compute the score accurately. In this case it does a best
+ guess based on the incomplete context and returns false; otherwise it
+ returns true.
-/* Compute *SCORE for context selector CTX. Return true if the score
- would be different depending on whether it is a declare simd clone or
- not. DECLARE_SIMD should be true for the case when it would be
- a declare simd clone. */
+ Cited text in the comments is from section 7.2 of the OpenMP 5.2
+ specification. */
static bool
-omp_context_compute_score (tree ctx, score_wide_int *score, bool declare_simd)
+omp_context_compute_score (tree ctx, tree construct_context,
+ bool complete_p, score_wide_int *score)
{
- tree selectors
- = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_CONSTRUCT);
- bool has_kind
- = (omp_get_context_selector (ctx, OMP_TRAIT_SET_DEVICE,
- OMP_TRAIT_DEVICE_KIND)
- || omp_get_context_selector (ctx, OMP_TRAIT_SET_TARGET_DEVICE,
- OMP_TRAIT_DEVICE_KIND));
- bool has_arch
- = (omp_get_context_selector (ctx, OMP_TRAIT_SET_DEVICE,
- OMP_TRAIT_DEVICE_ARCH)
- || omp_get_context_selector (ctx, OMP_TRAIT_SET_TARGET_DEVICE,
- OMP_TRAIT_DEVICE_ARCH));
- bool has_isa
- = (omp_get_context_selector (ctx, OMP_TRAIT_SET_DEVICE,
- OMP_TRAIT_DEVICE_ISA)
- || omp_get_context_selector (ctx, OMP_TRAIT_SET_TARGET_DEVICE,
- OMP_TRAIT_DEVICE_ISA));
- bool ret = false;
+ int l = list_length (construct_context);
+ bool retval = true;
+
+ /* "the final score is the sum of the values of all specified selectors
+ plus 1". */
*score = 1;
for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
- if (OMP_TSS_TRAIT_SELECTORS (tss) != selectors)
- for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts; ts = TREE_CHAIN (ts))
+ {
+ if (OMP_TSS_CODE (tss) == OMP_TRAIT_SET_CONSTRUCT)
{
- tree s = OMP_TS_SCORE (ts);
- if (s && TREE_CODE (s) == INTEGER_CST)
- *score += score_wide_int::from (wi::to_wide (s),
- TYPE_SIGN (TREE_TYPE (s)));
+ /* "Each trait selector for which the corresponding trait appears
+ in the context trait set in the OpenMP context..." */
+ score_wide_int tss_score;
+ omp_construct_traits_match (OMP_TSS_TRAIT_SELECTORS (tss),
+ construct_context, &tss_score);
+ *score += tss_score;
+ if (!complete_p)
+ retval = false;
+ }
+ else if (OMP_TSS_CODE (tss) == OMP_TRAIT_SET_DEVICE
+ || OMP_TSS_CODE (tss) == OMP_TRAIT_SET_TARGET_DEVICE)
+ {
+ /* "The kind, arch, and isa selectors, if specified, are given
+ the values 2**l, 2**(l+1), and 2**(l+2), respectively..."
+ FIXME: the spec isn't clear what should happen if there are
+ both "device" and "target_device" selector sets specified.
+ This implementation adds up the bits rather than ORs them. */
+ for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts;
+ ts = TREE_CHAIN (ts))
+ {
+ enum omp_ts_code code = OMP_TS_CODE (ts);
+ if (code == OMP_TRAIT_DEVICE_KIND)
+ *score += wi::shifted_mask <score_wide_int> (l, 1, false);
+ else if (code == OMP_TRAIT_DEVICE_ARCH)
+ *score += wi::shifted_mask <score_wide_int> (l + 1, 1, false);
+ else if (code == OMP_TRAIT_DEVICE_ISA)
+ *score += wi::shifted_mask <score_wide_int> (l + 2, 1, false);
+ }
+ if (!complete_p)
+ retval = false;
+ }
+ else
+ {
+ /* "Trait selectors for which a trait-score is specified..."
+ Note that there are no implementation-defined selectors, and
+ "other selectors are given a value of zero". */
+ for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts;
+ ts = TREE_CHAIN (ts))
+ {
+ tree s = OMP_TS_SCORE (ts);
+ if (s && TREE_CODE (s) == INTEGER_CST)
+ *score += score_wide_int::from (wi::to_wide (s),
+ TYPE_SIGN (TREE_TYPE (s)));
+ }
+ }
+ }
+ return retval;
+}
+
+/* CONSTRUCT_CONTEXT contains "the directive names, each being a trait,
+ of all enclosing constructs at that point in the program up to a target
+ construct", per section 7.1 of the 5.2 specification. The traits are
+ collected during gimplification and are listed outermost first.
+
+ This function attempts to apply the "if the point in the program is not
+ enclosed by a target construct, the following rules are applied in order"
+ requirements that follow in the same paragraph. This may not be possible,
+ depending on the compilation phase; in particular, "declare simd" clones
+ are not known until late resolution.
+
+ The augmented context is returned, and *COMPLETEP is set to true if
+ the context is known to be complete, false otherwise. */
+static tree
+omp_complete_construct_context (tree construct_context, bool *completep)
+{
+ /* The point in the program is enclosed by a target construct. */
+ if (construct_context
+ && OMP_TS_CODE (construct_context) == OMP_TRAIT_CONSTRUCT_TARGET)
+ *completep = true;
+
+ /* At parse time we have none of the information we need to collect
+ the missing pieces. */
+ else if (symtab->state == PARSING)
+ *completep = false;
+
+ else
+ {
+ tree attributes = DECL_ATTRIBUTES (current_function_decl);
+
+ /* Add simd trait when in a simd clone. This information is only
+ available during late resolution in the omp_device_lower pass,
+ however we can also rule out cases where we know earlier that
+ cfun is not a candidate for cloning. */
+ if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
+ {
+ cgraph_node *node = cgraph_node::get (cfun->decl);
+ if (node->simdclone)
+ construct_context = make_trait_selector (OMP_TRAIT_CONSTRUCT_SIMD,
+ NULL_TREE, NULL_TREE,
+ construct_context);
+ *completep = true;
+ }
+ else if (lookup_attribute ("omp declare simd", attributes))
+ *completep = false;
+ else
+ *completep = true;
+
+ /* Add construct selector set within a "declare variant" function. */
+ tree variant_attr
+ = lookup_attribute ("omp declare variant variant", attributes);
+ if (variant_attr)
+ {
+ tree temp = NULL_TREE;
+ for (tree t = TREE_VALUE (variant_attr); t; t = TREE_CHAIN (t))
+ temp = chainon (temp, copy_node (t));
+ construct_context = chainon (temp, construct_context);
}
- if (selectors || has_kind || has_arch || has_isa)
- {
- int nconstructs = list_length (selectors);
- enum tree_code *constructs = NULL;
- if (nconstructs)
- {
- constructs
- = (enum tree_code *) alloca (nconstructs
- * sizeof (enum tree_code));
- omp_construct_traits_to_codes (selectors, nconstructs, constructs);
- }
- int *scores
- = (int *) alloca ((2 * nconstructs + 2) * sizeof (int));
- if (omp_construct_selector_matches (constructs, nconstructs, scores)
- == 2)
- ret = true;
- int b = declare_simd ? nconstructs + 1 : 0;
- if (scores[b + nconstructs] + 4U < score->get_precision ())
- {
- for (int n = 0; n < nconstructs; ++n)
- {
- if (scores[b + n] < 0)
- {
- *score = -1;
- return ret;
- }
- *score += wi::shifted_mask <score_wide_int> (scores[b + n], 1, false);
- }
- if (has_kind)
- *score += wi::shifted_mask <score_wide_int> (scores[b + nconstructs],
- 1, false);
- if (has_arch)
- *score += wi::shifted_mask <score_wide_int> (scores[b + nconstructs] + 1,
- 1, false);
- if (has_isa)
- *score += wi::shifted_mask <score_wide_int> (scores[b + nconstructs] + 2,
- 1, false);
- }
- else /* FIXME: Implement this. */
- gcc_unreachable ();
+ /* Add target trait when in a target variant. */
+ if (lookup_attribute ("omp declare target block", attributes))
+ construct_context = make_trait_selector (OMP_TRAIT_CONSTRUCT_TARGET,
+ NULL_TREE, NULL_TREE,
+ construct_context);
}
- return ret;
+ return construct_context;
}
/* Class describing a single variant. */
@@ -2537,388 +2713,6 @@ omp_declare_variant_alt_hasher::equal (omp_declare_variant_base_entry *x,
static GTY(()) hash_table<omp_declare_variant_alt_hasher>
*omp_declare_variant_alt;
-/* Try to resolve declare variant after gimplification. */
-
-static tree
-omp_resolve_late_declare_variant (tree alt)
-{
- cgraph_node *node = cgraph_node::get (alt);
- cgraph_node *cur_node = cgraph_node::get (cfun->decl);
- if (node == NULL
- || !node->declare_variant_alt
- || !cfun->after_inlining)
- return alt;
-
- omp_declare_variant_base_entry entry;
- entry.base = NULL;
- entry.node = node;
- entry.variants = NULL;
- omp_declare_variant_base_entry *entryp
- = omp_declare_variant_alt->find_with_hash (&entry, DECL_UID (alt));
-
- unsigned int i, j;
- omp_declare_variant_entry *varentry1, *varentry2;
- auto_vec <bool, 16> matches;
- unsigned int nmatches = 0;
- FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
- {
- if (varentry1->matches)
- {
- /* This has been checked to be ok already. */
- matches.safe_push (true);
- nmatches++;
- continue;
- }
- switch (omp_context_selector_matches (varentry1->ctx, false, true))
- {
- case 0:
- matches.safe_push (false);
- break;
- case -1:
- return alt;
- default:
- matches.safe_push (true);
- nmatches++;
- break;
- }
- }
-
- if (nmatches == 0)
- return entryp->base->decl;
-
- /* A context selector that is a strict subset of another context selector
- has a score of zero. */
- FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
- if (matches[i])
- {
- for (j = i + 1;
- vec_safe_iterate (entryp->variants, j, &varentry2); ++j)
- if (matches[j])
- {
- int r = omp_context_selector_compare (varentry1->ctx,
- varentry2->ctx);
- if (r == -1)
- {
- /* ctx1 is a strict subset of ctx2, ignore ctx1. */
- matches[i] = false;
- break;
- }
- else if (r == 1)
- /* ctx2 is a strict subset of ctx1, remove ctx2. */
- matches[j] = false;
- }
- }
-
- score_wide_int max_score = -1;
- varentry2 = NULL;
- FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
- if (matches[i])
- {
- score_wide_int score
- = (cur_node->simdclone ? varentry1->score_in_declare_simd_clone
- : varentry1->score);
- if (score > max_score)
- {
- max_score = score;
- varentry2 = varentry1;
- }
- }
- return varentry2->variant->decl;
-}
-
-/* Hook to adjust hash tables on cgraph_node removal. */
-
-static void
-omp_declare_variant_remove_hook (struct cgraph_node *node, void *)
-{
- if (!node->declare_variant_alt)
- return;
-
- /* Drop this hash table completely. */
- omp_declare_variants = NULL;
- /* And remove node from the other hash table. */
- if (omp_declare_variant_alt)
- {
- omp_declare_variant_base_entry entry;
- entry.base = NULL;
- entry.node = node;
- entry.variants = NULL;
- omp_declare_variant_alt->remove_elt_with_hash (&entry,
- DECL_UID (node->decl));
- }
-}
-
-/* Try to resolve declare variant, return the variant decl if it should
- be used instead of base, or base otherwise. */
-
-tree
-omp_resolve_declare_variant (tree base)
-{
- tree variant1 = NULL_TREE, variant2 = NULL_TREE;
- if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
- return omp_resolve_late_declare_variant (base);
-
- auto_vec <tree, 16> variants;
- auto_vec <bool, 16> defer;
- bool any_deferred = false;
- for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
- {
- attr = lookup_attribute ("omp declare variant base", attr);
- if (attr == NULL_TREE)
- break;
- if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr))) != FUNCTION_DECL)
- continue;
- cgraph_node *node = cgraph_node::get (base);
- /* If this is already a magic decl created by this function,
- don't process it again. */
- if (node && node->declare_variant_alt)
- return base;
- switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr)),
- false, true))
- {
- case 0:
- /* No match, ignore. */
- break;
- case -1:
- /* Needs to be deferred. */
- any_deferred = true;
- variants.safe_push (attr);
- defer.safe_push (true);
- break;
- default:
- variants.safe_push (attr);
- defer.safe_push (false);
- break;
- }
- }
- if (variants.length () == 0)
- return base;
-
- if (any_deferred)
- {
- score_wide_int max_score1 = 0;
- score_wide_int max_score2 = 0;
- bool first = true;
- unsigned int i;
- tree attr1, attr2;
- omp_declare_variant_base_entry entry;
- entry.base = cgraph_node::get_create (base);
- entry.node = NULL;
- vec_alloc (entry.variants, variants.length ());
- FOR_EACH_VEC_ELT (variants, i, attr1)
- {
- score_wide_int score1;
- score_wide_int score2;
- bool need_two;
- tree ctx = TREE_VALUE (TREE_VALUE (attr1));
- need_two = omp_context_compute_score (ctx, &score1, false);
- if (need_two)
- omp_context_compute_score (ctx, &score2, true);
- else
- score2 = score1;
- if (first)
- {
- first = false;
- max_score1 = score1;
- max_score2 = score2;
- if (!defer[i])
- {
- variant1 = attr1;
- variant2 = attr1;
- }
- }
- else
- {
- if (max_score1 == score1)
- variant1 = NULL_TREE;
- else if (score1 > max_score1)
- {
- max_score1 = score1;
- variant1 = defer[i] ? NULL_TREE : attr1;
- }
- if (max_score2 == score2)
- variant2 = NULL_TREE;
- else if (score2 > max_score2)
- {
- max_score2 = score2;
- variant2 = defer[i] ? NULL_TREE : attr1;
- }
- }
- omp_declare_variant_entry varentry;
- varentry.variant
- = cgraph_node::get_create (TREE_PURPOSE (TREE_VALUE (attr1)));
- varentry.score = score1;
- varentry.score_in_declare_simd_clone = score2;
- varentry.ctx = ctx;
- varentry.matches = !defer[i];
- entry.variants->quick_push (varentry);
- }
-
- /* If there is a clear winner variant with the score which is not
- deferred, verify it is not a strict subset of any other context
- selector and if it is not, it is the best alternative no matter
- whether the others do or don't match. */
- if (variant1 && variant1 == variant2)
- {
- tree ctx1 = TREE_VALUE (TREE_VALUE (variant1));
- FOR_EACH_VEC_ELT (variants, i, attr2)
- {
- if (attr2 == variant1)
- continue;
- tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
- int r = omp_context_selector_compare (ctx1, ctx2);
- if (r == -1)
- {
- /* The winner is a strict subset of ctx2, can't
- decide now. */
- variant1 = NULL_TREE;
- break;
- }
- }
- if (variant1)
- {
- vec_free (entry.variants);
- return TREE_PURPOSE (TREE_VALUE (variant1));
- }
- }
-
- static struct cgraph_node_hook_list *node_removal_hook_holder;
- if (!node_removal_hook_holder)
- node_removal_hook_holder
- = symtab->add_cgraph_removal_hook (omp_declare_variant_remove_hook,
- NULL);
-
- if (omp_declare_variants == NULL)
- omp_declare_variants
- = hash_table<omp_declare_variant_hasher>::create_ggc (64);
- omp_declare_variant_base_entry **slot
- = omp_declare_variants->find_slot (&entry, INSERT);
- if (*slot != NULL)
- {
- vec_free (entry.variants);
- return (*slot)->node->decl;
- }
-
- *slot = ggc_cleared_alloc<omp_declare_variant_base_entry> ();
- (*slot)->base = entry.base;
- (*slot)->node = entry.base;
- (*slot)->variants = entry.variants;
- tree alt = build_decl (DECL_SOURCE_LOCATION (base), FUNCTION_DECL,
- DECL_NAME (base), TREE_TYPE (base));
- DECL_ARTIFICIAL (alt) = 1;
- DECL_IGNORED_P (alt) = 1;
- TREE_STATIC (alt) = 1;
- tree attributes = DECL_ATTRIBUTES (base);
- if (lookup_attribute ("noipa", attributes) == NULL)
- {
- attributes = tree_cons (get_identifier ("noipa"), NULL, attributes);
- if (lookup_attribute ("noinline", attributes) == NULL)
- attributes = tree_cons (get_identifier ("noinline"), NULL,
- attributes);
- if (lookup_attribute ("noclone", attributes) == NULL)
- attributes = tree_cons (get_identifier ("noclone"), NULL,
- attributes);
- if (lookup_attribute ("no_icf", attributes) == NULL)
- attributes = tree_cons (get_identifier ("no_icf"), NULL,
- attributes);
- }
- DECL_ATTRIBUTES (alt) = attributes;
- DECL_INITIAL (alt) = error_mark_node;
- (*slot)->node = cgraph_node::create (alt);
- (*slot)->node->declare_variant_alt = 1;
- (*slot)->node->create_reference (entry.base, IPA_REF_ADDR);
- omp_declare_variant_entry *varentry;
- FOR_EACH_VEC_SAFE_ELT (entry.variants, i, varentry)
- (*slot)->node->create_reference (varentry->variant, IPA_REF_ADDR);
- if (omp_declare_variant_alt == NULL)
- omp_declare_variant_alt
- = hash_table<omp_declare_variant_alt_hasher>::create_ggc (64);
- *omp_declare_variant_alt->find_slot_with_hash (*slot, DECL_UID (alt),
- INSERT) = *slot;
- return alt;
- }
-
- if (variants.length () == 1)
- return TREE_PURPOSE (TREE_VALUE (variants[0]));
-
- /* A context selector that is a strict subset of another context selector
- has a score of zero. */
- tree attr1, attr2;
- unsigned int i, j;
- FOR_EACH_VEC_ELT (variants, i, attr1)
- if (attr1)
- {
- tree ctx1 = TREE_VALUE (TREE_VALUE (attr1));
- FOR_EACH_VEC_ELT_FROM (variants, j, attr2, i + 1)
- if (attr2)
- {
- tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
- int r = omp_context_selector_compare (ctx1, ctx2);
- if (r == -1)
- {
- /* ctx1 is a strict subset of ctx2, remove
- attr1 from the vector. */
- variants[i] = NULL_TREE;
- break;
- }
- else if (r == 1)
- /* ctx2 is a strict subset of ctx1, remove attr2
- from the vector. */
- variants[j] = NULL_TREE;
- }
- }
- score_wide_int max_score1 = 0;
- score_wide_int max_score2 = 0;
- bool first = true;
- FOR_EACH_VEC_ELT (variants, i, attr1)
- if (attr1)
- {
- if (variant1)
- {
- score_wide_int score1;
- score_wide_int score2;
- bool need_two;
- tree ctx;
- if (first)
- {
- first = false;
- ctx = TREE_VALUE (TREE_VALUE (variant1));
- need_two = omp_context_compute_score (ctx, &max_score1, false);
- if (need_two)
- omp_context_compute_score (ctx, &max_score2, true);
- else
- max_score2 = max_score1;
- }
- ctx = TREE_VALUE (TREE_VALUE (attr1));
- need_two = omp_context_compute_score (ctx, &score1, false);
- if (need_two)
- omp_context_compute_score (ctx, &score2, true);
- else
- score2 = score1;
- if (score1 > max_score1)
- {
- max_score1 = score1;
- variant1 = attr1;
- }
- if (score2 > max_score2)
- {
- max_score2 = score2;
- variant2 = attr1;
- }
- }
- else
- {
- variant1 = attr1;
- variant2 = attr1;
- }
- }
- /* If there is a disagreement on which variant has the highest score
- depending on whether it will be in a declare simd clone or not,
- punt for now and defer until after IPA where we will know that. */
- return ((variant1 && variant1 == variant2)
- ? TREE_PURPOSE (TREE_VALUE (variant1)) : base);
-}
-
void
omp_lto_output_declare_variant_alt (lto_simple_output_block *ob,
cgraph_node *node,
@@ -3059,16 +2853,36 @@ sort_variant (const void * a, const void *b, void *)
}
/* Return a vector of dynamic replacement candidates for the directive
- candidates in ALL_VARIANTS. Return an empty vector if the metadirective
+ candidates in ALL_VARIANTS. Return an empty vector if the candidates
cannot be resolved. */
-static vec<struct omp_variant>
+vec<struct omp_variant>
omp_get_dynamic_candidates (vec <struct omp_variant> &all_variants,
- bool delay_p)
+ tree construct_context)
{
auto_vec <struct omp_variant> variants;
struct omp_variant default_variant;
bool default_found = false;
+ bool complete_p;
+
+ construct_context
+ = omp_complete_construct_context (construct_context, &complete_p);
+
+ if (dump_file)
+ {
+ fprintf (dump_file, "\nIn omp_get_dynamic_candidates:\n");
+ if (symtab->state == PARSING)
+ fprintf (dump_file, "invoked during parsing\n");
+ else if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0)
+ fprintf (dump_file, "invoked during gimplification\n");
+ else if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
+ fprintf (dump_file, "invoked during late resolution\n");
+ else
+ fprintf (dump_file, "confused about invocation context?!?\n");
+ fprintf (dump_file, "construct_context has %d traits (%s)\n",
+ (construct_context ? list_length (construct_context) : 0),
+ (complete_p ? "complete" : "incomplete"));
+ }
for (unsigned int i = 0; i < all_variants.length (); i++)
{
@@ -3097,19 +2911,24 @@ omp_get_dynamic_candidates (vec <struct omp_variant> &all_variants,
fprintf (dump_file, " as candidate - ");
}
- switch (omp_context_selector_matches (variant.selector, true, delay_p))
+ switch (omp_context_selector_matches (variant.selector,
+ construct_context, complete_p))
{
case -1:
- variant.resolvable_p = false;
+ /* Give up for now. This can only happen during early resolution,
+ prior to or during gimplification. */
if (dump_file)
- fprintf (dump_file, "unresolvable");
- /* FALLTHRU */
+ fprintf (dump_file, "unresolvable\n");
+ gcc_assert (!cfun || (cfun->curr_properties & PROP_gimple_any) == 0);
+ variants.truncate (0);
+ return variants.copy ();
case 1:
- /* TODO: Handle SIMD score?. */
- omp_context_compute_score (variant.selector, &variant.score, false);
+ variant.resolvable_p
+ = omp_context_compute_score (variant.selector, construct_context,
+ complete_p, &variant.score);
variant.dynamic_selector = omp_dynamic_cond (variant.selector);
variants.safe_push (variant);
- if (dump_file && variant.resolvable_p)
+ if (dump_file)
{
if (variant.dynamic_selector)
fprintf (dump_file, "matched, dynamic");
@@ -3130,6 +2949,21 @@ omp_get_dynamic_candidates (vec <struct omp_variant> &all_variants,
/* There must be one default variant. */
gcc_assert (default_found);
+ /* If there are no matching selectors, return the default. */
+ if (variants.length () == 0)
+ {
+ variants.safe_push (default_variant);
+ return variants.copy ();
+ }
+
+ /* If there is only one matching selector, use it. */
+ if (variants.length () == 1)
+ {
+ if (variants[0].dynamic_selector)
+ variants.safe_push (default_variant);
+ return variants.copy ();
+ }
+
/* A context selector that is a strict subset of another context selector
has a score of zero. */
for (unsigned int i = 0; i < variants.length (); i++)
@@ -3155,14 +2989,80 @@ omp_get_dynamic_candidates (vec <struct omp_variant> &all_variants,
/* Add the default as a final choice. */
variants.safe_push (default_variant);
+ if (dump_file)
+ {
+ fprintf (dump_file, "Sorted variants are:\n");
+ for (unsigned i = 0; i < variants.length (); i++)
+ {
+ HOST_WIDE_INT score = variants[i].score.to_shwi ();
+ fprintf (dump_file, "score %d ", (int)score);
+ if (variants[i].selector)
+ {
+ fprintf (dump_file, "selector ");
+ print_omp_context_selector (dump_file, variants[i].selector,
+ TDF_NONE);
+ fprintf (dump_file, "\n");
+ }
+ else
+ fprintf (dump_file, "default selector\n");
+ }
+ }
+
/* Build the dynamic candidate list. */
for (unsigned i = 0; i < variants.length (); i++)
{
- /* If one of the candidates is unresolvable, give up for now. */
+ /* In general, we can't proceed if we can't accurately score any
+ of the selectors, since the sorting may be incorrect. But, since
+ the actual score will never be lower than the guessed value, we
+ can use the variant if it is the first one and either the next
+ one is resolvable or we can make a direct comparison of the
+ isa/arch/kind bits. */
if (!variants[i].resolvable_p)
{
- variants.truncate (0);
- break;
+ bool ok = true;
+ if (i != 0)
+ ok = false;
+ else if (!variants[i+1].resolvable_p)
+ {
+ /* To keep comparisons simple, reject selectors that contain
+ sets other than device, target_device, or construct. */
+ for (tree tss = variants[i].selector;
+ tss && ok; tss = TREE_CHAIN (tss))
+ {
+ enum omp_tss_code code = OMP_TSS_CODE (tss);
+ if (code != OMP_TRAIT_SET_DEVICE
+ && code != OMP_TRAIT_SET_TARGET_DEVICE
+ && code != OMP_TRAIT_SET_CONSTRUCT)
+ ok = false;
+ }
+ for (tree tss = variants[i+1].selector;
+ tss && ok; tss = TREE_CHAIN (tss))
+ {
+ enum omp_tss_code code = OMP_TSS_CODE (tss);
+ if (code != OMP_TRAIT_SET_DEVICE
+ && code != OMP_TRAIT_SET_TARGET_DEVICE
+ && code != OMP_TRAIT_SET_CONSTRUCT)
+ ok = false;
+ }
+ /* Ignore the construct bits of the score. If the isa/arch/kind
+ bits are strictly ordered, we're good to go. Since
+ "the final score is the sum of the values of all specified
+ selectors plus 1", subtract that 1 from both scores before
+ getting rid of the low bits. */
+ if (ok)
+ {
+ size_t l = list_length (construct_context);
+ if ((variants[i].score - 1) >> l
+ <= (variants[i+1].score - 1) >> l)
+ ok = false;
+ }
+ }
+
+ if (!ok)
+ {
+ variants.truncate (0);
+ break;
+ }
}
if (dump_file)
@@ -3192,35 +3092,106 @@ omp_get_dynamic_candidates (vec <struct omp_variant> &all_variants,
return variants.copy ();
}
+/* Two attempts are made to resolve calls to "declare variant" functions:
+ early resolution in the gimplifier, and late resolution in the
+ omp_device_lower pass. If early resolution is not possible, the
+ original function call is gimplified into the same form as metadirective
+ and goes through the same late resolution code as metadirective. */
+
+/* Collect "declare variant" candidates for BASE. CONSTRUCT_CONTEXT
+ is the un-augmented context, or NULL_TREE if that information is not
+ available yet. */
+vec<struct omp_variant>
+omp_declare_variant_candidates (tree base, tree construct_context)
+{
+ auto_vec <struct omp_variant> candidates;
+ bool complete_p;
+ tree augmented_context
+ = omp_complete_construct_context (construct_context, &complete_p);
+
+ /* The variants are stored on (possible multiple) "omp declare variant base"
+ attributes on the base function. */
+ for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
+ {
+ attr = lookup_attribute ("omp declare variant base", attr);
+ if (attr == NULL_TREE)
+ break;
+
+ tree fndecl = TREE_PURPOSE (TREE_VALUE (attr));
+ tree selector = TREE_VALUE (TREE_VALUE (attr));
+
+ if (TREE_CODE (fndecl) != FUNCTION_DECL)
+ continue;
+
+ /* Ignore this variant if its selector is known not to match. */
+ if (!omp_context_selector_matches (selector, augmented_context,
+ complete_p))
+ continue;
+
+ struct omp_variant candidate;
+ candidate.selector = selector;
+ candidate.dynamic_selector = NULL_TREE;
+ candidate.alternative = fndecl;
+ candidates.safe_push (candidate);
+ }
+
+ /* Add a default that is the base function. */
+ struct omp_variant v;
+ v.selector = NULL_TREE;
+ v.dynamic_selector = NULL_TREE;
+ v.alternative = base;
+ candidates.safe_push (v);
+ return candidates.copy ();
+}
+
+/* Collect metadirective candidates for METADIRECTIVE. CONSTRUCT_CONTEXT
+ is the un-augmented context, or NULL_TREE if that information is not
+ available yet. */
+vec<struct omp_variant>
+omp_metadirective_candidates (tree metadirective, tree construct_context)
+{
+ auto_vec <struct omp_variant> candidates;
+ tree variant = OMP_METADIRECTIVE_VARIANTS (metadirective);
+ bool complete_p;
+ tree augmented_context
+ = omp_complete_construct_context (construct_context, &complete_p);
+
+ gcc_assert (variant);
+ for (; variant; variant = TREE_CHAIN (variant))
+ {
+ tree selector = OMP_METADIRECTIVE_VARIANT_SELECTOR (variant);
+
+ /* Ignore this variant if its selector is known not to match. */
+ if (!omp_context_selector_matches (selector, augmented_context,
+ complete_p))
+ continue;
+
+ struct omp_variant candidate;
+ candidate.selector = selector;
+ candidate.dynamic_selector = NULL_TREE;
+ candidate.alternative = OMP_METADIRECTIVE_VARIANT_DIRECTIVE (variant);
+ candidate.body = OMP_METADIRECTIVE_VARIANT_BODY (variant);
+ candidates.safe_push (candidate);
+ }
+ return candidates.copy ();
+}
+
/* Return a vector of dynamic replacement candidates for the metadirective
statement in METADIRECTIVE. Return an empty vector if the metadirective
- cannot be resolved. */
+ cannot be resolved. This function is intended to be called from the
+ front ends, prior to gimplification. */
vec<struct omp_variant>
omp_early_resolve_metadirective (tree metadirective)
{
- auto_vec <struct omp_variant> candidates;
- tree variant = OMP_METADIRECTIVE_VARIANTS (metadirective);
-
- gcc_assert (variant);
- while (variant)
- {
- struct omp_variant candidate;
-
- candidate.selector = OMP_METADIRECTIVE_VARIANT_SELECTOR (variant);
- candidate.alternative = OMP_METADIRECTIVE_VARIANT_DIRECTIVE (variant);
- candidate.body = OMP_METADIRECTIVE_VARIANT_BODY (variant);
-
- candidates.safe_push (candidate);
- variant = TREE_CHAIN (variant);
- }
-
- return omp_get_dynamic_candidates (candidates, true);
+ vec <struct omp_variant> candidates
+ = omp_metadirective_candidates (metadirective, NULL_TREE);
+ return omp_get_dynamic_candidates (candidates, NULL_TREE);
}
/* Return a vector of dynamic replacement candidates for the metadirective
- Gimple statement in GS. Return an empty vector if the metadirective
- cannot be resolved. */
+ Gimple statement in GS. This version is called during late resolution
+ in the ompdevlow pass. */
vec<struct omp_variant>
omp_late_resolve_metadirective (gimple *gs)
@@ -3232,12 +3203,14 @@ omp_late_resolve_metadirective (gimple *gs)
struct omp_variant variant;
variant.selector = gimple_op (gs, i);
+ variant.dynamic_selector = NULL_TREE;
variant.alternative = gimple_omp_metadirective_label (gs, i);
variants.safe_push (variant);
}
- return omp_get_dynamic_candidates (variants, false);
+ tree construct_context = gimple_omp_metadirective_context (gs);
+ return omp_get_dynamic_candidates (variants, construct_context);
}
/* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
@@ -95,16 +95,27 @@ struct omp_for_data
desirable to be the same on all targets. */
typedef generic_wide_int <fixed_wide_int_storage <1024> > score_wide_int;
-/* A structure describing a variant in a metadirective. */
-
+/* A structure describing a variant alternative in a metadirective or
+ variant function, used for matching and scoring during resolution. */
struct GTY(()) omp_variant
{
- score_wide_int score;
+ /* Context selector. This is NULL_TREE for the default. */
tree selector;
+ /* For early resolution of "metadirective", contains the nested directive.
+ For early resolution of "declare variant", contains the function decl
+ for this alternative. For late resolution of both, contains the label
+ that is the branch target for this alternative. */
tree alternative;
+ /* Common body, used for metadirective. */
tree body;
+ /* If the selector is dynamic, this is the dynamic part; otherwise
+ NULL_TREE. Filled in during resolution. */
tree dynamic_selector;
+ /* A selector can match but not be resolvable due to its score not
+ being computable yet. */
bool resolvable_p : 1;
+ /* The score, if resolvable_p is true. */
+ score_wide_int score;
};
#define OACC_FN_ATTRIB "oacc function"
@@ -190,17 +201,18 @@ extern tree find_combined_omp_for (tree *, int *, void *);
extern poly_uint64 omp_max_vf (void);
extern int omp_max_simt_vf (void);
extern const char *omp_context_name_list_prop (tree);
-extern void omp_construct_traits_to_codes (tree, int, enum tree_code *);
extern tree omp_check_context_selector (location_t loc, tree ctx,
bool metadirective_p);
extern void omp_mark_declare_variant (location_t loc, tree variant,
tree construct);
-extern int omp_context_selector_matches (tree, bool, bool);
-extern int omp_context_selector_set_compare (enum omp_tss_code, tree, tree);
+extern int omp_context_selector_matches (tree, tree, bool);
extern tree omp_get_context_selector (tree, enum omp_tss_code,
enum omp_ts_code);
extern tree omp_get_context_selector_list (tree, enum omp_tss_code);
-extern tree omp_resolve_declare_variant (tree);
+extern vec<struct omp_variant> omp_declare_variant_candidates (tree, tree);
+extern vec<struct omp_variant> omp_metadirective_candidates (tree, tree);
+extern vec<struct omp_variant>
+omp_get_dynamic_candidates (vec<struct omp_variant>&, tree);
extern vec<struct omp_variant> omp_early_resolve_metadirective (tree);
extern vec<struct omp_variant> omp_late_resolve_metadirective (gimple *);
extern tree oacc_launch_pack (unsigned code, tree device, unsigned op);
@@ -2736,16 +2736,8 @@ execute_omp_device_lower ()
continue;
if (!gimple_call_internal_p (stmt))
{
- if (calls_declare_variant_alt)
- if (tree fndecl = gimple_call_fndecl (stmt))
- {
- tree new_fndecl = omp_resolve_declare_variant (fndecl);
- if (new_fndecl != fndecl)
- {
- gimple_call_set_fndecl (stmt, new_fndecl);
- update_stmt (stmt);
- }
- }
+ /* FIXME: this is a leftover of obsolete code. */
+ gcc_assert (!calls_declare_variant_alt);
#ifdef ACCEL_COMPILER
if (omp_redirect_indirect_calls
&& gimple_call_fndecl (stmt) == NULL_TREE)
@@ -29,29 +29,29 @@ void f13 (void);
void f14 (void);
void f15 (void);
void f16 (void);
-#pragma omp declare variant (f14) match (construct={teams,parallel,for}) /* 16+8+4 */
-#pragma omp declare variant (f15) match (construct={parallel},user={condition(score(19):1)}) /* 8+19 */
-#pragma omp declare variant (f16) match (implementation={atomic_default_mem_order(score(27):seq_cst)})
+#pragma omp declare variant (f14) match (construct={teams,parallel,for}) /* 1+8+16 */
+#pragma omp declare variant (f15) match (construct={parallel},user={condition(score(16):1)}) /* 8+16 */
+#pragma omp declare variant (f16) match (implementation={atomic_default_mem_order(score(24):seq_cst)})
void f17 (void);
void f18 (void);
void f19 (void);
void f20 (void);
-#pragma omp declare variant (f18) match (construct={teams,parallel,for}) /* 16+8+4 */
+#pragma omp declare variant (f18) match (construct={teams,parallel,for}) /* 1+8+6 */
#pragma omp declare variant (f19) match (construct={for},user={condition(score(25):1)}) /* 4+25 */
#pragma omp declare variant (f20) match (implementation={atomic_default_mem_order(score(28):seq_cst)})
void f21 (void);
void f22 (void);
void f23 (void);
void f24 (void);
-#pragma omp declare variant (f22) match (construct={parallel,for}) /* 2+1 */
+#pragma omp declare variant (f22) match (construct={parallel,for}) /* 8+16 */
#pragma omp declare variant (f23) match (construct={for}) /* 0 */
#pragma omp declare variant (f24) match (implementation={atomic_default_mem_order(score(2):seq_cst)})
void f25 (void);
void f26 (void);
void f27 (void);
void f28 (void);
-#pragma omp declare variant (f26) match (construct={parallel,for}) /* 2+1 */
-#pragma omp declare variant (f27) match (construct={for},user={condition(1)}) /* 4 */
+#pragma omp declare variant (f26) match (construct={parallel,for}) /* 8+16 */
+#pragma omp declare variant (f27) match (construct={for},user={condition(score(25):1)}) /* 16 + 25 */
#pragma omp declare variant (f28) match (implementation={atomic_default_mem_order(score(3):seq_cst)})
void f29 (void);
@@ -1,5 +1,5 @@
/* { dg-do compile { target vect_simd_clones } } */
-/* { dg-additional-options "-fdump-tree-gimple" } */
+/* { dg-additional-options "-fdump-tree-ompdevlow" } */
/* { dg-additional-options "-mno-sse3" { target { i?86-*-* x86_64-*-* } } } */
int f01 (int);
@@ -20,5 +20,7 @@ test1 (int x)
isa has score 2^2 or 2^3. We can't decide on whether avx512f will match or
not, that also depends on whether it is a declare simd clone or not and which
one, but the f03 variant has a higher score anyway. */
- return f05 (x); /* { dg-final { scan-tree-dump-times "f03 \\\(x" 1 "gimple" } } */
+ return f05 (x);
+ /* { dg-final { scan-tree-dump "f03 \\\(x" "ompdevlow" } } */
+ /* { dg-final { scan-tree-dump-not "f05 \\\(x" "ompdevlow" } } */
}
@@ -38,8 +38,8 @@ void f18 (void);
void f19 (void);
#pragma omp declare variant (f1) match(user={condition()}) /* { dg-error "expected \[^\n\r]*expression before '\\)' token" } */
void f20 (void);
-#pragma omp declare variant (f1) match(user={condition(f1)}) /* { dg-error "property must be constant integer expression" "" { target { c || c++11 } } } */
-void f21 (void); /* { dg-error "cannot appear in a constant-expression" "" { target c++98_only } .-1 } */
+#pragma omp declare variant (f1) match(user={condition(f1)}) /* { dg-error "property must be integer expression" } */
+void f21 (void);
#pragma omp declare variant (f1) match(user={condition(1, 2, 3)}) /* { dg-error "expected '\\)' before ',' token" } */
void f22 (void);
#pragma omp declare variant (f1) match(construct={master}) /* { dg-warning "unknown selector 'master' for context selector set 'construct'" } */
new file mode 100644
@@ -0,0 +1,29 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-foffload=disable" } */
+/* { dg-additional-options "-mavx512bw -mavx512vl" { target { i?86-*-* x86_64-*-* } } } */
+
+/* References to function parameters in dynamic selector expressions for
+ "declare variant" isn't supported yet; see PR 113904. Check to see that
+ a proper error is diagnosed meanwhile and GCC doesn't just wander off
+ into the weeds and ICE. */
+
+extern int frob (int);
+
+void f01 (int, int);
+void f02 (int, int);
+void f03 (int, int);
+#pragma omp declare variant (f01) match (target_device={device_num (devnum), isa("avx512f","avx512vl")}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
+#pragma omp declare variant (f02) match (implementation={vendor(score(15):gnu)})
+#pragma omp declare variant (f03) match (user={condition(score(11):frob (ok + 42))}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
+void f04 (int devnum, int ok);
+
+void
+test1 (void)
+{
+ int i;
+ #pragma omp parallel for
+ for (i = 0; i < 1; i++)
+ f04 (17, 1);
+}
+
+
new file mode 100644
@@ -0,0 +1,26 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+extern int foo_p (int);
+extern int bar;
+
+int f01 (int);
+int f02 (int);
+int f03 (int);
+int f04 (int);
+#pragma omp declare variant (f01) match (device={isa("avx512f")}) /* 4 */
+#pragma omp declare variant (f02) match (implementation={vendor(score(3):gnu)},device={kind(cpu)}) /* 1 + 3 */
+#pragma omp declare variant (f03) match (user={condition(score(9):foo_p (bar))})
+#pragma omp declare variant (f04) match (implementation={vendor(score(6):gnu)},device={kind(host)}) /* 1 + 6 */
+int f05 (int);
+
+
+int
+test1 (int x)
+{
+ return f05 (x);
+}
+
+/* { dg-final { scan-tree-dump "f03 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump "f04 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "f05 \\\(x" "gimple" } } */
new file mode 100644
@@ -0,0 +1,30 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+extern int foo_p (int);
+extern int bar;
+extern int omp_get_default_device (void);
+
+int f01 (int);
+int f02 (int);
+int f03 (int);
+int f04 (int);
+#pragma omp declare variant (f01) match (target_device={device_num(omp_get_default_device()), isa("avx512f")}) /* 4 */
+#pragma omp declare variant (f02) match (user={condition(score(6):0)})
+#pragma omp declare variant (f03) match (user={condition(score(5):foo_p (bar))})
+#pragma omp declare variant (f04) match (user={condition(score(3):0)})
+int f05 (int);
+
+int
+test1 (int x)
+{
+ return f05 (x);
+}
+
+/* f01 and f03 are the dynamic selectors, the fall-through is f05.
+ f02 and f04 are static selectors and do not match. */
+/* { dg-final { scan-tree-dump "f01 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump "f03 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump "f05 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "f02 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "f04 \\\(x" "gimple" } } */
@@ -1,7 +1,5 @@
/* { dg-do compile } */
-/* { dg-additional-options "-fdump-tree-original" } */
/* { dg-additional-options "-fdump-tree-gimple" } */
-/* { dg-additional-options "-fdump-tree-optimized" } */
#define N 100
@@ -17,15 +15,7 @@ void f (int x[], int y[], int z[])
z[i] = x[i] * y[i];
}
-/* The metadirective should be resolved after Gimplification. */
-
-/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "when \\(device = .*arch.*nvptx.*\\):" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "default:" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } } */
-
-/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } } */
-
-/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } } */
+/* If offload device "nvptx" isn't supported, the front end can eliminate
+ that alternative and not produce a metadirective at all. Otherwise this
+ is supposed to be resolvable during gimplification. */
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
@@ -1,7 +1,5 @@
// { dg-do compile { target c++11 } }
-/* { dg-additional-options "-fdump-tree-original" } */
/* { dg-additional-options "-fdump-tree-gimple" } */
-/* { dg-additional-options "-fdump-tree-optimized" } */
#define N 100
@@ -18,14 +16,4 @@ void f (int x[], int y[], int z[])
}
/* The metadirective should be resolved after Gimplification. */
-
-/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "when \\(device = .*arch.*nvptx.*\\):" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "default:" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */
-/* { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } } */
-
-/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } } */
-
-/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } } */
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
new file mode 100644
@@ -0,0 +1,32 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-foffload=disable" } */
+/* { dg-additional-options "-mavx512bw -mavx512vl" { target { i?86-*-* x86_64-*-* } } } */
+
+/* References to function parameters in dynamic selector expressions for
+ "declare variant" isn't supported yet; see PR 113904. Check to see that
+ a proper error is diagnosed meanwhile and GCC doesn't just wander off
+ into the weeds and ICE. */
+
+extern int frob (int);
+
+class junk
+{
+ public:
+ int data;
+ static void f01 (int, int);
+ static void f02 (int, int);
+ static void f03 (int, int);
+#pragma omp declare variant (f01) match (target_device={device_num (devnum), isa("avx512f","avx512vl")}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
+#pragma omp declare variant (f02) match (implementation={vendor(score(15):gnu)})
+#pragma omp declare variant (f03) match (user={condition(score(11):frob (ok + 42))}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
+ static void f04 (int devnum, int ok);
+};
+
+void
+test1 (void)
+{
+ int i;
+ #pragma omp parallel for
+ for (i = 0; i < 1; i++)
+ junk::f04 (17, 1);
+}
new file mode 100644
@@ -0,0 +1,37 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-foffload=disable" } */
+/* { dg-additional-options "-mavx512bw -mavx512vl" { target { i?86-*-* x86_64-*-* } } } */
+
+/* References to function parameters in dynamic selector expressions for
+ "declare variant" isn't supported yet; see PR 113904. Check to see that
+ a proper error is diagnosed meanwhile and GCC doesn't just wander off
+ into the weeds and ICE. */
+
+extern int frob (int);
+extern int frobmore (class junk *);
+
+class junk
+{
+ public:
+ int data;
+ void f01 (int, int);
+ void f02 (int, int);
+ void f03 (int, int);
+ void f04 (int, int);
+ void f05 (int, int);
+#pragma omp declare variant (f01) match (target_device={device_num (devnum), isa("avx512f","avx512vl")}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
+#pragma omp declare variant (f02) match (implementation={vendor(score(15):gnu)})
+#pragma omp declare variant (f03) match (user={condition(score(11):frob (ok + 42))}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
+#pragma omp declare variant (f04) match (user={condition(score(11):data)}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
+#pragma omp declare variant (f05) match (user={condition(score(11):frobmore (this))}) /* { dg-message "sorry, unimplemented: reference to function parameter" } */
+ void f06 (int devnum, int ok);
+};
+
+void
+test1 (junk *j)
+{
+ int i;
+ #pragma omp parallel for
+ for (i = 0; i < 1; i++)
+ j->f06 (17, 1);
+}
@@ -64,9 +64,9 @@ contains
end subroutine
subroutine f17 ()
- !$omp declare variant (f14) match (construct={teams,parallel,do}) ! 16+8+4
- !$omp declare variant (f15) match (construct={parallel},user={condition(score(19):.true.)}) ! 8+19
- !$omp declare variant (f16) match (implementation={atomic_default_mem_order(score(27):seq_cst)})
+ !$omp declare variant (f14) match (construct={teams,parallel,do}) ! 1+8+16
+ !$omp declare variant (f15) match (construct={parallel},user={condition(score(16):.true.)}) ! 8+16
+ !$omp declare variant (f16) match (implementation={atomic_default_mem_order(score(24):seq_cst)})
end subroutine
subroutine f18 ()
@@ -79,7 +79,7 @@ contains
end subroutine
subroutine f21 ()
- !$omp declare variant (f18) match (construct={teams,parallel,do}) ! 16+8+4
+ !$omp declare variant (f18) match (construct={teams,parallel,do}) ! 1+8+16
!$omp declare variant (f19) match (construct={do},user={condition(score(25):.true.)}) ! 4+25
!$omp declare variant (f20) match (implementation={atomic_default_mem_order(score(28):seq_cst)})
end subroutine
@@ -94,7 +94,7 @@ contains
end subroutine
subroutine f25 ()
- !$omp declare variant (f22) match (construct={parallel,do}) ! 2+1
+ !$omp declare variant (f22) match (construct={parallel,do}) ! 8+16
!$omp declare variant (f23) match (construct={do}) ! 0
!$omp declare variant (f24) match (implementation={atomic_default_mem_order(score(2):seq_cst)})
end subroutine
@@ -109,8 +109,8 @@ contains
end subroutine
subroutine f29 ()
- !$omp declare variant (f26) match (construct={parallel,do}) ! 2+1
- !$omp declare variant (f27) match (construct={do},user={condition(.true.)}) ! 4
+ !$omp declare variant (f26) match (construct={parallel,do}) ! 8+16
+ !$omp declare variant (f27) match (construct={do},user={condition(score(25):.true.)}) ! 16+25
!$omp declare variant (f28) match (implementation={atomic_default_mem_order(score(3):seq_cst)})
end subroutine
@@ -1,28 +1,27 @@
! { dg-do compile { target vect_simd_clones } }
-! { dg-additional-options "-fdump-tree-gimple" }
+! { dg-additional-options "-O0 -fdump-tree-ompdevlow" }
! { dg-additional-options "-mno-sse3" { target { i?86-*-* x86_64-*-* } } }
-program main
- implicit none
+module foo
contains
integer function f01 (x)
integer, intent(in) :: x
- f01 = x
+ f01 = x + 1
end function
integer function f02 (x)
integer, intent(in) :: x
- f02 = x
+ f02 = x + 2
end function
integer function f03 (x)
integer, intent(in) :: x
- f03 = x
+ f03 = x + 3
end function
integer function f04 (x)
integer, intent(in) :: x
- f04 = x
+ f04 = x + 4
end function
integer function f05 (x)
@@ -32,9 +31,17 @@ contains
!$omp declare variant (f02) match (implementation={vendor(score(3):gnu)},device={kind(cpu)}) ! (1 or 2) + 3
!$omp declare variant (f03) match (user={condition(score(9):.true.)})
!$omp declare variant (f04) match (implementation={vendor(score(6):gnu)},device={kind(host)}) ! (1 or 2) + 6
- f05 = x
+ f05 = x + 5
end function
+end module
+program main
+ use :: foo
+ implicit none
+
+ if (test1 (42) /= 42 + 3) stop 100
+
+contains
integer function test1 (x)
!$omp declare simd
integer, intent(in) :: x
@@ -43,6 +50,9 @@ contains
! isa has score 2^2 or 2^3. We can't decide on whether avx512f will match or
! not, that also depends on whether it is a declare simd clone or not and which
! one, but the f03 variant has a higher score anyway. */
- test1 = f05 (x) ! { dg-final { scan-tree-dump-times "f03 \\\(x" 1 "gimple" } }
+ test1 = f05 (x)
+ ! { dg-final { scan-tree-dump "f03 \\\(x" "ompdevlow" } }
+ ! { dg-final { scan-tree-dump-not "f05 \\\(x" "ompdevlow" } }
end function
+
end program
@@ -44,8 +44,26 @@ program main
do i = 1, N
c(i) = a(i) * b(i)
end do
-
- !$omp begin metadirective &
+
+ !$omp metadirective &
+ !$omp& default (teams loop) &
+ !$omp& when (device={arch("nvptx")} parallel loop) ! { dg-error "expected .:." }
+ do i = 1, N
+ c(i) = a(i) * b(i)
+ end do
+
+ ! Test improperly nested metadirectives - even though the second
+ ! metadirective resolves to 'omp nothing', that is not the same as there
+ ! being literally nothing there.
+ !$omp metadirective &
+ !$omp& when (implementation={vendor("gnu")}: parallel do)
+ !$omp metadirective &
+ !$omp& when (implementation={vendor("cray")}: parallel do) ! { dg-error "Unexpected !.OMP METADIRECTIVE statement" }
+ do i = 1, N
+ c(i) = a(i) * b(i)
+ end do
+
+!$omp begin metadirective &
!$omp& when (device={arch("nvptx")}: parallel do) &
!$omp& default (barrier) ! { dg-error "variant directive used in OMP BEGIN METADIRECTIVE at .1. must have a corresponding end directive" }
do i = 1, N
@@ -1,7 +1,5 @@
! { dg-do compile }
-! { dg-additional-options "-fdump-tree-original" }
! { dg-additional-options "-fdump-tree-gimple" }
-! { dg-additional-options "-fdump-tree-optimized" }
module test
integer, parameter :: N = 100
@@ -20,15 +18,7 @@ contains
end subroutine
end module
-! The metadirective should be resolved after Gimplification.
-
-! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } }
-! { dg-final { scan-tree-dump-times "when \\(device =.*arch.*nvptx.*\\):" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } }
-! { dg-final { scan-tree-dump-times "default:" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } }
-
-! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } }
-
-! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } }
+! If offload device "nvptx" isn't supported, the front end can eliminate
+! that alternative and not produce a metadirective at all. Otherwise this
+! is supposed to be resolvable during gimplification.
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
@@ -1678,6 +1678,7 @@ remap_gimple_stmt (gimple *stmt, copy_body_data *id)
gimple *first_variant = NULL;
gimple **prev_next = &first_variant;
gimple_seq variant_seq = gimple_omp_variants (stmt);
+ tree context = gimple_omp_metadirective_context (stmt);
for (gimple_stmt_iterator gsi = gsi_start (variant_seq);
!gsi_end_p (gsi); gsi_next (&gsi))
{
@@ -1689,6 +1690,7 @@ remap_gimple_stmt (gimple *stmt, copy_body_data *id)
prev_next = &new_variant->next;
}
gimple_omp_metadirective_set_variants (copy, first_variant);
+ gimple_omp_metadirective_set_context (copy, context);
}
memset (&wi, 0, sizeof (wi));