From 135df52cc3f1ef90092ab02a01c4dabc7fd0ed18 Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Fri, 25 Oct 2019 00:29:09 +0200 Subject: [PATCH] gimplify.h (omp_construct_selector_matches): Declare. * 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 --- gcc/ChangeLog | 12 + gcc/c-family/ChangeLog | 7 + gcc/c-family/c-attribs.c | 6 + gcc/c-family/c-common.h | 1 - gcc/c-family/c-omp.c | 188 ----------- gcc/c/ChangeLog | 9 + gcc/c/c-decl.c | 8 +- gcc/c/c-parser.c | 2 +- gcc/cp/ChangeLog | 4 + gcc/cp/decl2.c | 8 +- gcc/gimplify.c | 102 ++++++ gcc/gimplify.h | 2 + gcc/omp-general.c | 298 ++++++++++++++++++ gcc/omp-general.h | 3 + gcc/testsuite/ChangeLog | 4 + .../c-c++-common/gomp/declare-variant-8.c | 125 ++++++++ 16 files changed, 585 insertions(+), 194 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-variant-8.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 9d199750036..4d9e4709da2 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,5 +1,17 @@ 2019-10-24 Jakub Jelinek + * 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. diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog index aeae1497b92..3b3d5744bac 100644 --- a/gcc/c-family/ChangeLog +++ b/gcc/c-family/ChangeLog @@ -1,3 +1,10 @@ +2019-10-24 Jakub Jelinek + + * 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 * c-lex.c (c_common_has_attribute): Update nodiscard value. diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c index ea273f89a8e..1c9f28587fb 100644 --- a/gcc/c-family/c-attribs.c +++ b/gcc/c-family/c-attribs.c @@ -456,6 +456,12 @@ const struct attribute_spec c_common_attribute_table[] = 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 }, diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 3bc021b65d9..70771834502 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1193,7 +1193,6 @@ 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); -extern int c_omp_context_selector_matches (tree); /* Return next tree in the chain for chain_next walking of tree nodes. */ static inline tree diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index 339818817e2..4f5a6ed175a 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -34,9 +34,6 @@ along with GCC; see the file COPYING3. If not see #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 @@ -2388,188 +2385,3 @@ c_omp_mark_declare_variant (location_t loc, tree variant, tree construct) error_at (loc, "%qD used as a variant with incompatible % " "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; -} diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 0cf32f6e638..e8518f27494 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,12 @@ +2019-10-24 Jakub Jelinek + + * 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 * c-parser.c (c_parser_attribute_any_word): Rename to diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c index f67033b82d2..c6c4a4d7aeb 100644 --- a/gcc/c/c-decl.c +++ b/gcc/c/c-decl.c @@ -4832,8 +4832,12 @@ c_decl_attributes (tree *node, tree attributes, int flags) 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 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 7618a46c8bc..9589cc68c25 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -19489,7 +19489,7 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms) 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"), diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 33420f1289d..78bbfb5b937 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,5 +1,9 @@ 2019-10-24 Jakub Jelinek + * 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. diff --git a/gcc/cp/decl2.c b/gcc/cp/decl2.c index 6d5e973b487..cff11baef46 100644 --- a/gcc/cp/decl2.c +++ b/gcc/cp/decl2.c @@ -1555,8 +1555,12 @@ cplus_decl_attributes (tree *decl, tree attributes, int flags) 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) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 914bb8eb8d6..05ae2f1552b 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -219,6 +219,7 @@ struct gimplify_omp_ctx 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; @@ -3385,6 +3386,13 @@ 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)); + 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. */ @@ -8137,6 +8145,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, int nowait = -1; ctx = new_omp_context (region_type); + ctx->code = code; outer_ctx = ctx->outer_context; if (code == OMP_TARGET) { @@ -10324,6 +10333,99 @@ 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 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 diff --git a/gcc/gimplify.h b/gcc/gimplify.h index 6c997a769cd..601b82b12bd 100644 --- a/gcc/gimplify.h +++ b/gcc/gimplify.h @@ -75,6 +75,8 @@ 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); +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); diff --git a/gcc/omp-general.c b/gcc/omp-general.c index 1a78a70bd57..9397b1951f8 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -35,6 +35,11 @@ along with GCC; see the file COPYING3. If not see #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; @@ -538,6 +543,299 @@ omp_max_simt_vf (void) 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. */ diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 7cd1d216fc0..c0c294d29f4 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -84,6 +84,9 @@ extern void omp_extract_for_data (gomp_for *for_stmt, struct omp_for_data *fd, 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); diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index e999a4781d1..8e652ab3cbf 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,7 @@ +2019-10-24 Jakub Jelinek + + * c-c++-common/gomp/declare-variant-8.c: New test. + 2019-10-24 Andreas Krebbel * gcc.dg/ipa/ipa-sra-19.c: Remove dg-skip-if. Add argument type to diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-8.c b/gcc/testsuite/c-c++-common/gomp/declare-variant-8.c new file mode 100644 index 00000000000..792e56ccd97 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-8.c @@ -0,0 +1,125 @@ +/* { 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" } } */ +} -- 2.30.2