@@ -75,7 +75,7 @@ extern void omp_firstprivatize_variable
extern enum gimplify_status gimplify_expr (tree *, gimple_seq *, gimple_seq *,
bool (*) (tree), fallback_t);
-HOST_WIDE_INT omp_construct_selector_matches (enum tree_code *, int);
+int omp_construct_selector_matches (enum tree_code *, int, int *);
extern void gimplify_type_sizes (tree, gimple_seq *);
extern void gimplify_one_sizepos (tree *, gimple_seq *);
@@ -10381,14 +10381,24 @@ gimplify_adjust_omp_clauses (gimple_seq
/* 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 positive number if they do, the number is then the number of constructs
- in the OpenMP context. */
+ 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. */
-HOST_WIDE_INT
-omp_construct_selector_matches (enum tree_code *constructs, int nconstructs)
+int
+omp_construct_selector_matches (enum tree_code *constructs, int nconstructs,
+ int *scores)
{
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;
for (struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; ctx;)
{
if (((ctx->region_type & ORT_PARALLEL) && ctx->code == OMP_PARALLEL)
@@ -10401,7 +10411,9 @@ omp_construct_selector_matches (enum tre
&& !omp_find_clause (ctx->clauses, OMP_CLAUSE_BIND)))
{
++cnt;
- if (matched < nconstructs && ctx->code == constructs[matched])
+ if (scores)
+ codes.safe_push (ctx->code);
+ else if (matched < nconstructs && ctx->code == constructs[matched])
{
if (ctx->code == OMP_SIMD)
{
@@ -10412,7 +10424,12 @@ omp_construct_selector_matches (enum tre
++matched;
}
if (ctx->code == OMP_TARGET)
- return matched < nconstructs ? 0 : simd_seen ? -1 : cnt;
+ {
+ if (scores == NULL)
+ return matched < nconstructs ? 0 : simd_seen ? -1 : 1;
+ target_seen = true;
+ break;
+ }
}
else if (ctx->region_type == ORT_WORKSHARE
&& ctx->code == OMP_LOOP
@@ -10424,31 +10441,40 @@ omp_construct_selector_matches (enum tre
ctx = ctx->outer_context->outer_context;
ctx = ctx->outer_context;
}
- if (cnt == 0
- && constructs[0] == OMP_SIMD
+ 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.c added clones and not to the base function. */
- gcc_assert (matched == 0);
- ++cnt;
- simd_seen = true;
- if (++matched == nconstructs)
- return -1;
+ 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)))
{
enum tree_code variant_constructs[5];
- int variant_nconstructs
- = omp_constructor_traits_to_codes (TREE_VALUE (attr),
- variant_constructs);
+ int variant_nconstructs = 0;
+ if (!target_seen)
+ variant_nconstructs
+ = omp_constructor_traits_to_codes (TREE_VALUE (attr),
+ variant_constructs);
for (int i = 0; i < variant_nconstructs; i++)
{
++cnt;
- if (matched < nconstructs
- && variant_constructs[i] == constructs[matched])
+ if (scores)
+ codes.safe_push (variant_constructs[i]);
+ else if (matched < nconstructs
+ && variant_constructs[i] == constructs[matched])
{
if (variant_constructs[i] == OMP_SIMD)
{
@@ -10460,15 +10486,38 @@ omp_construct_selector_matches (enum tre
}
}
}
- if (lookup_attribute ("omp declare target block",
- DECL_ATTRIBUTES (current_function_decl)))
+ if (!target_seen
+ && lookup_attribute ("omp declare target block",
+ DECL_ATTRIBUTES (current_function_decl)))
{
- ++cnt;
- if (matched < nconstructs && constructs[matched] == OMP_TARGET)
+ 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 : cnt;
+ return simd_seen ? -1 : 1;
return 0;
}
@@ -87,6 +87,7 @@ extern int omp_max_simt_vf (void);
extern int omp_constructor_traits_to_codes (tree, enum tree_code *);
extern int omp_context_selector_matches (tree);
extern int omp_context_selector_set_compare (const char *, tree, tree);
+extern tree omp_get_context_selector (tree, const char *, const char *);
extern tree omp_resolve_declare_variant (tree);
extern tree oacc_launch_pack (unsigned code, tree device, unsigned op);
extern tree oacc_replace_fn_attrib_attr (tree attribs, tree dims);
@@ -639,7 +639,7 @@ omp_maybe_offloaded (void)
if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0)
{
enum tree_code construct = OMP_TARGET;
- if (omp_construct_selector_matches (&construct, 1))
+ if (omp_construct_selector_matches (&construct, 1, NULL))
return true;
}
return false;
@@ -677,8 +677,8 @@ omp_context_selector_matches (tree ctx)
enum tree_code constructs[5];
int nconstructs
= omp_constructor_traits_to_codes (TREE_VALUE (t1), constructs);
- HOST_WIDE_INT r
- = omp_construct_selector_matches (constructs, nconstructs);
+ int r = omp_construct_selector_matches (constructs, nconstructs,
+ NULL);
if (r == 0)
return 0;
if (r == -1)
@@ -1261,13 +1261,93 @@ omp_context_selector_compare (tree ctx1,
return swapped ? -ret : ret;
}
+/* From context selector CTX, return trait-selector with name SEL in
+ trait-selector-set with name SET if any, or NULL_TREE if not found.
+ If SEL is NULL, return the list of trait-selectors in SET. */
+
+tree
+omp_get_context_selector (tree ctx, const char *set, const char *sel)
+{
+ tree setid = get_identifier (set);
+ tree selid = sel ? get_identifier (sel) : NULL_TREE;
+ for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
+ if (TREE_PURPOSE (t1) == setid)
+ {
+ if (sel == NULL)
+ return TREE_VALUE (t1);
+ for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
+ if (TREE_PURPOSE (t2) == selid)
+ return t2;
+ }
+ return NULL_TREE;
+}
+
+/* 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. */
+
+static bool
+omp_context_compute_score (tree ctx, widest_int *score, bool declare_simd)
+{
+ tree construct = omp_get_context_selector (ctx, "construct", NULL);
+ bool has_kind = omp_get_context_selector (ctx, "device", "kind");
+ bool has_arch = omp_get_context_selector (ctx, "device", "arch");
+ bool has_isa = omp_get_context_selector (ctx, "device", "isa");
+ bool ret = false;
+ *score = 1;
+ for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
+ for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
+ if (tree t3 = TREE_VALUE (t2))
+ if (TREE_PURPOSE (t3)
+ && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3)), " score") == 0
+ && TREE_CODE (TREE_VALUE (t3)) == INTEGER_CST)
+ *score += wi::to_widest (TREE_VALUE (t3));
+ if (construct || has_kind || has_arch || has_isa)
+ {
+ int scores[12];
+ enum tree_code constructs[5];
+ int nconstructs = 0;
+ if (construct)
+ nconstructs = omp_constructor_traits_to_codes (construct, constructs);
+ 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 = 0;
+ return ret;
+ }
+ *score += wi::shifted_mask <widest_int> (scores[b + n], 1, false);
+ }
+ if (has_kind)
+ *score += wi::shifted_mask <widest_int> (scores[b + nconstructs],
+ 1, false);
+ if (has_arch)
+ *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 1,
+ 1, false);
+ if (has_isa)
+ *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 2,
+ 1, false);
+ }
+ else /* FIXME: Implement this. */
+ gcc_unreachable ();
+ }
+ return ret;
+}
+
/* 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 variant = NULL_TREE;
+ tree variant1 = NULL_TREE, variant2 = NULL_TREE;
auto_vec <tree, 16> variants;
for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
{
@@ -1319,16 +1399,56 @@ omp_resolve_declare_variant (tree base)
variants[j] = NULL_TREE;
}
}
- /* FIXME: Scoring not implemented yet, so just resolve it
- if there is a single variant left. */
+ widest_int max_score1 = 0;
+ widest_int max_score2 = 0;
+ bool first = true;
FOR_EACH_VEC_ELT (variants, i, attr1)
if (attr1)
{
- if (variant)
- return base;
- variant = TREE_PURPOSE (TREE_VALUE (attr1));
+ if (variant1)
+ {
+ widest_int score1;
+ widest_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;
+ }
}
- return variant ? variant : base;
+ /* 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);
}
@@ -1193,7 +1193,6 @@ extern void c_omp_declare_simd_clauses_t
extern bool c_omp_predefined_variable (tree);
extern enum omp_clause_default_kind c_omp_predetermined_sharing (tree);
extern tree c_omp_check_context_selector (location_t, tree);
-extern tree c_omp_get_context_selector (tree, const char *, const char *);
extern void c_omp_mark_declare_variant (location_t, tree, tree);
/* Return next tree in the chain for chain_next walking of tree nodes. */
@@ -2237,27 +2237,6 @@ c_omp_check_context_selector (location_t
return ctx;
}
-/* From context selector CTX, return trait-selector with name SEL in
- trait-selector-set with name SET if any, or NULL_TREE if not found.
- If SEL is NULL, return the list of trait-selectors in SET. */
-
-tree
-c_omp_get_context_selector (tree ctx, const char *set, const char *sel)
-{
- tree setid = get_identifier (set);
- tree selid = sel ? get_identifier (sel) : NULL_TREE;
- for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
- if (TREE_PURPOSE (t1) == setid)
- {
- if (sel == NULL)
- return TREE_VALUE (t1);
- for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
- if (TREE_PURPOSE (t2) == selid)
- return t2;
- }
- return NULL_TREE;
-}
-
/* Register VARIANT as variant of some base function marked with
#pragma omp declare variant. CONSTRUCT is corresponding construct
selector set. */
@@ -19465,8 +19465,7 @@ c_finish_omp_declare_variant (c_parser *
error_at (token->location, "variant %qD is not a function", variant);
variant = error_mark_node;
}
- else if (c_omp_get_context_selector (ctx, "construct", "simd")
- == NULL_TREE
+ else if (omp_get_context_selector (ctx, "construct", "simd") == NULL_TREE
&& !comptypes (TREE_TYPE (fndecl), TREE_TYPE (variant)))
{
error_at (token->location, "variant %qD and base %qD have "
@@ -19487,7 +19486,7 @@ c_finish_omp_declare_variant (c_parser *
if (variant != error_mark_node)
{
C_DECL_USED (variant) = 1;
- tree construct = c_omp_get_context_selector (ctx, "construct", NULL);
+ tree construct = omp_get_context_selector (ctx, "construct", NULL);
c_omp_mark_declare_variant (match_loc, variant, construct);
if (omp_context_selector_matches (ctx))
{
@@ -7103,7 +7103,7 @@ omp_declare_variant_finalize_one (tree d
DECL_ARGUMENTS (decl), NULL);
tree ctx = TREE_VALUE (TREE_VALUE (attr));
- tree simd = c_omp_get_context_selector (ctx, "construct", "simd");
+ tree simd = omp_get_context_selector (ctx, "construct", "simd");
if (simd)
{
TREE_VALUE (simd)
@@ -7202,7 +7202,7 @@ omp_declare_variant_finalize_one (tree d
}
else
{
- tree construct = c_omp_get_context_selector (ctx, "construct", NULL);
+ tree construct = omp_get_context_selector (ctx, "construct", NULL);
c_omp_mark_declare_variant (match_loc, variant, construct);
if (!omp_context_selector_matches (ctx))
return true;
@@ -0,0 +1,88 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-foffload=disable -fdump-tree-gimple" } */
+/* { dg-additional-options "-mavx512bw -mavx512vl" { target { i?86-*-* x86_64-*-* } } } */
+
+#pragma omp requires atomic_default_mem_order(seq_cst)
+void f01 (void);
+void f02 (void);
+void f03 (void);
+#pragma omp declare variant (f01) match (device={isa(avx512f,avx512vl)}) /* 16 */
+#pragma omp declare variant (f02) match (implementation={vendor(score(15):gnu)})
+#pragma omp declare variant (f03) match (user={condition(score(11):1)})
+void f04 (void);
+void f05 (void);
+void f06 (void);
+void f07 (void);
+#pragma omp declare variant (f05) match (device={isa(avx512f,avx512vl)}) /* 16 */
+#pragma omp declare variant (f06) match (implementation={vendor(score(15):gnu)})
+#pragma omp declare variant (f07) match (user={condition(score(17):1)})
+void f08 (void);
+void f09 (void);
+void f10 (void);
+void f11 (void);
+void f12 (void);
+#pragma omp declare variant (f09) match (device={arch(x86_64)},user={condition(score(65):1)}) /* 64+65 */
+#pragma omp declare variant (f10) match (implementation={vendor(score(127):gnu)})
+#pragma omp declare variant (f11) match (device={isa(ssse3)}) /* 128 */
+#pragma omp declare variant (f12) match (implementation={atomic_default_mem_order(score(126):seq_cst)})
+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)})
+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 (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 (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 (f28) match (implementation={atomic_default_mem_order(score(3):seq_cst)})
+void f29 (void);
+
+void
+test1 (void)
+{
+ int i, j;
+ #pragma omp parallel for /* 2 constructs in OpenMP context, isa has score 2^4. */
+ for (i = 0; i < 1; i++)
+ f04 (); /* { dg-final { scan-tree-dump-times "f01 \\\(\\\);" 1 "gimple" { target i?86-*-* x86_64-*-* } } } */
+ /* { dg-final { scan-tree-dump-times "f02 \\\(\\\);" 1 "gimple" { target { ! { i?86-*-* x86_64-*-* } } } } } */
+ #pragma omp target teams /* 2 constructs in OpenMP context, isa has score 2^4. */
+ f08 (); /* { dg-final { scan-tree-dump-times "f07 \\\(\\\);" 1 "gimple" } } */
+ #pragma omp teams
+ #pragma omp parallel for
+ for (i = 0; i < 1; i++)
+ #pragma omp parallel for /* 5 constructs in OpenMP context, arch is 2^6, isa 2^7. */
+ for (j = 0; j < 1; j++)
+ {
+ f13 (); /* { dg-final { scan-tree-dump-times "f09 \\\(\\\);" 1 "gimple" { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+ /* { dg-final { scan-tree-dump-times "f11 \\\(\\\);" 1 "gimple" { target { { i?86-*-* x86_64-*-* } && { ! lp64 } } } } } */
+ /* { dg-final { scan-tree-dump-times "f10 \\\(\\\);" 1 "gimple" { target { ! { i?86-*-* x86_64-*-* } } } } } */
+ f17 (); /* { dg-final { scan-tree-dump-times "f14 \\\(\\\);" 1 "gimple" } } */
+ f21 (); /* { dg-final { scan-tree-dump-times "f19 \\\(\\\);" 1 "gimple" } } */
+ }
+ #pragma omp for
+ for (i = 0; i < 1; i++)
+ #pragma omp parallel for
+ for (j = 0; j < 1; j++)
+ {
+ f25 (); /* { dg-final { scan-tree-dump-times "f22 \\\(\\\);" 1 "gimple" } } */
+ f29 (); /* { dg-final { scan-tree-dump-times "f27 \\\(\\\);" 1 "gimple" } } */
+ }
+}