Hi!
The following patch extends the omp declare variant handling to perform
another check during gimplification, where it already can redirect various
calls to base functions to their corresponding variants.
The scoring is still unimplemented, so right now it does the redirection
only if there is a single matching variant.
As the C++ omp declare variant patch has not been committed yet,
it is for now limited to C.
Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk.
2019-10-24 Jakub Jelinek <ja...@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.
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.
--- gcc/gimplify.h.jj 2019-10-18 00:16:11.196526233 +0200
+++ gcc/gimplify.h 2019-10-24 18:52:28.897994585 +0200
@@ -75,6 +75,8 @@ 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);
+
extern void gimplify_type_sizes (tree, gimple_seq *);
extern void gimplify_one_sizepos (tree *, gimple_seq *);
extern gbind *gimplify_body (tree, bool);
--- gcc/gimplify.c.jj 2019-10-18 00:16:11.149526930 +0200
+++ gcc/gimplify.c 2019-10-24 20:43:05.097408585 +0200
@@ -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
/* 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,
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
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
--- gcc/omp-general.h.jj 2019-10-11 14:10:44.988386966 +0200
+++ gcc/omp-general.h 2019-10-24 19:16:33.579877753 +0200
@@ -84,6 +84,9 @@ extern void omp_extract_for_data (gomp_f
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);
--- gcc/omp-general.c.jj 2019-10-11 14:10:44.988386966 +0200
+++ gcc/omp-general.c 2019-10-24 20:29:42.650690218 +0200
@@ -35,6 +35,11 @@ along with GCC; see the file COPYING3.
#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. */
--- gcc/c-family/c-common.h.jj 2019-10-24 14:46:34.919752034 +0200
+++ gcc/c-family/c-common.h 2019-10-24 14:55:01.431954218 +0200
@@ -1193,7 +1193,6 @@ extern enum omp_clause_default_kind c_om
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
--- gcc/c-family/c-omp.c.jj 2019-10-12 10:26:18.074940950 +0200
+++ gcc/c-family/c-omp.c 2019-10-24 14:59:43.047620888 +0200
@@ -34,9 +34,6 @@ along with GCC; see the file COPYING3.
#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 l
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;
-}
--- gcc/c-family/c-attribs.c.jj 2019-10-12 10:26:18.100940560 +0200
+++ gcc/c-family/c-attribs.c 2019-10-24 20:27:32.010689694 +0200
@@ -456,6 +456,12 @@ const struct attribute_spec c_common_att
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 },
--- gcc/c/c-parser.c.jj 2019-10-15 18:34:29.658082081 +0200
+++ gcc/c/c-parser.c 2019-10-24 14:58:26.733795158 +0200
@@ -19489,7 +19489,7 @@ c_finish_omp_declare_variant (c_parser *
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"),
--- gcc/c/c-decl.c.jj 2019-10-11 20:48:24.446177211 +0200
+++ gcc/c/c-decl.c 2019-10-24 15:52:09.209316303 +0200
@@ -4832,8 +4832,12 @@ c_decl_attributes (tree *node, tree attr
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
--- gcc/cp/decl2.c.jj 2019-10-12 10:21:30.280262887 +0200
+++ gcc/cp/decl2.c 2019-10-24 15:53:17.661270531 +0200
@@ -1555,8 +1555,12 @@ cplus_decl_attributes (tree *decl, tree
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)
--- gcc/testsuite/c-c++-common/gomp/declare-variant-8.c.jj 2019-10-24
20:25:52.623210839 +0200
+++ gcc/testsuite/c-c++-common/gomp/declare-variant-8.c 2019-10-24
20:44:48.070832547 +0200
@@ -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"
} } */
+}
Jakub
