* gimplify.h (omp_construct_selector_matches): Declare.
* gimplify.c (struct gimplify_omp_ctx): Add code member.
(gimplify_call_expr): Call omp_resolve_declare_variant and remap
called function if needed for flag_openmp.
(gimplify_scan_omp_clauses): Set ctx->code.
(omp_construct_selector_matches): New function.
* omp-general.h (omp_constructor_traits_to_codes,
omp_context_selector_matches, omp_resolve_declare_variant): Declare.
* omp-general.c (omp_constructor_traits_to_codes,
omp_context_selector_matches, omp_resolve_declare_variant): New
functions.
c-family/
* c-common.h (c_omp_context_selector_matches): Remove.
* c-omp.c (c_omp_context_selector_matches): Remove.
* c-attribs.c (c_common_attribute_table): Add
"omp declare target {host,nohost,block}" attributes.
c/
* c-parser.c (c_finish_omp_declare_variant): Use
omp_context_selector_matches instead of
c_omp_context_selector_matches.
* c-decl.c (c_decl_attributes): Add "omp declare target block"
attribute in between declare target and end declare target
pragmas.
cp/
* decl2.c (cplus_decl_attributes): Add "omp declare target block"
attribute in between declare target and end declare target
pragmas.
testsuite/
* c-c++-common/gomp/declare-variant-8.c: New test.
From-SVN: r277427
2019-10-24 Jakub Jelinek <jakub@redhat.com>
+ * gimplify.h (omp_construct_selector_matches): Declare.
+ * gimplify.c (struct gimplify_omp_ctx): Add code member.
+ (gimplify_call_expr): Call omp_resolve_declare_variant and remap
+ called function if needed for flag_openmp.
+ (gimplify_scan_omp_clauses): Set ctx->code.
+ (omp_construct_selector_matches): New function.
+ * omp-general.h (omp_constructor_traits_to_codes,
+ omp_context_selector_matches, omp_resolve_declare_variant): Declare.
+ * omp-general.c (omp_constructor_traits_to_codes,
+ omp_context_selector_matches, omp_resolve_declare_variant): New
+ functions.
+
* config/arc/arc.c (hwloop_optimize): Add missing space in string
literal.
* config/rx/rx.c (rx_print_operand): Likewise.
+2019-10-24 Jakub Jelinek <jakub@redhat.com>
+
+ * c-common.h (c_omp_context_selector_matches): Remove.
+ * c-omp.c (c_omp_context_selector_matches): Remove.
+ * c-attribs.c (c_common_attribute_table): Add
+ "omp declare target {host,nohost,block}" attributes.
+
2019-10-17 JeanHeyd Meneide <phdofthehouse@gmail.com>
* c-lex.c (c_common_has_attribute): Update nodiscard value.
handle_omp_declare_target_attribute, NULL },
{ "omp declare target implicit", 0, 0, true, false, false, false,
handle_omp_declare_target_attribute, NULL },
+ { "omp declare target host", 0, 0, true, false, false, false,
+ handle_omp_declare_target_attribute, NULL },
+ { "omp declare target nohost", 0, 0, true, false, false, false,
+ handle_omp_declare_target_attribute, NULL },
+ { "omp declare target block", 0, 0, true, false, false, false,
+ handle_omp_declare_target_attribute, NULL },
{ "alloc_align", 1, 1, false, true, true, false,
handle_alloc_align_attribute,
attr_alloc_exclusions },
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);
-extern int c_omp_context_selector_matches (tree);
/* Return next tree in the chain for chain_next walking of tree nodes. */
static inline tree
#include "memmodel.h"
#include "attribs.h"
#include "gimplify.h"
-#include "cgraph.h"
-#include "symbol-summary.h"
-#include "hsa-common.h"
/* Complete a #pragma oacc wait construct. LOC is the location of
error_at (loc, "%qD used as a variant with incompatible %<constructor%> "
"selector sets", variant);
}
-
-/* Return 1 if context selector 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
- IPA, others until vectorization. */
-
-int
-c_omp_context_selector_matches (tree ctx)
-{
- int ret = 1;
- for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
- {
- char set = IDENTIFIER_POINTER (TREE_PURPOSE (t1))[0];
- if (set == 'c')
- {
- /* 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. */
- ret = -1;
- continue;
- }
- for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
- {
- const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
- switch (*sel)
- {
- case 'v':
- if (set == 'i' && !strcmp (sel, "vendor"))
- for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
- {
- const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
- if (!strcmp (prop, " score") || !strcmp (prop, "gnu"))
- continue;
- return 0;
- }
- break;
- case 'e':
- if (set == 'i' && !strcmp (sel, "extension"))
- /* We don't support any extensions right now. */
- return 0;
- break;
- case 'a':
- if (set == 'i' && !strcmp (sel, "atomic_default_mem_order"))
- {
- 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. */
- ret = -1;
- break;
- }
- tree t3 = TREE_VALUE (t2);
- const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
- if (!strcmp (prop, " score"))
- {
- t3 = TREE_CHAIN (t3);
- prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
- }
- 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;
- }
- if (set == 'd' && !strcmp (sel, "arch"))
- /* For now, need a target hook. */
- ret = -1;
- break;
- case 'u':
- if (set == 'i' && !strcmp (sel, "unified_address"))
- {
- if ((omp_requires_mask & OMP_REQUIRES_UNIFIED_ADDRESS) == 0)
- ret = -1;
- break;
- }
- if (set == 'i' && !strcmp (sel, "unified_shared_memory"))
- {
- if ((omp_requires_mask
- & OMP_REQUIRES_UNIFIED_SHARED_MEMORY) == 0)
- ret = -1;
- break;
- }
- break;
- case 'd':
- if (set == 'i' && !strcmp (sel, "dynamic_allocators"))
- {
- if ((omp_requires_mask
- & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
- ret = -1;
- break;
- }
- break;
- case 'r':
- if (set == 'i' && !strcmp (sel, "reverse_offload"))
- {
- if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
- ret = -1;
- break;
- }
- break;
- case 'k':
- if (set == 'd' && !strcmp (sel, "kind"))
- for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
- {
- const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
- if (!strcmp (prop, "any"))
- continue;
- if (!strcmp (prop, "fpga"))
- return 0; /* Right now GCC doesn't support any fpgas. */
- if (!strcmp (prop, "host"))
- {
- if (ENABLE_OFFLOADING || hsa_gen_requested_p ())
- ret = -1;
- continue;
- }
- if (!strcmp (prop, "nohost"))
- {
- if (ENABLE_OFFLOADING || hsa_gen_requested_p ())
- ret = -1;
- else
- return 0;
- continue;
- }
- if (!strcmp (prop, "cpu") || !strcmp (prop, "gpu"))
- {
- bool maybe_gpu = false;
- if (hsa_gen_requested_p ())
- maybe_gpu = true;
- else if (ENABLE_OFFLOADING)
- for (const char *c = getenv ("OFFLOAD_TARGET_NAMES");
- c; )
- {
- if (!strncmp (c, "nvptx", strlen ("nvptx"))
- || !strncmp (c, "amdgcn", strlen ("amdgcn")))
- {
- maybe_gpu = true;
- break;
- }
- else if ((c = strchr (c, ',')))
- c++;
- }
- if (!maybe_gpu)
- {
- if (prop[0] == 'g')
- return 0;
- }
- else
- ret = -1;
- continue;
- }
- /* Any other kind doesn't match. */
- return 0;
- }
- break;
- case 'i':
- if (set == 'd' && !strcmp (sel, "isa"))
- /* For now, need a target hook. */
- ret = -1;
- break;
- case 'c':
- if (set == 'u' && !strcmp (sel, "condition"))
- for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
- if (TREE_PURPOSE (t3) == NULL_TREE
- && integer_zerop (TREE_VALUE (t3)))
- return 0;
- break;
- default:
- break;
- }
- }
- }
- return ret;
-}
+2019-10-24 Jakub Jelinek <jakub@redhat.com>
+
+ * c-parser.c (c_finish_omp_declare_variant): Use
+ omp_context_selector_matches instead of
+ c_omp_context_selector_matches.
+ * c-decl.c (c_decl_attributes): Add "omp declare target block"
+ attribute in between declare target and end declare target
+ pragmas.
+
2019-10-15 Joseph Myers <joseph@codesourcery.com>
* c-parser.c (c_parser_attribute_any_word): Rename to
attributes = tree_cons (get_identifier ("omp declare target implicit"),
NULL_TREE, attributes);
else
- attributes = tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, attributes);
+ {
+ attributes = tree_cons (get_identifier ("omp declare target"),
+ NULL_TREE, attributes);
+ attributes = tree_cons (get_identifier ("omp declare target block"),
+ NULL_TREE, attributes);
+ }
}
/* Look up the current declaration with all the attributes merged
C_DECL_USED (variant) = 1;
tree construct = c_omp_get_context_selector (ctx, "construct", NULL);
c_omp_mark_declare_variant (match_loc, variant, construct);
- if (c_omp_context_selector_matches (ctx))
+ if (omp_context_selector_matches (ctx))
{
tree attr
= tree_cons (get_identifier ("omp declare variant base"),
2019-10-24 Jakub Jelinek <jakub@redhat.com>
+ * decl2.c (cplus_decl_attributes): Add "omp declare target block"
+ attribute in between declare target and end declare target
+ pragmas.
+
* call.c (convert_arg_to_ellipsis): Add missing space in string
literal.
attributes = tree_cons (get_identifier ("omp declare target implicit"),
NULL_TREE, attributes);
else
- attributes = tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, attributes);
+ {
+ attributes = tree_cons (get_identifier ("omp declare target"),
+ NULL_TREE, attributes);
+ attributes = tree_cons (get_identifier ("omp declare target block"),
+ NULL_TREE, attributes);
+ }
}
if (processing_template_decl)
location_t location;
enum omp_clause_default_kind default_kind;
enum omp_region_type region_type;
+ enum tree_code code;
bool combined_loop;
bool distribute;
bool target_firstprivatize_array_bases;
/* Remember the original function pointer type. */
fnptrtype = TREE_TYPE (CALL_EXPR_FN (*expr_p));
+ if (flag_openmp && fndecl)
+ {
+ tree variant = omp_resolve_declare_variant (fndecl);
+ if (variant != fndecl)
+ CALL_EXPR_FN (*expr_p) = build1 (ADDR_EXPR, fnptrtype, variant);
+ }
+
/* There is a sequence point before the call, so any side effects in
the calling expression must occur before the actual call. Force
gimplify_expr to use an internal post queue. */
int nowait = -1;
ctx = new_omp_context (region_type);
+ ctx->code = code;
outer_ctx = ctx->outer_context;
if (code == OMP_TARGET)
{
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 positive number if they do, the number is then the number of constructs
+ in the OpenMP context. */
+
+HOST_WIDE_INT
+omp_construct_selector_matches (enum tree_code *constructs, int nconstructs)
+{
+ int matched = 0, cnt = 0;
+ bool simd_seen = false;
+ 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)))
+ {
+ ++cnt;
+ 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)
+ return matched < nconstructs ? 0 : simd_seen ? -1 : cnt;
+ }
+ else if (ctx->region_type == ORT_WORKSHARE
+ && ctx->code == OMP_LOOP
+ && ctx->outer_context
+ && ctx->outer_context->region_type == ORT_COMBINED_PARALLEL
+ && ctx->outer_context->outer_context
+ && ctx->outer_context->outer_context->code == OMP_LOOP
+ && ctx->outer_context->outer_context->distribute)
+ ctx = ctx->outer_context->outer_context;
+ ctx = ctx->outer_context;
+ }
+ if (cnt == 0
+ && constructs[0] == OMP_SIMD
+ && 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;
+ }
+ 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);
+ for (int i = 0; i < variant_nconstructs; i++)
+ {
+ ++cnt;
+ if (matched < nconstructs
+ && variant_constructs[i] == constructs[matched])
+ {
+ if (variant_constructs[i] == OMP_SIMD)
+ {
+ if (matched)
+ return 0;
+ simd_seen = true;
+ }
+ ++matched;
+ }
+ }
+ }
+ if (lookup_attribute ("omp declare target block",
+ DECL_ATTRIBUTES (current_function_decl)))
+ {
+ ++cnt;
+ if (matched < nconstructs && constructs[matched] == OMP_TARGET)
+ ++matched;
+ }
+ if (matched == nconstructs)
+ return simd_seen ? -1 : cnt;
+ return 0;
+}
+
/* Gimplify OACC_CACHE. */
static void
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);
+
extern void gimplify_type_sizes (tree, gimple_seq *);
extern void gimplify_one_sizepos (tree *, gimple_seq *);
extern gbind *gimplify_body (tree, bool);
#include "omp-general.h"
#include "stringpool.h"
#include "attribs.h"
+#include "gimplify.h"
+#include "cgraph.h"
+#include "symbol-summary.h"
+#include "hsa-common.h"
+#include "tree-pass.h"
enum omp_requires omp_requires_mask;
return 0;
}
+/* Store the construct selectors as tree codes from last to first,
+ return their number. */
+
+int
+omp_constructor_traits_to_codes (tree ctx, enum tree_code *constructs)
+{
+ int nconstructs = list_length (ctx);
+ int i = nconstructs - 1;
+ for (tree t2 = ctx; t2; t2 = TREE_CHAIN (t2), i--)
+ {
+ const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
+ if (!strcmp (sel, "target"))
+ constructs[i] = OMP_TARGET;
+ else if (!strcmp (sel, "teams"))
+ constructs[i] = OMP_TEAMS;
+ else if (!strcmp (sel, "parallel"))
+ constructs[i] = OMP_PARALLEL;
+ else if (!strcmp (sel, "for") || !strcmp (sel, "do"))
+ constructs[i] = OMP_FOR;
+ else if (!strcmp (sel, "simd"))
+ constructs[i] = OMP_SIMD;
+ else
+ gcc_unreachable ();
+ }
+ gcc_assert (i == -1);
+ return nconstructs;
+}
+
+/* Return 1 if context selector 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
+ IPA, others until vectorization. */
+
+int
+omp_context_selector_matches (tree ctx)
+{
+ int ret = 1;
+ for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
+ {
+ char set = IDENTIFIER_POINTER (TREE_PURPOSE (t1))[0];
+ if (set == 'c')
+ {
+ /* 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. */
+ if (symtab->state == PARSING
+ || (cfun->curr_properties & PROP_gimple_any) != 0)
+ {
+ ret = -1;
+ continue;
+ }
+
+ 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);
+ if (r == 0)
+ return 0;
+ if (r == -1)
+ ret = -1;
+ continue;
+ }
+ for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
+ {
+ const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
+ switch (*sel)
+ {
+ case 'v':
+ if (set == 'i' && !strcmp (sel, "vendor"))
+ for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
+ {
+ const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
+ if (!strcmp (prop, " score") || !strcmp (prop, "gnu"))
+ continue;
+ return 0;
+ }
+ break;
+ case 'e':
+ if (set == 'i' && !strcmp (sel, "extension"))
+ /* We don't support any extensions right now. */
+ return 0;
+ break;
+ case 'a':
+ if (set == 'i' && !strcmp (sel, "atomic_default_mem_order"))
+ {
+ 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 t3 = TREE_VALUE (t2);
+ const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
+ if (!strcmp (prop, " score"))
+ {
+ t3 = TREE_CHAIN (t3);
+ prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
+ }
+ 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;
+ }
+ if (set == 'd' && !strcmp (sel, "arch"))
+ /* For now, need a target hook. */
+ ret = -1;
+ break;
+ case 'u':
+ if (set == 'i' && !strcmp (sel, "unified_address"))
+ {
+ if ((omp_requires_mask & OMP_REQUIRES_UNIFIED_ADDRESS) == 0)
+ {
+ if (symtab->state == PARSING)
+ ret = -1;
+ else
+ return 0;
+ }
+ break;
+ }
+ if (set == 'i' && !strcmp (sel, "unified_shared_memory"))
+ {
+ if ((omp_requires_mask
+ & OMP_REQUIRES_UNIFIED_SHARED_MEMORY) == 0)
+ {
+ if (symtab->state == PARSING)
+ ret = -1;
+ else
+ return 0;
+ }
+ break;
+ }
+ break;
+ case 'd':
+ if (set == 'i' && !strcmp (sel, "dynamic_allocators"))
+ {
+ if ((omp_requires_mask
+ & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
+ {
+ if (symtab->state == PARSING)
+ ret = -1;
+ else
+ return 0;
+ }
+ break;
+ }
+ break;
+ case 'r':
+ if (set == 'i' && !strcmp (sel, "reverse_offload"))
+ {
+ if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
+ {
+ if (symtab->state == PARSING)
+ ret = -1;
+ else
+ return 0;
+ }
+ break;
+ }
+ break;
+ case 'k':
+ if (set == 'd' && !strcmp (sel, "kind"))
+ for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
+ {
+ const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
+ if (!strcmp (prop, "any"))
+ continue;
+ if (!strcmp (prop, "fpga"))
+ return 0; /* Right now GCC doesn't support any fpgas. */
+ if (!strcmp (prop, "host"))
+ {
+ if (ENABLE_OFFLOADING || hsa_gen_requested_p ())
+ ret = -1;
+ continue;
+ }
+ if (!strcmp (prop, "nohost"))
+ {
+ if (ENABLE_OFFLOADING || hsa_gen_requested_p ())
+ ret = -1;
+ else
+ return 0;
+ continue;
+ }
+ if (!strcmp (prop, "cpu") || !strcmp (prop, "gpu"))
+ {
+ bool maybe_gpu = false;
+ if (hsa_gen_requested_p ())
+ maybe_gpu = true;
+ else if (ENABLE_OFFLOADING)
+ for (const char *c = getenv ("OFFLOAD_TARGET_NAMES");
+ c; )
+ {
+ if (!strncmp (c, "nvptx", strlen ("nvptx"))
+ || !strncmp (c, "amdgcn", strlen ("amdgcn")))
+ {
+ maybe_gpu = true;
+ break;
+ }
+ else if ((c = strchr (c, ',')))
+ c++;
+ }
+ if (!maybe_gpu)
+ {
+ if (prop[0] == 'g')
+ return 0;
+ }
+ else
+ ret = -1;
+ continue;
+ }
+ /* Any other kind doesn't match. */
+ return 0;
+ }
+ break;
+ case 'i':
+ if (set == 'd' && !strcmp (sel, "isa"))
+ /* For now, need a target hook. */
+ ret = -1;
+ break;
+ case 'c':
+ if (set == 'u' && !strcmp (sel, "condition"))
+ for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
+ if (TREE_PURPOSE (t3) == NULL_TREE)
+ {
+ if (integer_zerop (TREE_VALUE (t3)))
+ return 0;
+ if (integer_nonzerop (TREE_VALUE (t3)))
+ break;
+ ret = -1;
+ }
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ 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;
+ for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
+ {
+ attr = lookup_attribute ("omp declare variant base", attr);
+ if (attr == NULL_TREE)
+ break;
+ switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr))))
+ {
+ case 0:
+ /* No match, ignore. */
+ break;
+ case -1:
+ /* Needs to be deferred. */
+ return base;
+ default:
+ /* FIXME: Scoring not implemented yet, so just resolve it
+ if there is a single variant only. */
+ if (variant)
+ return base;
+ if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr))) == FUNCTION_DECL)
+ variant = TREE_PURPOSE (TREE_VALUE (attr));
+ else
+ return base;
+ }
+ }
+ return variant ? variant : base;
+}
+
+
/* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
macro on gomp-constants.h. We do not check for overflow. */
extern gimple *omp_build_barrier (tree lhs);
extern poly_uint64 omp_max_vf (void);
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 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);
extern void oacc_replace_fn_attrib (tree fn, tree dims);
+2019-10-24 Jakub Jelinek <jakub@redhat.com>
+
+ * c-c++-common/gomp/declare-variant-8.c: New test.
+
2019-10-24 Andreas Krebbel <krebbel@linux.ibm.com>
* gcc.dg/ipa/ipa-sra-19.c: Remove dg-skip-if. Add argument type to
--- /dev/null
+/* { dg-do compile { target c } } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void f01 (void);
+#pragma omp declare variant (f01) match (user={condition(6 == 7)},implementation={vendor(gnu)})
+void f02 (void);
+void f03 (void);
+#pragma omp declare variant (f03) match (user={condition(6 == 6)},implementation={atomic_default_mem_order(seq_cst)})
+void f04 (void);
+void f05 (void);
+#pragma omp declare variant (f05) match (user={condition(1)},implementation={atomic_default_mem_order(relaxed)})
+void f06 (void);
+#pragma omp requires atomic_default_mem_order(seq_cst)
+void f07 (void);
+#pragma omp declare variant (f07) match (construct={parallel,for},device={kind(any)})
+void f08 (void);
+void f09 (void);
+#pragma omp declare variant (f09) match (construct={parallel,for},implementation={vendor(gnu)})
+void f10 (void);
+void f11 (void);
+#pragma omp declare variant (f11) match (construct={parallel,for})
+void f12 (void);
+void f13 (void);
+#pragma omp declare variant (f13) match (construct={parallel,for})
+void f14 (void);
+#pragma omp declare target to (f13, f14)
+void f15 (void);
+#pragma omp declare variant (f15) match (implementation={vendor(llvm)})
+void f16 (void);
+void f17 (void);
+#pragma omp declare variant (f17) match (construct={target,parallel})
+void f18 (void);
+void f19 (void);
+#pragma omp declare variant (f19) match (construct={target,parallel})
+void f20 (void);
+void f21 (void);
+#pragma omp declare variant (f21) match (construct={teams,parallel})
+void f22 (void);
+void f23 (void);
+#pragma omp declare variant (f23) match (construct={teams,parallel,for})
+void f24 (void);
+void f25 (void);
+#pragma omp declare variant (f25) match (construct={teams,parallel})
+void f26 (void);
+void f27 (void);
+#pragma omp declare variant (f27) match (construct={teams,parallel,for})
+void f28 (void);
+void f29 (void);
+#pragma omp declare variant (f29) match (implementation={vendor(gnu)})
+void f30 (void);
+void f31 (void);
+#pragma omp declare variant (f31) match (construct={teams,parallel,for})
+void f32 (void);
+
+void
+test1 (void)
+{
+ int i;
+ f02 (); /* { dg-final { scan-tree-dump-times "f02 \\\(\\\);" 1 "gimple" } } */
+ f04 (); /* { dg-final { scan-tree-dump-times "f03 \\\(\\\);" 1 "gimple" } } */
+ f06 (); /* { dg-final { scan-tree-dump-times "f06 \\\(\\\);" 1 "gimple" } } */
+ #pragma omp parallel
+ #pragma omp for
+ for (i = 0; i < 1; i++)
+ f08 (); /* { dg-final { scan-tree-dump-times "f07 \\\(\\\);" 1 "gimple" } } */
+ #pragma omp parallel for
+ for (i = 0; i < 1; i++)
+ f10 (); /* { dg-final { scan-tree-dump-times "f09 \\\(\\\);" 1 "gimple" } } */
+ #pragma omp for
+ for (i = 0; i < 1; i++)
+ #pragma omp parallel
+ f12 (); /* { dg-final { scan-tree-dump-times "f12 \\\(\\\);" 1 "gimple" } } */
+ #pragma omp parallel
+ #pragma omp target
+ #pragma omp for
+ for (i = 0; i < 1; i++)
+ f14 (); /* { dg-final { scan-tree-dump-times "f14 \\\(\\\);" 1 "gimple" } } */
+ f16 (); /* { dg-final { scan-tree-dump-times "f16 \\\(\\\);" 1 "gimple" } } */
+}
+
+#pragma omp declare target
+void
+test2 (void)
+{
+ #pragma omp parallel
+ f18 (); /* { dg-final { scan-tree-dump-times "f17 \\\(\\\);" 1 "gimple" } } */
+}
+#pragma omp end declare target
+
+void test3 (void);
+#pragma omp declare target to (test3)
+
+void
+test3 (void)
+{
+ #pragma omp parallel
+ f20 (); /* { dg-final { scan-tree-dump-times "f20 \\\(\\\);" 1 "gimple" } } */
+}
+
+void
+f21 (void)
+{
+ int i;
+ #pragma omp for
+ for (i = 0; i < 1; i++)
+ f24 (); /* { dg-final { scan-tree-dump-times "f23 \\\(\\\);" 1 "gimple" } } */
+}
+
+void
+f26 (void)
+{
+ int i;
+ #pragma omp for
+ for (i = 0; i < 1; i++)
+ f28 (); /* { dg-final { scan-tree-dump-times "f28 \\\(\\\);" 1 "gimple" } } */
+}
+
+void
+f29 (void)
+{
+ int i;
+ #pragma omp for
+ for (i = 0; i < 1; i++)
+ f32 (); /* { dg-final { scan-tree-dump-times "f32 \\\(\\\);" 1 "gimple" } } */
+}
projects / gcc.git / commitdiff