OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

gimplify.h (omp_construct_selector_matches): Declare.

Browse Source 
* 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
This commit is contained in:
Jakub Jelinek 2019-10-25 00:29:09 +02:00 committed by Jakub Jelinek
parent f8cb8bcde1
commit 135df52cc3
16 changed files with 585 additions and 194 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
gcc/ChangeLog
View File

@ -1,5 +1,17 @@
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.

7
gcc/c-family/ChangeLog
View File

@ -1,3 +1,10 @@
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.

6
gcc/c-family/c-attribs.c
View File

@ -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 },

1
gcc/c-family/c-common.h
View File

@ -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

188
gcc/c-family/c-omp.c
View File

@ -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 %<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;
}

9
gcc/c/ChangeLog
View File

@ -1,3 +1,12 @@
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

8
gcc/c/c-decl.c
View File

@ -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

2
gcc/c/c-parser.c
View File

@ -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"),

4
gcc/cp/ChangeLog
View File

@ -1,5 +1,9 @@
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.

8
gcc/cp/decl2.c
View File

@ -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)

102
gcc/gimplify.c
View File

@ -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

2
gcc/gimplify.h
View File

@ -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);

298
gcc/omp-general.c
View File

@ -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. */

3
gcc/omp-general.h
View File

@ -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);

4
gcc/testsuite/ChangeLog
View File

@ -1,3 +1,7 @@
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

125
gcc/testsuite/c-c++-common/gomp/declare-variant-8.c Normal file
View File

@ -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" } } */
}