procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

OpenMP: Fortran support for metadirectives and dynamic selectors

Browse source 
gcc/fortran/ChangeLog
	PR middle-end/112779
	PR middle-end/113904
	* decl.cc (gfc_match_end): Handle COMP_OMP_BEGIN_METADIRECTIVE and
	COMP_OMP_METADIRECTIVE.
	* dump-parse-tree.cc (show_omp_node): Handle EXEC_OMP_METADIRECTIVE.
	(show_code_node): Likewise.
	* gfortran.h (enum gfc_statement): Add ST_OMP_METADIRECTIVE,
	ST_OMP_BEGIN_METADIRECTIVE, and ST_OMP_END_METADIRECTIVE.
	(struct gfc_omp_clauses): Rename target_first_st_is_teams to
	target_first_st_is_teams_or_meta.
	(struct gfc_omp_variant): New.
	(gfc_get_omp_variant): New.
	(struct gfc_st_label): Add omp_region field.
	(enum gfc_exec_op): Add EXEC_OMP_METADIRECTIVE.
	(struct gfc_code): Add omp_variants fields.
	(gfc_free_omp_variants): Declare.
	(match_omp_directive): Declare.
	(is_omp_declarative_stmt): Declare.
	* io.cc (format_asterisk): Adjust initializer.
	* match.h (gfc_match_omp_begin_metadirective): Declare.
	(gfc_match_omp_metadirective): Declare.
	* openmp.cc (gfc_omp_directives): Uncomment metadirective.
	(gfc_match_omp_eos): Adjust to match context selectors.
	(gfc_free_omp_variants): New.
	(gfc_match_omp_clauses): Remove context_selector parameter and adjust
	to use gfc_match_omp_eos instead.
	(match_omp): Adjust call to gfc_match_omp_clauses.
	(gfc_match_omp_context_selector): Add metadirective_p parameter and
	adjust error-checking.  Adjust matching of simd clauses.
	(gfc_match_omp_context_selector_specification): Adjust parameters
	so it can be used for metadirective as well as declare variant.
	(match_omp_metadirective): New.
	(gfc_match_omp_begin_metadirective): New.
	(gfc_match_omp_metadirective): New.
	(resolve_omp_metadirective): New.
	(resolve_omp_target): Handle metadirectives.
	(gfc_resolve_omp_directive): Handle EXEC_OMP_METADIRECTIVE.
	* parse.cc (gfc_matching_omp_context_selector): New.
	(gfc_in_omp_metadirective_body): New.
	(gfc_omp_region_count): New.
	(decode_omp_directive): Handle ST_OMP_BEGIN_METADIRECTIVE and
	ST_OMP_METADIRECTIVE.
	(match_omp_directive): New.
	(case_omp_structured_block): Define.
	(case_omp_do): Define.
	(gfc_ascii_statement): Handle ST_OMP_BEGIN_METADIRECTIVE,
	ST_OMP_END_METADIRECTIVE, and ST_OMP_METADIRECTIVE.
	(accept_statement):  Handle ST_OMP_METADIRECTIVE and
	ST_OMP_BEGIN_METADIRECTIVE.
	(gfc_omp_end_stmt): New, split from...
	(parse_omp_do): ...here, and...
	(parse_omp_structured_block): ...here.  Handle metadirectives,
	plus "allocate", "atomic", and "dispatch" which were missing.
	(parse_omp_oacc_atomic): Handle "end metadirective".
	(parse_openmp_allocate_block): Likewise.
	(parse_omp_dispatch): Likewise.
	(parse_omp_metadirective_body): New.
	(parse_executable): Handle metadirective.  Use new case macros
	defined above.
	(gfc_parse_file): Initialize metadirective state.
	(is_omp_declarative_stmt): New.
	* parse.h (enum gfc_compile_state): Add COMP_OMP_METADIRECTIVE
	and COMP_OMP_BEGIN_METADIRECTIVE.
	(gfc_omp_end_stmt): Declare.
	(gfc_matching_omp_context_selector): Declare.
	(gfc_in_omp_metadirective_body): Declare.
	(gfc_omp_metadirective_region_count): Declare.
	* resolve.cc (gfc_resolve_code): Handle EXEC_OMP_METADIRECTIVE.
	* st.cc (gfc_free_statement): Likewise.
	* symbol.cc (compare_st_labels): Handle labels within a metadirective
	body.
	(gfc_get_st_label): Likewise.
	* trans-decl.cc (gfc_get_label_decl): Encode the metadirective region
	in the label_name.
	* trans-openmp.cc (gfc_trans_omp_directive): Handle
	EXEC_OMP_METADIRECTIVE.
	(gfc_trans_omp_set_selector): New, split/adapted from code....
	(gfc_trans_omp_declare_variant): ...here.
	(gfc_trans_omp_metadirective): New.
	* trans-stmt.h 	(gfc_trans_omp_metadirective): Declare.
	* trans.cc (trans_code): Handle EXEC_OMP_METADIRECTIVE.

gcc/testsuite/ChangeLog
	PR middle-end/112779
	PR middle-end/113904
	* gfortran.dg/gomp/metadirective-1.f90: New.
	* gfortran.dg/gomp/metadirective-10.f90: New.
	* gfortran.dg/gomp/metadirective-11.f90: New.
	* gfortran.dg/gomp/metadirective-12.f90: New.
	* gfortran.dg/gomp/metadirective-13.f90: New.
	* gfortran.dg/gomp/metadirective-2.f90: New.
	* gfortran.dg/gomp/metadirective-3.f90: New.
	* gfortran.dg/gomp/metadirective-4.f90: New.
	* gfortran.dg/gomp/metadirective-5.f90: New.
	* gfortran.dg/gomp/metadirective-6.f90: New.
	* gfortran.dg/gomp/metadirective-7.f90: New.
	* gfortran.dg/gomp/metadirective-8.f90: New.
	* gfortran.dg/gomp/metadirective-9.f90: New.
	* gfortran.dg/gomp/metadirective-construct.f90: New.
	* gfortran.dg/gomp/metadirective-no-score.f90: New.
	* gfortran.dg/gomp/pure-1.f90 (func_metadirective): New.
	(func_metadirective_2): New.
	(func_metadirective_3): New.
	* gfortran.dg/gomp/pure-2.f90 (func_metadirective): Delete.

libgomp/ChangeLog
	PR middle-end/112779
	PR middle-end/113904
	* testsuite/libgomp.fortran/metadirective-1.f90: New.
	* testsuite/libgomp.fortran/metadirective-2.f90: New.
	* testsuite/libgomp.fortran/metadirective-3.f90: New.
	* testsuite/libgomp.fortran/metadirective-4.f90: New.
	* testsuite/libgomp.fortran/metadirective-5.f90: New.
	* testsuite/libgomp.fortran/metadirective-6.f90: New.

Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
Co-Authored-By: Paul-Antoine Arras <pa@codesourcery.com>
This commit is contained in:
Sandra Loosemore 2025-01-30 17:03:06 +00:00
parent 6a6df260c7
commit 8fbccdb342
38 changed files with 2101 additions and 370 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

29
gcc/fortran/decl.cc
View file

@ -8457,6 +8457,7 @@ gfc_match_end (gfc_statement *st)
case COMP_CONTAINS:
case COMP_DERIVED_CONTAINS:
case COMP_OMP_BEGIN_METADIRECTIVE:
state = gfc_state_stack->previous->state;
block_name = gfc_state_stack->previous->sym == NULL
? NULL : gfc_state_stack->previous->sym->name;
@ -8464,6 +8465,28 @@ gfc_match_end (gfc_statement *st)
&& gfc_state_stack->previous->sym->abr_modproc_decl;
break;
case COMP_OMP_METADIRECTIVE:
{
/* Metadirectives can be nested, so we need to drill down to the
first state that is not COMP_OMP_METADIRECTIVE. */
gfc_state_data *state_data = gfc_state_stack;
do
{
state_data = state_data->previous;
state = state_data->state;
block_name = (state_data->sym == NULL
? NULL : state_data->sym->name);
abbreviated_modproc_decl = (state_data->sym
&& state_data->sym->abr_modproc_decl);
}
while (state == COMP_OMP_METADIRECTIVE);
if (block_name && startswith (block_name, "block@"))
block_name = NULL;
}
break;
default:
break;
}
@ -8609,6 +8632,12 @@ gfc_match_end (gfc_statement *st)
gfc_free_enum_history ();
break;
case COMP_OMP_BEGIN_METADIRECTIVE:
*st = ST_OMP_END_METADIRECTIVE;
target = " metadirective";
eos_ok = 0;
break;
default:
gfc_error ("Unexpected END statement at %C");
goto cleanup;

20
gcc/fortran/dump-parse-tree.cc
View file

@ -2377,6 +2377,7 @@ show_omp_node (int level, gfc_code *c)
case EXEC_OMP_MASTER: name = "MASTER"; break;
case EXEC_OMP_MASTER_TASKLOOP: name = "MASTER TASKLOOP"; break;
case EXEC_OMP_MASTER_TASKLOOP_SIMD: name = "MASTER TASKLOOP SIMD"; break;
case EXEC_OMP_METADIRECTIVE: name = "METADIRECTIVE"; break;
case EXEC_OMP_ORDERED: name = "ORDERED"; break;
case EXEC_OMP_DEPOBJ: name = "DEPOBJ"; break;
case EXEC_OMP_PARALLEL: name = "PARALLEL"; break;
@ -2581,6 +2582,24 @@ show_omp_node (int level, gfc_code *c)
d = d->block;
}
}
else if (c->op == EXEC_OMP_METADIRECTIVE)
{
gfc_omp_variant *variant = c->ext.omp_variants;
while (variant)
{
code_indent (level + 1, 0);
if (variant->selectors)
fputs ("WHEN ()\n", dumpfile);
else
fputs ("DEFAULT ()\n", dumpfile);
/* TODO: Print selector. */
show_code (level + 2, variant->code);
if (variant->next)
fputs ("\n", dumpfile);
variant = variant->next;
}
}
else
show_code (level + 1, c->block->next);
if (c->op == EXEC_OMP_ATOMIC)
@ -3821,6 +3840,7 @@ show_code_node (int level, gfc_code *c)
case EXEC_OMP_MASTER:
case EXEC_OMP_MASTER_TASKLOOP:
case EXEC_OMP_MASTER_TASKLOOP_SIMD:
case EXEC_OMP_METADIRECTIVE:
case EXEC_OMP_ORDERED:
case EXEC_OMP_PARALLEL:
case EXEC_OMP_PARALLEL_DO:

21
gcc/fortran/gfortran.h
View file

@ -318,6 +318,7 @@ enum gfc_statement
ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD, ST_OMP_MASKED_TASKLOOP,
ST_OMP_END_MASKED_TASKLOOP, ST_OMP_MASKED_TASKLOOP_SIMD,
ST_OMP_END_MASKED_TASKLOOP_SIMD, ST_OMP_SCOPE, ST_OMP_END_SCOPE,
ST_OMP_METADIRECTIVE, ST_OMP_BEGIN_METADIRECTIVE, ST_OMP_END_METADIRECTIVE,
ST_OMP_ERROR, ST_OMP_ASSUME, ST_OMP_END_ASSUME, ST_OMP_ASSUMES,
ST_OMP_ALLOCATE, ST_OMP_ALLOCATE_EXEC,
ST_OMP_ALLOCATORS, ST_OMP_END_ALLOCATORS,
@ -1634,7 +1635,7 @@ typedef struct gfc_omp_clauses
unsigned order_unconstrained:1, order_reproducible:1, capture:1;
unsigned grainsize_strict:1, num_tasks_strict:1, compare:1, weak:1;
unsigned non_rectangular:1, order_concurrent:1;
unsigned contains_teams_construct:1, target_first_st_is_teams:1;
unsigned contains_teams_construct:1, target_first_st_is_teams_or_meta:1;
unsigned contained_in_target_construct:1, indirect:1;
unsigned full:1, erroneous:1;
ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3;
@ -1757,6 +1758,17 @@ typedef struct gfc_omp_declare_variant
gfc_omp_declare_variant;
#define gfc_get_omp_declare_variant() XCNEW (gfc_omp_declare_variant)
typedef struct gfc_omp_variant
{
struct gfc_omp_variant *next;
locus where; /* Where the metadirective clause occurred. */
gfc_omp_set_selector *selectors;
enum gfc_statement stmt;
struct gfc_code *code;
} gfc_omp_variant;
#define gfc_get_omp_variant() XCNEW (gfc_omp_variant)
typedef struct gfc_omp_udr
{
@ -1805,6 +1817,7 @@ typedef struct gfc_st_label
locus where;
gfc_namespace *ns;
int omp_region;
}
gfc_st_label;
@ -3108,7 +3121,7 @@ enum gfc_exec_op
EXEC_OMP_TARGET_TEAMS_LOOP, EXEC_OMP_MASKED, EXEC_OMP_PARALLEL_MASKED,
EXEC_OMP_PARALLEL_MASKED_TASKLOOP, EXEC_OMP_PARALLEL_MASKED_TASKLOOP_SIMD,
EXEC_OMP_MASKED_TASKLOOP, EXEC_OMP_MASKED_TASKLOOP_SIMD, EXEC_OMP_SCOPE,
EXEC_OMP_UNROLL, EXEC_OMP_TILE, EXEC_OMP_INTEROP,
EXEC_OMP_UNROLL, EXEC_OMP_TILE, EXEC_OMP_INTEROP, EXEC_OMP_METADIRECTIVE,
EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS, EXEC_OMP_DISPATCH
};
@ -3154,6 +3167,7 @@ typedef struct gfc_code
gfc_omp_clauses *omp_clauses;
const char *omp_name;
gfc_omp_namelist *omp_namelist;
gfc_omp_variant *omp_variants;
bool omp_bool;
int stop_code;
@ -3802,6 +3816,7 @@ void gfc_free_omp_declare_variant_list (gfc_omp_declare_variant *list);
void gfc_free_omp_declare_simd (gfc_omp_declare_simd *);
void gfc_free_omp_declare_simd_list (gfc_omp_declare_simd *);
void gfc_free_omp_udr (gfc_omp_udr *);
void gfc_free_omp_variants (gfc_omp_variant *);
gfc_omp_udr *gfc_omp_udr_find (gfc_symtree *, gfc_typespec *);
void gfc_resolve_omp_allocate (gfc_namespace *, gfc_omp_namelist *);
void gfc_resolve_omp_assumptions (gfc_omp_assumptions *);
@ -4089,6 +4104,8 @@ void debug (gfc_expr *);
bool gfc_parse_file (void);
void gfc_global_used (gfc_gsymbol *, locus *);
gfc_namespace* gfc_build_block_ns (gfc_namespace *);
gfc_statement match_omp_directive (void);
bool is_omp_declarative_stmt (gfc_statement);
/* dependency.cc */
int gfc_dep_compare_functions (gfc_expr *, gfc_expr *, bool);

2
gcc/fortran/io.cc
View file

@ -29,7 +29,7 @@ along with GCC; see the file COPYING3. If not see
gfc_st_label
format_asterisk = {0, NULL, NULL, -1, ST_LABEL_FORMAT, ST_LABEL_FORMAT, NULL,
0, {NULL, NULL}, NULL};
0, {NULL, NULL}, NULL, 0};
typedef struct
{

2
gcc/fortran/match.h
View file

@ -155,6 +155,7 @@ match gfc_match_omp_assume (void);
match gfc_match_omp_assumes (void);
match gfc_match_omp_atomic (void);
match gfc_match_omp_barrier (void);
match gfc_match_omp_begin_metadirective (void);
match gfc_match_omp_cancel (void);
match gfc_match_omp_cancellation_point (void);
match gfc_match_omp_critical (void);
@ -180,6 +181,7 @@ match gfc_match_omp_masked_taskloop_simd (void);
match gfc_match_omp_master (void);
match gfc_match_omp_master_taskloop (void);
match gfc_match_omp_master_taskloop_simd (void);
match gfc_match_omp_metadirective (void);
match gfc_match_omp_nothing (void);
match gfc_match_omp_ordered (void);
match gfc_match_omp_ordered_depend (void);

312
gcc/fortran/openmp.cc
View file

@ -83,7 +83,7 @@ static const struct gfc_omp_directive gfc_omp_directives[] = {
{"interop", GFC_OMP_DIR_EXECUTABLE, ST_OMP_INTEROP},
{"loop", GFC_OMP_DIR_EXECUTABLE, ST_OMP_LOOP},
{"masked", GFC_OMP_DIR_EXECUTABLE, ST_OMP_MASKED},
/* {"metadirective", GFC_OMP_DIR_META, ST_OMP_METADIRECTIVE}, */
{"metadirective", GFC_OMP_DIR_META, ST_OMP_METADIRECTIVE},
/* Note: gfc_match_omp_nothing returns ST_NONE. */
{"nothing", GFC_OMP_DIR_UTILITY, ST_OMP_NOTHING},
/* Special case; for now map to the first one.
@ -116,7 +116,8 @@ static const struct gfc_omp_directive gfc_omp_directives[] = {
/* Match an end of OpenMP directive. End of OpenMP directive is optional
whitespace, followed by '\n' or comment '!'. */
whitespace, followed by '\n' or comment '!'. In the special case where a
context selector is being matched, match against ')' instead. */
static match
gfc_match_omp_eos (void)
@ -127,17 +128,25 @@ gfc_match_omp_eos (void)
old_loc = gfc_current_locus;
gfc_gobble_whitespace ();
c = gfc_next_ascii_char ();
switch (c)
if (gfc_matching_omp_context_selector)
{
case '!':
do
c = gfc_next_ascii_char ();
while (c != '\n');
/* Fall through */
if (gfc_peek_ascii_char () == ')')
return MATCH_YES;
}
else
{
c = gfc_next_ascii_char ();
switch (c)
{
case '!':
do
c = gfc_next_ascii_char ();
while (c != '\n');
/* Fall through */
case '\n':
return MATCH_YES;
case '\n':
return MATCH_YES;
}
}
gfc_current_locus = old_loc;
@ -349,6 +358,19 @@ gfc_free_omp_udr (gfc_omp_udr *omp_udr)
}
}
/* Free variants of an !$omp metadirective construct. */
void
gfc_free_omp_variants (gfc_omp_variant *variant)
{
while (variant)
{
gfc_omp_variant *next_variant = variant->next;
gfc_free_omp_set_selector_list (variant->selectors);
free (variant);
variant = next_variant;
}
}
static gfc_omp_udr *
gfc_find_omp_udr (gfc_namespace *ns, const char *name, gfc_typespec *ts)
@ -2321,8 +2343,7 @@ gfc_match_dupl_atomic (bool not_dupl, const char *name)
static match
gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
bool first = true, bool needs_space = true,
bool openacc = false, bool context_selector = false,
bool openmp_target = false)
bool openacc = false, bool openmp_target = false)
{
bool error = false;
gfc_omp_clauses *c = gfc_get_omp_clauses ();
@ -4384,9 +4405,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
end:
if (error
|| (context_selector && gfc_peek_ascii_char () != ')')
|| (!context_selector && gfc_match_omp_eos () != MATCH_YES))
if (error || gfc_match_omp_eos () != MATCH_YES)
{
if (!gfc_error_flag_test ())
gfc_error ("Failed to match clause at %C");
@ -5100,7 +5119,7 @@ static match
match_omp (gfc_exec_op op, const omp_mask mask)
{
gfc_omp_clauses *c;
if (gfc_match_omp_clauses (&c, mask, true, true, false, false,
if (gfc_match_omp_clauses (&c, mask, true, true, false,
op == EXEC_OMP_TARGET) != MATCH_YES)
return MATCH_ERROR;
new_st.op = op;
@ -6295,7 +6314,8 @@ gfc_match_omp_interop (void)
score(score-expression) */
match
gfc_match_omp_context_selector (gfc_omp_set_selector *oss)
gfc_match_omp_context_selector (gfc_omp_set_selector *oss,
bool metadirective_p)
{
do
{
@ -6455,14 +6475,31 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss)
|| (property_kind == OMP_TRAIT_PROPERTY_DEV_NUM_EXPR
&& otp->expr->ts.type != BT_INTEGER)
|| otp->expr->rank != 0
|| otp->expr->expr_type != EXPR_CONSTANT)
|| (!metadirective_p
&& otp->expr->expr_type != EXPR_CONSTANT))
{
if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR)
gfc_error ("property must be a constant logical expression "
"at %C");
if (metadirective_p)
{
if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR)
gfc_error ("property must be a "
"logical expression at %L",
&otp->expr->where);
else
gfc_error ("property must be an "
"integer expression at %L",
&otp->expr->where);
}
else
gfc_error ("property must be a constant integer expression "
"at %C");
{
if (property_kind == OMP_TRAIT_PROPERTY_BOOL_EXPR)
gfc_error ("property must be a constant "
"logical expression at %L",
&otp->expr->where);
else
gfc_error ("property must be a constant "
"integer expression at %L",
&otp->expr->where);
}
return MATCH_ERROR;
}
/* Device number must be conforming, which includes
@ -6482,14 +6519,17 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss)
{
if (os->code == OMP_TRAIT_CONSTRUCT_SIMD)
{
gfc_matching_omp_context_selector = true;
if (gfc_match_omp_clauses (&otp->clauses,
OMP_DECLARE_SIMD_CLAUSES,
true, false, false, true)
true, false, false)
!= MATCH_YES)
{
gfc_matching_omp_context_selector = false;
gfc_error ("expected simd clause at %C");
return MATCH_ERROR;
}
gfc_matching_omp_context_selector = false;
}
else if (os->code == OMP_TRAIT_IMPLEMENTATION_REQUIRES)
{
@ -6546,7 +6586,8 @@ gfc_match_omp_context_selector (gfc_omp_set_selector *oss)
user */
match
gfc_match_omp_context_selector_specification (gfc_omp_declare_variant *odv)
gfc_match_omp_context_selector_specification (gfc_omp_set_selector **oss_head,
bool metadirective_p)
{
do
{
@ -6579,11 +6620,11 @@ gfc_match_omp_context_selector_specification (gfc_omp_declare_variant *odv)
}
gfc_omp_set_selector *oss = gfc_get_omp_set_selector ();
oss->next = odv->set_selectors;
oss->next = *oss_head;
oss->code = set;
odv->set_selectors = oss;
*oss_head = oss;
if (gfc_match_omp_context_selector (oss) != MATCH_YES)
if (gfc_match_omp_context_selector (oss, metadirective_p) != MATCH_YES)
return MATCH_ERROR;
m = gfc_match (" }");
@ -6714,7 +6755,8 @@ gfc_match_omp_declare_variant (void)
return MATCH_ERROR;
}
has_match = true;
if (gfc_match_omp_context_selector_specification (odv)
if (gfc_match_omp_context_selector_specification (&odv->set_selectors,
false)
!= MATCH_YES)
return MATCH_ERROR;
if (gfc_match (" )") != MATCH_YES)
@ -6831,6 +6873,167 @@ gfc_match_omp_declare_variant (void)
}
static match
match_omp_metadirective (bool begin_p)
{
locus old_loc = gfc_current_locus;
gfc_omp_variant *variants_head;
gfc_omp_variant **next_variant = &variants_head;
bool default_seen = false;
/* Parse the context selectors. */
for (;;)
{
bool default_p = false;
gfc_omp_set_selector *selectors = NULL;
gfc_gobble_whitespace ();
if (gfc_match_eos () == MATCH_YES)
break;
gfc_match_char (',');
gfc_gobble_whitespace ();
locus variant_locus = gfc_current_locus;
if (gfc_match (" default ( ") == MATCH_YES)
default_p = true;
else if (gfc_match (" otherwise ( ") == MATCH_YES)
default_p = true;
else if (gfc_match (" when ( ") != MATCH_YES)
{
gfc_error ("expected %<when%>, %<otherwise%>, or %<default%> at %C");
gfc_current_locus = old_loc;
return MATCH_ERROR;
}
if (default_p && default_seen)
{
gfc_error ("too many %<otherwise%> or %<default%> clauses "
"in %<metadirective%> at %C");
gfc_current_locus = old_loc;
return MATCH_ERROR;
}
else if (default_seen)
{
gfc_error ("%<otherwise%> or %<default%> clause "
"must appear last in %<metadirective%> at %C");
gfc_current_locus = old_loc;
return MATCH_ERROR;
}
if (!default_p)
{
if (gfc_match_omp_context_selector_specification (&selectors, true)
!= MATCH_YES)
return MATCH_ERROR;
if (gfc_match (" : ") != MATCH_YES)
{
gfc_error ("expected %<:%> at %C");
gfc_current_locus = old_loc;
return MATCH_ERROR;
}
gfc_commit_symbols ();
}
gfc_matching_omp_context_selector = true;
gfc_statement directive = match_omp_directive ();
gfc_matching_omp_context_selector = false;
if (is_omp_declarative_stmt (directive))
sorry ("declarative directive variants are not supported");
if (gfc_error_flag_test ())
{
gfc_current_locus = old_loc;
return MATCH_ERROR;
}
if (gfc_match (" )") != MATCH_YES)
{
gfc_error ("Expected %<)%> at %C");
gfc_current_locus = old_loc;
return MATCH_ERROR;
}
gfc_commit_symbols ();
if (begin_p
&& directive != ST_NONE
&& gfc_omp_end_stmt (directive) == ST_NONE)
{
gfc_error ("variant directive used in OMP BEGIN METADIRECTIVE "
"at %C must have a corresponding end directive");
gfc_current_locus = old_loc;
return MATCH_ERROR;
}
if (default_p)
default_seen = true;
gfc_omp_variant *omv = gfc_get_omp_variant ();
omv->selectors = selectors;
omv->stmt = directive;
omv->where = variant_locus;
if (directive == ST_NONE)
{
/* The directive was a 'nothing' directive. */
omv->code = gfc_get_code (EXEC_CONTINUE);
omv->code->ext.omp_clauses = NULL;
}
else
{
omv->code = gfc_get_code (new_st.op);
omv->code->ext.omp_clauses = new_st.ext.omp_clauses;
/* Prevent the OpenMP clauses from being freed via NEW_ST. */
new_st.ext.omp_clauses = NULL;
}
*next_variant = omv;
next_variant = &omv->next;
}
if (gfc_match_omp_eos () != MATCH_YES)
{
gfc_error ("Unexpected junk after OMP METADIRECTIVE at %C");
gfc_current_locus = old_loc;
return MATCH_ERROR;
}
/* Add a 'default (nothing)' clause if no default is explicitly given. */
if (!default_seen)
{
gfc_omp_variant *omv = gfc_get_omp_variant ();
omv->stmt = ST_NONE;
omv->code = gfc_get_code (EXEC_CONTINUE);
omv->code->ext.omp_clauses = NULL;
omv->where = old_loc;
omv->selectors = NULL;
*next_variant = omv;
next_variant = &omv->next;
}
new_st.op = EXEC_OMP_METADIRECTIVE;
new_st.ext.omp_variants = variants_head;
return MATCH_YES;
}
match
gfc_match_omp_begin_metadirective (void)
{
return match_omp_metadirective (true);
}
match
gfc_match_omp_metadirective (void)
{
return match_omp_metadirective (false);
}
match
gfc_match_omp_threadprivate (void)
{
@ -11987,6 +12190,19 @@ resolve_omp_do (gfc_code *code)
non_generated_count);
}
static void
resolve_omp_metadirective (gfc_code *code, gfc_namespace *ns)
{
gfc_omp_variant *variant = code->ext.omp_variants;
while (variant)
{
gfc_code *variant_code = variant->code;
gfc_resolve_code (variant_code, ns);
variant = variant->next;
}
}
static gfc_statement
omp_code_to_statement (gfc_code *code)
@ -12538,13 +12754,32 @@ resolve_omp_target (gfc_code *code)
gfc_code *c = code->block->next;
if (c->op == EXEC_BLOCK)
c = c->ext.block.ns->code;
if (code->ext.omp_clauses->target_first_st_is_teams
&& ((GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL)
|| (c->op == EXEC_BLOCK
&& c->next
&& GFC_IS_TEAMS_CONSTRUCT (c->next->op)
&& c->next->next == NULL)))
return;
if (code->ext.omp_clauses->target_first_st_is_teams_or_meta)
{
if (c->op == EXEC_OMP_METADIRECTIVE)
{
struct gfc_omp_variant *mc
= c->ext.omp_variants;
/* All mc->(next...->)code should be identical with regards
to the diagnostic below. */
do
{
if (mc->stmt != ST_NONE
&& GFC_IS_TEAMS_CONSTRUCT (mc->code->op))
{
if (c->next == NULL && mc->code->next == NULL)
return;
c = mc->code;
break;
}
mc = mc->next;
}
while (mc);
}
else if (GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL)
return;
}
while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op))
c = c->next;
if (c)
@ -12714,6 +12949,9 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns)
resolve_omp_clauses (code, code->ext.omp_clauses, ns);
resolve_omp_dispatch (code);
break;
case EXEC_OMP_METADIRECTIVE:
resolve_omp_metadirective (code, ns);
break;
default:
break;
}

621
gcc/fortran/parse.cc
View file

@ -48,6 +48,16 @@ gfc_state_data *gfc_state_stack;
static bool last_was_use_stmt = false;
bool in_exec_part;
/* True when matching an OpenMP context selector. */
bool gfc_matching_omp_context_selector;
/* True when parsing the body of an OpenMP metadirective. */
bool gfc_in_omp_metadirective_body;
/* Each metadirective body in the translation unit is given a unique
number, used to ensure that labels in the body have unique names. */
int gfc_omp_metadirective_region_count;
/* TODO: Re-order functions to kill these forward decls. */
static void check_statement_label (gfc_statement);
static void undo_new_statement (void);
@ -993,6 +1003,12 @@ decode_omp_directive (void)
matcho ("assumes", gfc_match_omp_assumes, ST_OMP_ASSUMES);
matchs ("assume", gfc_match_omp_assume, ST_OMP_ASSUME);
break;
case 'b':
matcho ("begin metadirective", gfc_match_omp_begin_metadirective,
ST_OMP_BEGIN_METADIRECTIVE);
break;
case 'd':
matchds ("declare reduction", gfc_match_omp_declare_reduction,
ST_OMP_DECLARE_REDUCTION);
@ -1005,11 +1021,19 @@ decode_omp_directive (void)
break;
case 'e':
matchs ("end assume", gfc_match_omp_eos_error, ST_OMP_END_ASSUME);
matcho ("end metadirective", gfc_match_omp_eos_error,
ST_OMP_END_METADIRECTIVE);
matchs ("end simd", gfc_match_omp_eos_error, ST_OMP_END_SIMD);
matchs ("end tile", gfc_match_omp_eos_error, ST_OMP_END_TILE);
matchs ("end unroll", gfc_match_omp_eos_error, ST_OMP_END_UNROLL);
matcho ("error", gfc_match_omp_error, ST_OMP_ERROR);
break;
case 'm':
matcho ("metadirective", gfc_match_omp_metadirective,
ST_OMP_METADIRECTIVE);
break;
case 'n':
matcho ("nothing", gfc_match_omp_nothing, ST_NONE);
break;
@ -1309,6 +1333,10 @@ decode_omp_directive (void)
gfc_error_now ("Unclassifiable OpenMP directive at %C");
}
/* If parsing a metadirective, let the caller deal with the cleanup. */
if (gfc_matching_omp_context_selector)
return ST_NONE;
reject_statement ();
gfc_error_recovery ();
@ -1430,6 +1458,12 @@ decode_omp_directive (void)
return ST_GET_FCN_CHARACTERISTICS;
}
gfc_statement
match_omp_directive (void)
{
return decode_omp_directive ();
}
static gfc_statement
decode_gcc_attribute (void)
{
@ -1955,6 +1989,44 @@ next_statement (void)
case ST_OMP_DECLARE_VARIANT: case ST_OMP_ALLOCATE: case ST_OMP_ASSUMES: \
case ST_OMP_REQUIRES: case ST_OACC_ROUTINE: case ST_OACC_DECLARE
/* OpenMP statements that are followed by a structured block. */
#define case_omp_structured_block case ST_OMP_ASSUME: case ST_OMP_PARALLEL: \
case ST_OMP_PARALLEL_MASKED: case ST_OMP_PARALLEL_MASTER: \
case ST_OMP_PARALLEL_SECTIONS: case ST_OMP_ORDERED: \
case ST_OMP_CRITICAL: case ST_OMP_MASKED: case ST_OMP_MASTER: \
case ST_OMP_SCOPE: case ST_OMP_SECTIONS: case ST_OMP_SINGLE: \
case ST_OMP_TARGET: case ST_OMP_TARGET_DATA: case ST_OMP_TARGET_PARALLEL: \
case ST_OMP_TARGET_TEAMS: case ST_OMP_TEAMS: case ST_OMP_TASK: \
case ST_OMP_TASKGROUP: \
case ST_OMP_WORKSHARE: case ST_OMP_PARALLEL_WORKSHARE
/* OpenMP statements that are followed by a do loop. */
#define case_omp_do case ST_OMP_DISTRIBUTE: \
case ST_OMP_DISTRIBUTE_PARALLEL_DO: \
case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: case ST_OMP_DISTRIBUTE_SIMD: \
case ST_OMP_DO: case ST_OMP_DO_SIMD: case ST_OMP_LOOP: \
case ST_OMP_PARALLEL_DO: case ST_OMP_PARALLEL_DO_SIMD: \
case ST_OMP_PARALLEL_LOOP: case ST_OMP_PARALLEL_MASKED_TASKLOOP: \
case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD: \
case ST_OMP_PARALLEL_MASTER_TASKLOOP: \
case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD: \
case ST_OMP_MASKED_TASKLOOP: case ST_OMP_MASKED_TASKLOOP_SIMD: \
case ST_OMP_MASTER_TASKLOOP: case ST_OMP_MASTER_TASKLOOP_SIMD: \
case ST_OMP_SIMD: \
case ST_OMP_TARGET_PARALLEL_DO: case ST_OMP_TARGET_PARALLEL_DO_SIMD: \
case ST_OMP_TARGET_PARALLEL_LOOP: case ST_OMP_TARGET_SIMD: \
case ST_OMP_TARGET_TEAMS_DISTRIBUTE: \
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: \
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: \
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: case ST_OMP_TARGET_TEAMS_LOOP: \
case ST_OMP_TASKLOOP: case ST_OMP_TASKLOOP_SIMD: \
case ST_OMP_TEAMS_DISTRIBUTE: case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: \
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: \
case ST_OMP_TEAMS_DISTRIBUTE_SIMD: case ST_OMP_TEAMS_LOOP: \
case ST_OMP_TILE: case ST_OMP_UNROLL
/* Block end statements. Errors associated with interchanging these
are detected in gfc_match_end(). */
@ -2592,6 +2664,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel)
case ST_OMP_BARRIER:
p = "!$OMP BARRIER";
break;
case ST_OMP_BEGIN_METADIRECTIVE:
p = "!$OMP BEGIN METADIRECTIVE";
break;
case ST_OMP_CANCEL:
p = "!$OMP CANCEL";
break;
@ -2697,6 +2772,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel)
case ST_OMP_END_MASTER_TASKLOOP_SIMD:
p = "!$OMP END MASTER TASKLOOP SIMD";
break;
case ST_OMP_END_METADIRECTIVE:
p = "!$OMP END METADIRECTIVE";
break;
case ST_OMP_END_ORDERED:
p = "!$OMP END ORDERED";
break;
@ -2850,6 +2928,9 @@ gfc_ascii_statement (gfc_statement st, bool strip_sentinel)
case ST_OMP_MASTER_TASKLOOP_SIMD:
p = "!$OMP MASTER TASKLOOP SIMD";
break;
case ST_OMP_METADIRECTIVE:
p = "!$OMP METADIRECTIVE";
break;
case ST_OMP_ORDERED:
case ST_OMP_ORDERED_DEPEND:
p = "!$OMP ORDERED";
@ -3116,6 +3197,8 @@ accept_statement (gfc_statement st)
break;
case ST_ENTRY:
case ST_OMP_METADIRECTIVE:
case ST_OMP_BEGIN_METADIRECTIVE:
case_executable:
case_exec_markers:
add_statement ();
@ -5511,6 +5594,150 @@ loop:
accept_statement (st);
}
/* Get the corresponding ending statement type for the OpenMP directive
OMP_ST. If it does not have one, return ST_NONE. */
gfc_statement
gfc_omp_end_stmt (gfc_statement omp_st,
bool omp_do_p, bool omp_structured_p)
{
if (omp_do_p)
{
switch (omp_st)
{
case ST_OMP_DISTRIBUTE: return ST_OMP_END_DISTRIBUTE;
case ST_OMP_DISTRIBUTE_PARALLEL_DO:
return ST_OMP_END_DISTRIBUTE_PARALLEL_DO;
case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
return ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD;
case ST_OMP_DISTRIBUTE_SIMD:
return ST_OMP_END_DISTRIBUTE_SIMD;
case ST_OMP_DO: return ST_OMP_END_DO;
case ST_OMP_DO_SIMD: return ST_OMP_END_DO_SIMD;
case ST_OMP_LOOP: return ST_OMP_END_LOOP;
case ST_OMP_PARALLEL_DO: return ST_OMP_END_PARALLEL_DO;
case ST_OMP_PARALLEL_DO_SIMD:
return ST_OMP_END_PARALLEL_DO_SIMD;
case ST_OMP_PARALLEL_LOOP:
return ST_OMP_END_PARALLEL_LOOP;
case ST_OMP_SIMD: return ST_OMP_END_SIMD;
case ST_OMP_TARGET_PARALLEL_DO:
return ST_OMP_END_TARGET_PARALLEL_DO;
case ST_OMP_TARGET_PARALLEL_DO_SIMD:
return ST_OMP_END_TARGET_PARALLEL_DO_SIMD;
case ST_OMP_TARGET_PARALLEL_LOOP:
return ST_OMP_END_TARGET_PARALLEL_LOOP;
case ST_OMP_TARGET_SIMD: return ST_OMP_END_TARGET_SIMD;
case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE;
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO;
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
return ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD;
case ST_OMP_TARGET_TEAMS_LOOP:
return ST_OMP_END_TARGET_TEAMS_LOOP;
case ST_OMP_TASKLOOP: return ST_OMP_END_TASKLOOP;
case ST_OMP_TASKLOOP_SIMD: return ST_OMP_END_TASKLOOP_SIMD;
case ST_OMP_MASKED_TASKLOOP: return ST_OMP_END_MASKED_TASKLOOP;
case ST_OMP_MASKED_TASKLOOP_SIMD:
return ST_OMP_END_MASKED_TASKLOOP_SIMD;
case ST_OMP_MASTER_TASKLOOP: return ST_OMP_END_MASTER_TASKLOOP;
case ST_OMP_MASTER_TASKLOOP_SIMD:
return ST_OMP_END_MASTER_TASKLOOP_SIMD;
case ST_OMP_PARALLEL_MASKED_TASKLOOP:
return ST_OMP_END_PARALLEL_MASKED_TASKLOOP;
case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD:
return ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD;
case ST_OMP_PARALLEL_MASTER_TASKLOOP:
return ST_OMP_END_PARALLEL_MASTER_TASKLOOP;
case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD:
return ST_OMP_END_PARALLEL_MASTER_TASKLOOP_SIMD;
case ST_OMP_TEAMS_DISTRIBUTE:
return ST_OMP_END_TEAMS_DISTRIBUTE;
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
return ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO;
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
return ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
return ST_OMP_END_TEAMS_DISTRIBUTE_SIMD;
case ST_OMP_TEAMS_LOOP:
return ST_OMP_END_TEAMS_LOOP;
case ST_OMP_TILE:
return ST_OMP_END_TILE;
case ST_OMP_UNROLL:
return ST_OMP_END_UNROLL;
default:
break;
}
}
if (omp_structured_p)
{
switch (omp_st)
{
case ST_OMP_ALLOCATORS:
return ST_OMP_END_ALLOCATORS;
case ST_OMP_ASSUME:
return ST_OMP_END_ASSUME;
case ST_OMP_ATOMIC:
return ST_OMP_END_ATOMIC;
case ST_OMP_DISPATCH:
return ST_OMP_END_DISPATCH;
case ST_OMP_PARALLEL:
return ST_OMP_END_PARALLEL;
case ST_OMP_PARALLEL_MASKED:
return ST_OMP_END_PARALLEL_MASKED;
case ST_OMP_PARALLEL_MASTER:
return ST_OMP_END_PARALLEL_MASTER;
case ST_OMP_PARALLEL_SECTIONS:
return ST_OMP_END_PARALLEL_SECTIONS;
case ST_OMP_SCOPE:
return ST_OMP_END_SCOPE;
case ST_OMP_SECTIONS:
return ST_OMP_END_SECTIONS;
case ST_OMP_ORDERED:
return ST_OMP_END_ORDERED;
case ST_OMP_CRITICAL:
return ST_OMP_END_CRITICAL;
case ST_OMP_MASKED:
return ST_OMP_END_MASKED;
case ST_OMP_MASTER:
return ST_OMP_END_MASTER;
case ST_OMP_SINGLE:
return ST_OMP_END_SINGLE;
case ST_OMP_TARGET:
return ST_OMP_END_TARGET;
case ST_OMP_TARGET_DATA:
return ST_OMP_END_TARGET_DATA;
case ST_OMP_TARGET_PARALLEL:
return ST_OMP_END_TARGET_PARALLEL;
case ST_OMP_TARGET_TEAMS:
return ST_OMP_END_TARGET_TEAMS;
case ST_OMP_TASK:
return ST_OMP_END_TASK;
case ST_OMP_TASKGROUP:
return ST_OMP_END_TASKGROUP;
case ST_OMP_TEAMS:
return ST_OMP_END_TEAMS;
case ST_OMP_TEAMS_DISTRIBUTE:
return ST_OMP_END_TEAMS_DISTRIBUTE;
case ST_OMP_DISTRIBUTE:
return ST_OMP_END_DISTRIBUTE;
case ST_OMP_WORKSHARE:
return ST_OMP_END_WORKSHARE;
case ST_OMP_PARALLEL_WORKSHARE:
return ST_OMP_END_PARALLEL_WORKSHARE;
case ST_OMP_BEGIN_METADIRECTIVE:
return ST_OMP_END_METADIRECTIVE;
default:
break;
}
}
return ST_NONE;
}
/* Parse the statements of OpenMP do/parallel do. */
@ -5571,94 +5798,16 @@ parse_omp_do (gfc_statement omp_st, int nested)
st = next_statement ();
do_end:
gfc_statement omp_end_st = ST_OMP_END_DO;
switch (omp_st)
{
case ST_OMP_DISTRIBUTE: omp_end_st = ST_OMP_END_DISTRIBUTE; break;
case ST_OMP_DISTRIBUTE_PARALLEL_DO:
omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO;
break;
case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD;
break;
case ST_OMP_DISTRIBUTE_SIMD:
omp_end_st = ST_OMP_END_DISTRIBUTE_SIMD;
break;
case ST_OMP_DO: omp_end_st = ST_OMP_END_DO; break;
case ST_OMP_DO_SIMD: omp_end_st = ST_OMP_END_DO_SIMD; break;
case ST_OMP_LOOP: omp_end_st = ST_OMP_END_LOOP; break;
case ST_OMP_PARALLEL_DO: omp_end_st = ST_OMP_END_PARALLEL_DO; break;
case ST_OMP_PARALLEL_DO_SIMD:
omp_end_st = ST_OMP_END_PARALLEL_DO_SIMD;
break;
case ST_OMP_PARALLEL_LOOP:
omp_end_st = ST_OMP_END_PARALLEL_LOOP;
break;
case ST_OMP_SIMD: omp_end_st = ST_OMP_END_SIMD; break;
case ST_OMP_TARGET_PARALLEL_DO:
omp_end_st = ST_OMP_END_TARGET_PARALLEL_DO;
break;
case ST_OMP_TARGET_PARALLEL_DO_SIMD:
omp_end_st = ST_OMP_END_TARGET_PARALLEL_DO_SIMD;
break;
case ST_OMP_TARGET_PARALLEL_LOOP:
omp_end_st = ST_OMP_END_TARGET_PARALLEL_LOOP;
break;
case ST_OMP_TARGET_SIMD: omp_end_st = ST_OMP_END_TARGET_SIMD; break;
case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE;
break;
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO;
break;
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
break;
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD;
break;
case ST_OMP_TARGET_TEAMS_LOOP:
omp_end_st = ST_OMP_END_TARGET_TEAMS_LOOP;
break;
case ST_OMP_TASKLOOP: omp_end_st = ST_OMP_END_TASKLOOP; break;
case ST_OMP_TASKLOOP_SIMD: omp_end_st = ST_OMP_END_TASKLOOP_SIMD; break;
case ST_OMP_MASKED_TASKLOOP: omp_end_st = ST_OMP_END_MASKED_TASKLOOP; break;
case ST_OMP_MASKED_TASKLOOP_SIMD:
omp_end_st = ST_OMP_END_MASKED_TASKLOOP_SIMD;
break;
case ST_OMP_MASTER_TASKLOOP: omp_end_st = ST_OMP_END_MASTER_TASKLOOP; break;
case ST_OMP_MASTER_TASKLOOP_SIMD:
omp_end_st = ST_OMP_END_MASTER_TASKLOOP_SIMD;
break;
case ST_OMP_PARALLEL_MASKED_TASKLOOP:
omp_end_st = ST_OMP_END_PARALLEL_MASKED_TASKLOOP;
break;
case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD:
omp_end_st = ST_OMP_END_PARALLEL_MASKED_TASKLOOP_SIMD;
break;
case ST_OMP_PARALLEL_MASTER_TASKLOOP:
omp_end_st = ST_OMP_END_PARALLEL_MASTER_TASKLOOP;
break;
case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD:
omp_end_st = ST_OMP_END_PARALLEL_MASTER_TASKLOOP_SIMD;
break;
case ST_OMP_TEAMS_DISTRIBUTE:
omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE;
break;
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO;
break;
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
break;
case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_SIMD;
break;
case ST_OMP_TEAMS_LOOP: omp_end_st = ST_OMP_END_TEAMS_LOOP; break;
case ST_OMP_TILE: omp_end_st = ST_OMP_END_TILE; break;
case ST_OMP_UNROLL: omp_end_st = ST_OMP_END_UNROLL; break;
default: gcc_unreachable ();
}
gfc_statement omp_end_st = gfc_omp_end_stmt (omp_st, true, false);
if (omp_st == ST_NONE)
gcc_unreachable ();
/* If handling a metadirective variant, treat 'omp end metadirective'
as the expected end statement for the current construct. */
if (st == ST_OMP_END_METADIRECTIVE
&& gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE)
st = omp_end_st;
if (st == omp_end_st)
{
if (new_st.op == EXEC_OMP_END_NOWAIT)
@ -5693,7 +5842,10 @@ parse_omp_oacc_atomic (bool omp_p)
if (omp_p)
{
st_atomic = ST_OMP_ATOMIC;
st_end_atomic = ST_OMP_END_ATOMIC;
if (gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE)
st_end_atomic = ST_OMP_END_METADIRECTIVE;
else
st_end_atomic = ST_OMP_END_ATOMIC;
}
else
{
@ -5944,7 +6096,10 @@ parse_openmp_allocate_block (gfc_statement omp_st)
accept_statement (st);
pop_state ();
st = next_statement ();
if (omp_st == ST_OMP_ALLOCATORS && st == ST_OMP_END_ALLOCATORS)
if (omp_st == ST_OMP_ALLOCATORS
&& (st == ST_OMP_END_ALLOCATORS
|| (st == ST_OMP_END_METADIRECTIVE
&& gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE)))
{
accept_statement (st);
st = next_statement ();
@ -5970,80 +6125,15 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
np->op = cp->op;
np->block = NULL;
switch (omp_st)
{
case ST_OMP_ASSUME:
omp_end_st = ST_OMP_END_ASSUME;
break;
case ST_OMP_PARALLEL:
omp_end_st = ST_OMP_END_PARALLEL;
break;
case ST_OMP_PARALLEL_MASKED:
omp_end_st = ST_OMP_END_PARALLEL_MASKED;
break;
case ST_OMP_PARALLEL_MASTER:
omp_end_st = ST_OMP_END_PARALLEL_MASTER;
break;
case ST_OMP_PARALLEL_SECTIONS:
omp_end_st = ST_OMP_END_PARALLEL_SECTIONS;
break;
case ST_OMP_SCOPE:
omp_end_st = ST_OMP_END_SCOPE;
break;
case ST_OMP_SECTIONS:
omp_end_st = ST_OMP_END_SECTIONS;
break;
case ST_OMP_ORDERED:
omp_end_st = ST_OMP_END_ORDERED;
break;
case ST_OMP_CRITICAL:
omp_end_st = ST_OMP_END_CRITICAL;
break;
case ST_OMP_MASKED:
omp_end_st = ST_OMP_END_MASKED;
break;
case ST_OMP_MASTER:
omp_end_st = ST_OMP_END_MASTER;
break;
case ST_OMP_SINGLE:
omp_end_st = ST_OMP_END_SINGLE;
break;
case ST_OMP_TARGET:
omp_end_st = ST_OMP_END_TARGET;
break;
case ST_OMP_TARGET_DATA:
omp_end_st = ST_OMP_END_TARGET_DATA;
break;
case ST_OMP_TARGET_PARALLEL:
omp_end_st = ST_OMP_END_TARGET_PARALLEL;
break;
case ST_OMP_TARGET_TEAMS:
omp_end_st = ST_OMP_END_TARGET_TEAMS;
break;
case ST_OMP_TASK:
omp_end_st = ST_OMP_END_TASK;
break;
case ST_OMP_TASKGROUP:
omp_end_st = ST_OMP_END_TASKGROUP;
break;
case ST_OMP_TEAMS:
omp_end_st = ST_OMP_END_TEAMS;
break;
case ST_OMP_TEAMS_DISTRIBUTE:
omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE;
break;
case ST_OMP_DISTRIBUTE:
omp_end_st = ST_OMP_END_DISTRIBUTE;
break;
case ST_OMP_WORKSHARE:
omp_end_st = ST_OMP_END_WORKSHARE;
break;
case ST_OMP_PARALLEL_WORKSHARE:
omp_end_st = ST_OMP_END_PARALLEL_WORKSHARE;
break;
default:
gcc_unreachable ();
}
omp_end_st = gfc_omp_end_stmt (omp_st, false, true);
if (omp_end_st == ST_NONE)
gcc_unreachable ();
/* If handling a metadirective variant, treat 'omp end metadirective'
as the expected end statement for the current construct. */
if (gfc_state_stack->previous != NULL
&& gfc_state_stack->previous->state == COMP_OMP_BEGIN_METADIRECTIVE)
omp_end_st = ST_OMP_END_METADIRECTIVE;
bool block_construct = false;
gfc_namespace *my_ns = NULL;
@ -6089,11 +6179,13 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
case ST_OMP_TEAMS_LOOP:
case ST_OMP_METADIRECTIVE:
case ST_OMP_BEGIN_METADIRECTIVE:
{
gfc_state_data *stk = gfc_state_stack->previous;
if (stk->state == COMP_OMP_STRICTLY_STRUCTURED_BLOCK)
stk = stk->previous;
stk->tail->ext.omp_clauses->target_first_st_is_teams = true;
stk->tail->ext.omp_clauses->target_first_st_is_teams_or_meta = true;
break;
}
default:
@ -6266,7 +6358,6 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
return st;
}
static gfc_statement
parse_omp_dispatch (void)
{
@ -6295,7 +6386,9 @@ parse_omp_dispatch (void)
}
pop_state ();
st = next_statement ();
if (st == ST_OMP_END_DISPATCH)
if (st == ST_OMP_END_DISPATCH
|| (st == ST_OMP_END_METADIRECTIVE
&& gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE))
{
if (cp->ext.omp_clauses->nowait && new_st.ext.omp_bool)
gfc_error_now ("Duplicated NOWAIT clause on !$OMP DISPATCH and !$OMP "
@ -6307,6 +6400,98 @@ parse_omp_dispatch (void)
return st;
}
static gfc_statement
parse_omp_metadirective_body (gfc_statement omp_st)
{
gfc_omp_variant *variant
= new_st.ext.omp_variants;
locus body_locus = gfc_current_locus;
accept_statement (omp_st);
gfc_statement next_st = ST_NONE;
while (variant)
{
gfc_current_locus = body_locus;
gfc_state_data s;
bool workshare_p
= (variant->stmt == ST_OMP_WORKSHARE
|| variant->stmt == ST_OMP_PARALLEL_WORKSHARE);
enum gfc_compile_state new_state
= (omp_st == ST_OMP_METADIRECTIVE
? COMP_OMP_METADIRECTIVE : COMP_OMP_BEGIN_METADIRECTIVE);
new_st = *variant->code;
push_state (&s, new_state, NULL);
gfc_statement st;
bool old_in_metadirective_body = gfc_in_omp_metadirective_body;
gfc_in_omp_metadirective_body = true;
gfc_omp_metadirective_region_count++;
switch (variant->stmt)
{
case_omp_structured_block:
st = parse_omp_structured_block (variant->stmt, workshare_p);
break;
case_omp_do:
st = parse_omp_do (variant->stmt, 0);
/* TODO: Does st == ST_IMPLIED_ENDDO need special handling? */
break;
case ST_OMP_ALLOCATORS:
st = parse_openmp_allocate_block (variant->stmt);
break;
case ST_OMP_ATOMIC:
st = parse_omp_oacc_atomic (true);
break;
case ST_OMP_DISPATCH:
st = parse_omp_dispatch ();
break;
default:
accept_statement (variant->stmt);
st = parse_executable (next_statement ());
break;
}
if (gfc_state_stack->state == COMP_OMP_METADIRECTIVE
&& startswith (gfc_ascii_statement (st), "!$OMP END "))
{
for (gfc_state_data *p = gfc_state_stack; p; p = p->previous)
if (p->state == COMP_OMP_STRUCTURED_BLOCK
|| p->state == COMP_OMP_BEGIN_METADIRECTIVE)
goto finish;
gfc_error ("Unexpected %s statement in OMP METADIRECTIVE "
"block at %C",
gfc_ascii_statement (st));
reject_statement ();
st = next_statement ();
}
finish:
gfc_in_omp_metadirective_body = old_in_metadirective_body;
if (gfc_state_stack->head)
*variant->code = *gfc_state_stack->head;
pop_state ();
gfc_commit_symbols ();
gfc_warning_check ();
if (variant->next)
gfc_clear_new_st ();
/* Sanity-check that each variant finishes parsing at the same place. */
if (next_st == ST_NONE)
next_st = st;
else
gcc_assert (st == next_st);
variant = variant->next;
}
return next_st;
}
/* Accept a series of executable statements. We return the first
statement that doesn't fit to the caller. Any block statements are
passed on to the correct handler, which usually passes the buck
@ -6316,6 +6501,7 @@ static gfc_statement
parse_executable (gfc_statement st)
{
int close_flag;
bool one_stmt_p = false;
in_exec_part = true;
if (st == ST_NONE)
@ -6323,6 +6509,12 @@ parse_executable (gfc_statement st)
for (;;)
{
/* Only parse one statement for the form of metadirective without
an explicit begin..end. */
if (gfc_state_stack->state == COMP_OMP_METADIRECTIVE && one_stmt_p)
return st;
one_stmt_p = true;
close_flag = check_do_closure ();
if (close_flag)
switch (st)
@ -6432,70 +6624,13 @@ parse_executable (gfc_statement st)
st = parse_openmp_allocate_block (st);
continue;
case ST_OMP_ASSUME:
case ST_OMP_PARALLEL:
case ST_OMP_PARALLEL_MASKED:
case ST_OMP_PARALLEL_MASTER:
case ST_OMP_PARALLEL_SECTIONS:
case ST_OMP_ORDERED:
case ST_OMP_CRITICAL:
case ST_OMP_MASKED:
case ST_OMP_MASTER:
case ST_OMP_SCOPE:
case ST_OMP_SECTIONS:
case ST_OMP_SINGLE:
case ST_OMP_TARGET:
case ST_OMP_TARGET_DATA:
case ST_OMP_TARGET_PARALLEL:
case ST_OMP_TARGET_TEAMS:
case ST_OMP_TEAMS:
case ST_OMP_TASK:
case ST_OMP_TASKGROUP:
st = parse_omp_structured_block (st, false);
case_omp_structured_block:
st = parse_omp_structured_block (st,
st == ST_OMP_WORKSHARE
|| st == ST_OMP_PARALLEL_WORKSHARE);
continue;
case ST_OMP_WORKSHARE:
case ST_OMP_PARALLEL_WORKSHARE:
st = parse_omp_structured_block (st, true);
continue;
case ST_OMP_DISTRIBUTE:
case ST_OMP_DISTRIBUTE_PARALLEL_DO:
case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
case ST_OMP_DISTRIBUTE_SIMD:
case ST_OMP_DO:
case ST_OMP_DO_SIMD:
case ST_OMP_LOOP:
case ST_OMP_PARALLEL_DO:
case ST_OMP_PARALLEL_DO_SIMD:
case ST_OMP_PARALLEL_LOOP:
case ST_OMP_PARALLEL_MASKED_TASKLOOP:
case ST_OMP_PARALLEL_MASKED_TASKLOOP_SIMD:
case ST_OMP_PARALLEL_MASTER_TASKLOOP:
case ST_OMP_PARALLEL_MASTER_TASKLOOP_SIMD:
case ST_OMP_MASKED_TASKLOOP:
case ST_OMP_MASKED_TASKLOOP_SIMD:
case ST_OMP_MASTER_TASKLOOP:
case ST_OMP_MASTER_TASKLOOP_SIMD:
case ST_OMP_SIMD:
case ST_OMP_TARGET_PARALLEL_DO:
case ST_OMP_TARGET_PARALLEL_DO_SIMD:
case ST_OMP_TARGET_PARALLEL_LOOP:
case ST_OMP_TARGET_SIMD:
case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
case ST_OMP_TARGET_TEAMS_LOOP:
case ST_OMP_TASKLOOP:
case ST_OMP_TASKLOOP_SIMD:
case ST_OMP_TEAMS_DISTRIBUTE:
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
case ST_OMP_TEAMS_LOOP:
case ST_OMP_TILE:
case ST_OMP_UNROLL:
case_omp_do:
st = parse_omp_do (st, 0);
if (st == ST_IMPLIED_ENDDO)
return st;
@ -6513,6 +6648,17 @@ parse_executable (gfc_statement st)
st = parse_omp_dispatch ();
continue;
case ST_OMP_METADIRECTIVE:
case ST_OMP_BEGIN_METADIRECTIVE:
st = parse_omp_metadirective_body (st);
continue;
case ST_OMP_END_METADIRECTIVE:
if (gfc_state_stack->state == COMP_OMP_BEGIN_METADIRECTIVE)
return next_statement ();
else
return st;
default:
return st;
}
@ -7278,6 +7424,10 @@ gfc_parse_file (void)
gfc_statement_label = NULL;
gfc_omp_metadirective_region_count = 0;
gfc_in_omp_metadirective_body = false;
gfc_matching_omp_context_selector = false;
if (setjmp (eof_buf))
return false; /* Come here on unexpected EOF */
@ -7589,3 +7739,16 @@ is_oacc (gfc_state_data *sd)
return false;
}
}
/* Return true if ST is a declarative OpenMP statement. */
bool
is_omp_declarative_stmt (gfc_statement st)
{
switch (st)
{
case_omp_decl:
return true;
default:
return false;
}
}

8
gcc/fortran/parse.h
View file

@ -31,7 +31,8 @@ enum gfc_compile_state
COMP_STRUCTURE, COMP_UNION, COMP_MAP,
COMP_DO, COMP_SELECT, COMP_FORALL, COMP_WHERE, COMP_CONTAINS, COMP_ENUM,
COMP_SELECT_TYPE, COMP_SELECT_RANK, COMP_OMP_STRUCTURED_BLOCK, COMP_CRITICAL,
COMP_DO_CONCURRENT, COMP_OMP_STRICTLY_STRUCTURED_BLOCK
COMP_DO_CONCURRENT, COMP_OMP_STRICTLY_STRUCTURED_BLOCK,
COMP_OMP_METADIRECTIVE, COMP_OMP_BEGIN_METADIRECTIVE
};
/* Stack element for the current compilation state. These structures
@ -67,10 +68,15 @@ bool gfc_check_do_variable (gfc_symtree *);
bool gfc_find_state (gfc_compile_state);
gfc_state_data *gfc_enclosing_unit (gfc_compile_state *);
const char *gfc_ascii_statement (gfc_statement, bool strip_sentinel = false) ;
gfc_statement gfc_omp_end_stmt (gfc_statement, bool = true, bool = true);
match gfc_match_enum (void);
match gfc_match_enumerator_def (void);
void gfc_free_enum_history (void);
extern bool gfc_matching_function;
extern bool gfc_matching_omp_context_selector;
extern bool gfc_in_omp_metadirective_body;
extern int gfc_omp_metadirective_region_count;
match gfc_match_prefix (gfc_typespec *);
bool is_oacc (gfc_state_data *);
#endif /* GFC_PARSE_H */

6
gcc/fortran/resolve.cc
View file

@ -13806,6 +13806,11 @@ gfc_resolve_code (gfc_code *code, gfc_namespace *ns)
gfc_resolve_forall (code, ns, forall_save);
forall_flag = 2;
}
else if (code->op == EXEC_OMP_METADIRECTIVE)
for (gfc_omp_variant *variant
= code->ext.omp_variants;
variant; variant = variant->next)
gfc_resolve_code (variant->code, ns);
else if (code->block)
{
omp_workshare_save = -1;
@ -14379,6 +14384,7 @@ start:
case EXEC_OMP_MASKED:
case EXEC_OMP_MASKED_TASKLOOP:
case EXEC_OMP_MASKED_TASKLOOP_SIMD:
case EXEC_OMP_METADIRECTIVE:
case EXEC_OMP_ORDERED:
case EXEC_OMP_SCAN:
case EXEC_OMP_SCOPE:

4
gcc/fortran/st.cc
View file

@ -306,6 +306,10 @@ gfc_free_statement (gfc_code *p)
case EXEC_OMP_TASKYIELD:
break;
case EXEC_OMP_METADIRECTIVE:
gfc_free_omp_variants (p->ext.omp_variants);
break;
default:
gfc_internal_error ("gfc_free_statement(): Bad statement");
}

26
gcc/fortran/symbol.cc
View file

@ -2697,10 +2697,13 @@ free_components (gfc_component *p)
static int
compare_st_labels (void *a1, void *b1)
{
int a = ((gfc_st_label *) a1)->value;
int b = ((gfc_st_label *) b1)->value;
gfc_st_label *a = (gfc_st_label *) a1;
gfc_st_label *b = (gfc_st_label *) b1;
return (b - a);
if (a->omp_region == b->omp_region)
return b->value - a->value;
else
return b->omp_region - a->omp_region;
}
@ -2750,6 +2753,8 @@ gfc_get_st_label (int labelno)
{
gfc_st_label *lp;
gfc_namespace *ns;
int omp_region = (gfc_in_omp_metadirective_body
? gfc_omp_metadirective_region_count : 0);
if (gfc_current_state () == COMP_DERIVED)
ns = gfc_current_block ()->f2k_derived;
@ -2766,10 +2771,16 @@ gfc_get_st_label (int labelno)
lp = ns->st_labels;
while (lp)
{
if (lp->value == labelno)
return lp;
if (lp->value < labelno)
if (lp->omp_region == omp_region)
{
if (lp->value == labelno)
return lp;
if (lp->value < labelno)
lp = lp->left;
else
lp = lp->right;
}
else if (lp->omp_region < omp_region)
lp = lp->left;
else
lp = lp->right;
@ -2781,6 +2792,7 @@ gfc_get_st_label (int labelno)
lp->defined = ST_LABEL_UNKNOWN;
lp->referenced = ST_LABEL_UNKNOWN;
lp->ns = ns;
lp->omp_region = omp_region;
gfc_insert_bbt (&ns->st_labels, lp, compare_st_labels);

5
gcc/fortran/trans-decl.cc
View file

@ -342,7 +342,10 @@ gfc_get_label_decl (gfc_st_label * lp)
gcc_assert (lp != NULL && lp->value <= MAX_LABEL_VALUE);
/* Build a mangled name for the label. */
sprintf (label_name, "__label_%.6d", lp->value);
if (lp->omp_region)
sprintf (label_name, "__label_%d_%.6d", lp->omp_region, lp->value);
else
sprintf (label_name, "__label_%.6d", lp->value);
/* Build the LABEL_DECL node. */
label_decl = gfc_build_label_decl (get_identifier (label_name));

233
gcc/fortran/trans-openmp.cc
View file

@ -8494,6 +8494,8 @@ gfc_trans_omp_directive (gfc_code *code)
case EXEC_OMP_MASTER_TASKLOOP:
case EXEC_OMP_MASTER_TASKLOOP_SIMD:
return gfc_trans_omp_master_masked_taskloop (code, code->op);
case EXEC_OMP_METADIRECTIVE:
return gfc_trans_omp_metadirective (code);
case EXEC_OMP_ORDERED:
return gfc_trans_omp_ordered (code);
case EXEC_OMP_PARALLEL:
@ -8587,6 +8589,100 @@ gfc_trans_omp_declare_simd (gfc_namespace *ns)
}
}
/* Translate the context selector list GFC_SELECTORS, using WHERE as the
locus for error messages. */
static tree
gfc_trans_omp_set_selector (gfc_omp_set_selector *gfc_selectors, locus where)
{
tree set_selectors = NULL_TREE;
gfc_omp_set_selector *oss;
for (oss = gfc_selectors; oss; oss = oss->next)
{
tree selectors = NULL_TREE;
gfc_omp_selector *os;
enum omp_tss_code set = oss->code;
gcc_assert (set != OMP_TRAIT_SET_INVALID);
for (os = oss->trait_selectors; os; os = os->next)
{
tree scoreval = NULL_TREE;
tree properties = NULL_TREE;
gfc_omp_trait_property *otp;
enum omp_ts_code sel = os->code;
/* Per the spec, "Implementations can ignore specified
selectors that are not those described in this section";
however, we must record such selectors because they
cause match failures. */
if (sel == OMP_TRAIT_INVALID)
{
selectors = make_trait_selector (sel, NULL_TREE, NULL_TREE,
selectors);
continue;
}
for (otp = os->properties; otp; otp = otp->next)
{
switch (otp->property_kind)
{
case OMP_TRAIT_PROPERTY_DEV_NUM_EXPR:
case OMP_TRAIT_PROPERTY_BOOL_EXPR:
{
tree expr = NULL_TREE;
gfc_se se;
gfc_init_se (&se, NULL);
gfc_conv_expr (&se, otp->expr);
expr = se.expr;
properties = make_trait_property (NULL_TREE, expr,
properties);
}
break;
case OMP_TRAIT_PROPERTY_ID:
properties
= make_trait_property (get_identifier (otp->name),
NULL_TREE, properties);
break;
case OMP_TRAIT_PROPERTY_NAME_LIST:
{
tree prop = OMP_TP_NAMELIST_NODE;
tree value = NULL_TREE;
if (otp->is_name)
value = get_identifier (otp->name);
else
value = gfc_conv_constant_to_tree (otp->expr);
properties = make_trait_property (prop, value,
properties);
}
break;
case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
properties = gfc_trans_omp_clauses (NULL, otp->clauses,
where, true);
break;
default:
gcc_unreachable ();
}
}
if (os->score)
{
gfc_se se;
gfc_init_se (&se, NULL);
gfc_conv_expr (&se, os->score);
scoreval = se.expr;
}
selectors = make_trait_selector (sel, scoreval,
properties, selectors);
}
set_selectors = make_trait_set_selector (set, selectors, set_selectors);
}
return set_selectors;
}
void
gfc_trans_omp_declare_variant (gfc_namespace *ns)
{
@ -8662,90 +8758,8 @@ gfc_trans_omp_declare_variant (gfc_namespace *ns)
&& strcmp (odv->base_proc_symtree->name, ns->proc_name->name)))
continue;
tree set_selectors = NULL_TREE;
gfc_omp_set_selector *oss;
for (oss = odv->set_selectors; oss; oss = oss->next)
{
tree selectors = NULL_TREE;
gfc_omp_selector *os;
enum omp_tss_code set = oss->code;
gcc_assert (set != OMP_TRAIT_SET_INVALID);
for (os = oss->trait_selectors; os; os = os->next)
{
tree scoreval = NULL_TREE;
tree properties = NULL_TREE;
gfc_omp_trait_property *otp;
enum omp_ts_code sel = os->code;
/* Per the spec, "Implementations can ignore specified
selectors that are not those described in this section";
however, we must record such selectors because they
cause match failures. */
if (sel == OMP_TRAIT_INVALID)
{
selectors = make_trait_selector (sel, NULL_TREE, NULL_TREE,
selectors);
continue;
}
for (otp = os->properties; otp; otp = otp->next)
{
switch (otp->property_kind)
{
case OMP_TRAIT_PROPERTY_DEV_NUM_EXPR:
case OMP_TRAIT_PROPERTY_BOOL_EXPR:
{
gfc_se se;
gfc_init_se (&se, NULL);
gfc_conv_expr (&se, otp->expr);
properties = make_trait_property (NULL_TREE, se.expr,
properties);
}
break;
case OMP_TRAIT_PROPERTY_ID:
properties
= make_trait_property (get_identifier (otp->name),
NULL_TREE, properties);
break;
case OMP_TRAIT_PROPERTY_NAME_LIST:
{
tree prop = OMP_TP_NAMELIST_NODE;
tree value = NULL_TREE;
if (otp->is_name)
value = get_identifier (otp->name);
else
value = gfc_conv_constant_to_tree (otp->expr);
properties = make_trait_property (prop, value,
properties);
}
break;
case OMP_TRAIT_PROPERTY_CLAUSE_LIST:
properties = gfc_trans_omp_clauses (NULL, otp->clauses,
odv->where, true);
break;
default:
gcc_unreachable ();
}
}
if (os->score)
{
gfc_se se;
gfc_init_se (&se, NULL);
gfc_conv_expr (&se, os->score);
scoreval = se.expr;
}
selectors = make_trait_selector (sel, scoreval,
properties, selectors);
}
set_selectors = make_trait_set_selector (set, selectors,
set_selectors);
}
tree set_selectors = gfc_trans_omp_set_selector (odv->set_selectors,
odv->where);
const char *variant_proc_name = odv->variant_proc_symtree->name;
gfc_symbol *variant_proc_sym = odv->variant_proc_symtree->n.sym;
if (variant_proc_sym == NULL || variant_proc_sym->attr.implicit_type)
@ -9048,3 +9062,54 @@ gfc_omp_call_is_alloc (tree ptr)
}
return build_call_expr_loc (input_location, fn, 1, ptr);
}
tree
gfc_trans_omp_metadirective (gfc_code *code)
{
gfc_omp_variant *variant = code->ext.omp_variants;
tree metadirective_tree = make_node (OMP_METADIRECTIVE);
SET_EXPR_LOCATION (metadirective_tree, gfc_get_location (&code->loc));
TREE_TYPE (metadirective_tree) = void_type_node;
OMP_METADIRECTIVE_VARIANTS (metadirective_tree) = NULL_TREE;
tree tree_body = NULL_TREE;
while (variant)
{
tree ctx = gfc_trans_omp_set_selector (variant->selectors,
variant->where);
ctx = omp_check_context_selector (gfc_get_location (&variant->where),
ctx, true);
if (ctx == error_mark_node)
return error_mark_node;
/* If the selector doesn't match, drop the whole variant. */
if (!omp_context_selector_matches (ctx, NULL_TREE, false))
{
variant = variant->next;
continue;
}
gfc_code *next_code = variant->code->next;
if (next_code && tree_body == NULL_TREE)
tree_body = gfc_trans_code (next_code);
if (next_code)
variant->code->next = NULL;
tree directive = gfc_trans_code (variant->code);
if (next_code)
variant->code->next = next_code;
tree body = next_code ? tree_body : NULL_TREE;
tree omp_variant = make_omp_metadirective_variant (ctx, directive, body);
OMP_METADIRECTIVE_VARIANTS (metadirective_tree)
= chainon (OMP_METADIRECTIVE_VARIANTS (metadirective_tree),
omp_variant);
variant = variant->next;
}
/* TODO: Resolve the metadirective here if possible. */
return metadirective_tree;
}

1
gcc/fortran/trans-stmt.h
View file

@ -71,6 +71,7 @@ tree gfc_trans_deallocate (gfc_code *);
tree gfc_trans_omp_directive (gfc_code *);
void gfc_trans_omp_declare_simd (gfc_namespace *);
void gfc_trans_omp_declare_variant (gfc_namespace *);
tree gfc_trans_omp_metadirective (gfc_code *code);
tree gfc_trans_oacc_directive (gfc_code *);
tree gfc_trans_oacc_declare (gfc_namespace *);

1
gcc/fortran/trans.cc
View file

@ -2588,6 +2588,7 @@ trans_code (gfc_code * code, tree cond)
case EXEC_OMP_MASTER:
case EXEC_OMP_MASTER_TASKLOOP:
case EXEC_OMP_MASTER_TASKLOOP_SIMD:
case EXEC_OMP_METADIRECTIVE:
case EXEC_OMP_ORDERED:
case EXEC_OMP_PARALLEL:
case EXEC_OMP_PARALLEL_DO:

80
gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 Normal file
View file

@ -0,0 +1,80 @@
! { dg-do compile }
program main
integer, parameter :: N = 10
integer, dimension(N) :: a
integer, dimension(N) :: b
integer, dimension(N) :: c
integer :: i
do i = 1, N
a(i) = i * 2
b(i) = i * 3
end do
!$omp metadirective &
!$omp& default (teams loop) &
!$omp& default (parallel loop) ! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." }
do i = 1, N
c(i) = a(i) * b(i)
end do
!$omp metadirective &
!$omp& otherwise (teams loop) &
!$omp& default (parallel loop) ! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." }
do i = 1, N
c(i) = a(i) * b(i)
end do
!$omp metadirective &
!$omp& otherwise (teams loop) &
!$omp& otherwise (parallel loop) ! { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective' at .1." }
do i = 1, N
c(i) = a(i) * b(i)
end do
!$omp metadirective default (xyz) ! { dg-error "Unclassifiable OpenMP directive at .1." }
do i = 1, N
c(i) = a(i) * b(i)
end do
!$omp metadirective &
!$omp& default (teams loop) &
!$omp& where (device={arch("nvptx")}: parallel loop) ! { dg-error "expected 'when', 'otherwise', or 'default' at .1." }
do i = 1, N
c(i) = a(i) * b(i)
end do
!$omp metadirective &
!$omp& otherwise (teams loop) &
!$omp& when (device={arch("nvptx")}: parallel loop) ! { dg-error "'otherwise' or 'default' clause must appear last" }
do i = 1, N
c(i) = a(i) * b(i)
end do
!$omp metadirective &
!$omp& when (device={arch("nvptx")} parallel loop) & ! { dg-error "expected .:." }
!$omp& default (teams loop)
do i = 1, N
c(i) = a(i) * b(i)
end do
! Test improperly nested metadirectives - even though the second
! metadirective resolves to 'omp nothing', that is not the same as there
! being literally nothing there.
!$omp metadirective &
!$omp& when (implementation={vendor("gnu")}: parallel do)
!$omp metadirective &
!$omp& when (implementation={vendor("cray")}: parallel do) ! { dg-error "Unexpected !.OMP METADIRECTIVE statement" }
do i = 1, N
c(i) = a(i) * b(i)
end do
!$omp begin metadirective &
!$omp& when (device={arch("nvptx")}: parallel do) &
!$omp& default (barrier) ! { dg-error "variant directive used in OMP BEGIN METADIRECTIVE at .1. must have a corresponding end directive" }
do i = 1, N
c(i) = a(i) * b(i)
end do
!$omp end metadirective ! { dg-error "Unexpected !.OMP END METADIRECTIVE statement at .1." }
end program

40
gcc/testsuite/gfortran.dg/gomp/metadirective-10.f90 Normal file
View file

@ -0,0 +1,40 @@
! { dg-do compile }
program metadirectives
implicit none
logical :: UseDevice
!$OMP metadirective &
!$OMP when ( user = { condition ( UseDevice ) } &
!$OMP : parallel ) &
!$OMP default ( parallel )
block
call bar()
end block
!$OMP metadirective &
!$OMP when ( user = { condition ( UseDevice ) } &
!$OMP : parallel ) &
!$OMP default ( parallel )
call bar()
!$omp end parallel ! Accepted, because all cases have 'parallel'
!$OMP begin metadirective &
!$OMP when ( user = { condition ( UseDevice ) } &
!$OMP : nothing ) &
!$OMP default ( parallel )
call bar()
block
call foo()
end block
!$OMP end metadirective
!$OMP begin metadirective &
!$OMP when ( user = { condition ( UseDevice ) } &
!$OMP : parallel ) &
!$OMP default ( parallel )
call bar()
!$omp end parallel ! { dg-error "Unexpected !.OMP END PARALLEL statement at .1." }
end program ! { dg-error "Unexpected END statement at .1." }
! { dg-error "Unexpected end of file" "" { target *-*-* } 0 }

33
gcc/testsuite/gfortran.dg/gomp/metadirective-11.f90 Normal file
View file

@ -0,0 +1,33 @@
! { dg-do compile }
! { dg-ice "Statements following a block in a metadirective" }
! PR fortran/107067
program metadirectives
implicit none
logical :: UseDevice
!$OMP begin metadirective &
!$OMP when ( user = { condition ( UseDevice ) } &
!$OMP : nothing ) &
!$OMP default ( parallel )
block
call foo()
end block
call bar() ! FIXME/XFAIL ICE in parse_omp_metadirective_body()
!$omp end metadirective
!$OMP begin metadirective &
!$OMP when ( user = { condition ( UseDevice ) } &
!$OMP : nothing ) &
!$OMP default ( parallel )
block
call bar()
end block
block ! FIXME/XFAIL ICE in parse_omp_metadirective_body()
call foo()
end block
!$omp end metadirective
end program

18
gcc/testsuite/gfortran.dg/gomp/metadirective-12.f90 Normal file
View file

@ -0,0 +1,18 @@
! { dg-do compile }
! PR112779 item H; this testcase used to ICE.
program test
implicit none
integer, parameter :: N = 100
integer :: x(N), y(N), z(N)
block
integer :: i
!$omp metadirective &
!$omp& when(device={arch("nvptx")}: teams loop) &
!$omp& default(parallel loop)
do i = 1, N
z(i) = x(i) * y(i)
enddo
end block
end

30
gcc/testsuite/gfortran.dg/gomp/metadirective-13.f90 Normal file
View file

@ -0,0 +1,30 @@
! { dg-do compile }
subroutine foo
implicit none
external f
!$omp dispatch
call f()
!$omp dispatch
call f()
!$omp end dispatch
!$omp begin metadirective when(construct={parallel} : nothing) otherwise(dispatch)
call f()
!$omp end metadirective
end
subroutine bar
implicit none
integer :: x
!$omp atomic update
x = x + 1
!$omp atomic update
x = x + 1
!$omp end atomic
!$omp begin metadirective when(construct={parallel} : nothing) otherwise(atomic update)
x = x + 1
!$omp end metadirective
end

72
gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 Normal file
View file

@ -0,0 +1,72 @@
! { dg-do compile }
program main
integer, parameter :: N = 100
integer :: x = 0
integer :: y = 0
integer :: i
! Test implicit default directive
!$omp metadirective &
!$omp& when (device={arch("nvptx")}: barrier)
x = 1
! Test implicit default directive combined with a directive that takes a
! do loop.
!$omp metadirective &
!$omp& when (device={arch("nvptx")}: parallel do)
do i = 1, N
x = x + i
end do
! Test with multiple standalone directives.
!$omp metadirective &
!$omp& when (device={arch("nvptx")}: barrier) &
!$omp& default (flush)
x = 1
! Test combining a standalone directive with one that takes a do loop.
!$omp metadirective &
!$omp& when (device={arch("nvptx")}: parallel do) &
!$omp& default (barrier)
do i = 1, N
x = x + i
end do
! Test combining a directive that takes a do loop with one that takes
! a statement body.
!$omp begin metadirective &
!$omp& when (device={arch("nvptx")}: parallel do) &
!$omp& default (parallel)
do i = 1, N
x = x + i
end do
!$omp end metadirective
! Test labels in the body.
!$omp begin metadirective &
!$omp& when (device={arch("nvptx")}: parallel do) &
!$omp& when (device={arch("gcn")}: parallel)
do i = 1, N
x = x + i
if (x .gt. N/2) goto 10
10 x = x + 1
goto 20
x = x + 2
20 continue
end do
!$omp end metadirective
! Test that commas are permitted before each clause.
!$omp begin metadirective, &
!$omp& when (device={arch("nvptx")}: parallel do) &
!$omp& , when (device={arch("gcn")}: parallel) &
!$omp& , default (parallel)
do i = 1, N
x = x + i
end do
!$omp end metadirective
! Test empty metadirective.
!$omp metadirective
end program

25
gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 Normal file
View file

@ -0,0 +1,25 @@
! { dg-do compile }
! { dg-additional-options "-fdump-tree-gimple" }
module test
integer, parameter :: N = 100
contains
subroutine f (x, y, z)
integer :: x(N), y(N), z(N)
!$omp target map (to: v1, v2) map(from: v3)
!$omp metadirective &
!$omp& when(device={arch("nvptx")}: teams loop) &
!$omp& default(parallel loop)
do i = 1, N
z(i) = x(i) * y(i)
enddo
!$omp end target
end subroutine
end module
! If offload device "nvptx" isn't supported, the front end can eliminate
! that alternative and not produce a metadirective at all. Otherwise this
! won't be resolved until late.
! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" { target { ! offload_nvptx } } } }
! { dg-final { scan-tree-dump "#pragma omp metadirective" "gimple" { target { offload_nvptx } } } }

37
gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 Normal file
View file

@ -0,0 +1,37 @@
! { dg-do compile }
! { dg-additional-options "-fdump-tree-original" }
! { dg-additional-options "-fdump-tree-gimple" }
program test
implicit none
integer, parameter :: N = 100
real :: a(N)
!$omp target map(from: a)
call f (a, 3.14159)
!$omp end target
call f (a, 2.71828)
contains
subroutine f (a, x)
integer :: i
real :: a(N), x
!$omp declare target
!$omp metadirective &
!$omp& when (construct={target}: distribute parallel do ) &
!$omp& default(parallel do simd)
do i = 1, N
a(i) = x * i
end do
end subroutine
end program
! The metadirective should be resolved during Gimplification.
! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } }
! { dg-final { scan-tree-dump-times "when \\(construct = .*target.*\\):" 1 "original" } }
! { dg-final { scan-tree-dump-times "otherwise:" 1 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } }
! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }

30
gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 Normal file
View file

@ -0,0 +1,30 @@
! { dg-do compile }
! { dg-additional-options "-fdump-tree-gimple" }
module test
integer, parameter :: N = 100
contains
subroutine f (a, flag)
integer :: a(N)
logical :: flag
integer :: i
!$omp metadirective &
!$omp& when (user={condition(flag)}: &
!$omp& target teams distribute parallel do map(from: a(1:N))) &
!$omp& default(parallel do)
do i = 1, N
a(i) = i
end do
end subroutine
end module
! The metadirective should be resolved at parse time, but is currently
! resolved during Gimplification
! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp target" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp distribute" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } }

26
gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 Normal file
View file

@ -0,0 +1,26 @@
! { dg-do compile }
! { dg-additional-options "-fdump-tree-gimple" }
module test
integer, parameter :: N = 100
contains
subroutine f (a, run_parallel, run_guided)
integer :: a(N)
logical :: run_parallel, run_guided
integer :: i
!$omp begin metadirective when(user={condition(run_parallel)}: parallel)
!$omp metadirective &
!$omp& when(construct={parallel}, user={condition(run_guided)}: &
!$omp& do schedule(guided)) &
!$omp& when(construct={parallel}: do schedule(static))
do i = 1, N
a(i) = i
end do
!$omp end metadirective
end subroutine
end module
! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } }

42
gcc/testsuite/gfortran.dg/gomp/metadirective-7.f90 Normal file
View file

@ -0,0 +1,42 @@
! { dg-do compile }
! { dg-additional-options "-fdump-tree-gimple -fdump-tree-ompdevlow" }
subroutine f (a, num)
integer, parameter :: N = 256
integer :: a(N)
integer :: num
integer :: i
!$omp metadirective &
!$omp& when (target_device={device_num(num), kind("gpu"), arch("nvptx")}: &
!$omp& target parallel do map(tofrom: a(1:N))) &
!$omp& when (target_device={device_num(num), kind("gpu"), &
!$omp& arch("amdgcn"), isa("gfx906")}: &
!$omp& target parallel do) &
!$omp& when (target_device={device_num(num), kind("cpu"), arch("x86_64")}: &
!$omp& parallel do)
do i = 1, N
a(i) = a(i) + i
end do
!$omp metadirective &
!$omp& when (target_device={kind("gpu"), arch("nvptx")}: &
!$omp& target parallel do map(tofrom: a(1:N)))
do i = 1, N
a(i) = a(i) + i
end do
end subroutine
! For configurations with offloading, we expect one "pragma omp target"
! with "device(num)" for each target_device selector that specifies
! "device_num(num)". Without offloading, there should be zero as the
! resolution happens during gimplification.
! { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(" 3 "gimple" { target offloading_enabled } } }
! { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(" 0 "gimple" { target { ! offloading_enabled } } } }
! For configurations with offloading, expect one OMP_TARGET_DEVICE_MATCHES
! for each kind/arch/isa selector. These are supposed to go away after
! ompdevlow.
! { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 9 "gimple" { target offloading_enabled } } }
! { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "gimple" { target { ! offloading_enabled } } } }
! { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "ompdevlow" { target offloading_enabled } } }

22
gcc/testsuite/gfortran.dg/gomp/metadirective-8.f90 Normal file
View file

@ -0,0 +1,22 @@
! { dg-do compile }
program test
integer :: i
integer, parameter :: N = 100
integer :: sum = 0
! The compiler should never consider a situation where both metadirectives
! match, but that does not matter because the spec says "Replacement of
! the metadirective with the directive variant associated with any of the
! dynamic replacement candidates must result in a conforming OpenMP
! program. So the second metadirective is rejected as not being
! a valid loop-nest even if the first one does not match.
!$omp metadirective when (implementation={vendor("ibm")}: &
!$omp& target teams distribute)
!$omp metadirective when (implementation={vendor("gnu")}: parallel do) ! { dg-error "Unexpected !.OMP METADIRECTIVE statement" }
do i = 1, N
sum = sum + i
end do
end program

30
gcc/testsuite/gfortran.dg/gomp/metadirective-9.f90 Normal file
View file

@ -0,0 +1,30 @@
! { dg-do compile }
program OpenMP_Metadirective_WrongEnd_Test
implicit none
integer :: &
iaVS, iV, jV, kV
integer, dimension ( 3 ) :: &
lV, uV
logical :: &
UseDevice
!$OMP metadirective &
!$OMP when ( user = { condition ( UseDevice ) } &
!$OMP : target teams distribute parallel do simd collapse ( 3 ) &
!$OMP private ( iaVS ) ) &
!$OMP default ( parallel do simd collapse ( 3 ) private ( iaVS ) )
do kV = lV ( 3 ), uV ( 3 )
do jV = lV ( 2 ), uV ( 2 )
do iV = lV ( 1 ), uV ( 1 )
end do
end do
end do
!$OMP end target teams distribute parallel do simd ! { dg-error "Unexpected !.OMP END TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD statement in OMP METADIRECTIVE block at .1." }
end program

260
gcc/testsuite/gfortran.dg/gomp/metadirective-construct.f90 Normal file
View file

@ -0,0 +1,260 @@
! { dg-do compile }
! { dg-additional-options "-foffload=disable -fdump-tree-original -fdump-tree-gimple" }
program main
implicit none
integer, parameter :: N = 10
double precision, parameter :: S = 2.0
double precision :: a(N)
call init (N, a)
call f1 (N, a, S)
call check (N, a, S)
call init (N, a)
call f2 (N, a, S)
call check (N, a, S)
call init (N, a)
call f3 (N, a, S)
call check (N, a, S)
call init (N, a)
call f4 (N, a, S)
call check (N, a, S)
call init (N, a)
call f5 (N, a, S)
call check (N, a, S)
call init (N, a)
call f6 (N, a, S)
call check (N, a, S)
call init (N, a)
call f7 (N, a, S)
call check (N, a, S)
call init (N, a)
call f8 (N, a, S)
call check (N, a, S)
call init (N, a)
call f9 (N, a, S)
call check (N, a, S)
contains
subroutine init (n, a)
implicit none
integer :: n
double precision :: a(n)
integer :: i
do i = 1, n
a(i) = i
end do
end subroutine
subroutine check (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
do i = 1, n
if (a(i) /= i * s) error stop
end do
end subroutine
! Check various combinations for enforcing correct ordering of
! construct matches.
subroutine f1 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target teams
!$omp parallel
!$omp metadirective &
!$omp & when (construct={target} &
!$omp & : do) &
!$omp & default (error at(execution) message("f1 match failed"))
do i = 1, n
a(i) = a(i) * s
end do
!$omp end parallel
!$omp end target teams
end subroutine
subroutine f2 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target teams
!$omp parallel
!$omp metadirective &
!$omp & when (construct={teams, parallel} &
!$omp & : do) &
!$omp & default (error at(execution) message("f2 match failed"))
do i = 1, n
a(i) = a(i) * s
end do
!$omp end parallel
!$omp end target teams
end subroutine
subroutine f3 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target teams
!$omp parallel
!$omp metadirective &
!$omp & when (construct={target, teams, parallel} &
!$omp & : do) &
!$omp & default (error at(execution) message("f3 match failed"))
do i = 1, n
a(i) = a(i) * s
end do
!$omp end parallel
!$omp end target teams
end subroutine
subroutine f4 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target teams
!$omp parallel
!$omp metadirective &
!$omp & when (construct={target, parallel} &
!$omp & : do) &
!$omp & default (error at(execution) message("f4 match failed"))
do i = 1, n
a(i) = a(i) * s
end do
!$omp end parallel
!$omp end target teams
end subroutine
subroutine f5 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target teams
!$omp parallel
!$omp metadirective &
!$omp & when (construct={target, teams} &
!$omp & : do) &
!$omp & default (error at(execution) message("f5 match failed"))
do i = 1, n
a(i) = a(i) * s
end do
!$omp end parallel
!$omp end target teams
end subroutine
! Next batch is for things where the construct doesn't match the context.
subroutine f6 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target
!$omp teams
!$omp metadirective &
!$omp & when (construct={parallel} &
!$omp & : error at(execution) message("f6 match failed")) &
!$omp & default (parallel do)
do i = 1, n
a(i) = a(i) * s
end do
!$omp end teams
!$omp end target
end subroutine
subroutine f7 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target
!$omp teams
!$omp metadirective &
!$omp & when (construct={target, parallel} &
!$omp & : error at(execution) message("f7 match failed")) &
!$omp & default (parallel do)
do i = 1, n
a(i) = a(i) * s
end do
!$omp end teams
!$omp end target
end subroutine
subroutine f8 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target
!$omp teams
!$omp metadirective &
!$omp & when (construct={parallel, target} &
!$omp & : error at(execution) message("f8 match failed")) &
!$omp & default (parallel do)
do i = 1, n
a(i) = a(i) * s
end do
!$omp end teams
!$omp end target
end subroutine
! Next test choosing the best alternative when there are multiple
! matches.
subroutine f9 (n, a, s)
implicit none
integer :: n
double precision :: a(n)
double precision :: s
integer :: i
!$omp target teams
!$omp parallel
!$omp metadirective &
!$omp & when (construct={teams, parallel} &
!$omp & : error at(execution) message("f9 match incorrect 1")) &
!$omp & when (construct={target, teams, parallel} &
!$omp & : do) &
!$omp & when (construct={target, teams} &
!$omp & : error at(execution) message("f9 match incorrect 2")) &
!$omp & default (error at(execution) message("f9 match failed"))
do i = 1, n
a(i) = a(i) * s
end do
!$omp end parallel
!$omp end target teams
end subroutine
end program
! Note there are no tests for the matching the extended simd clause
! syntax, which is only useful for "declare variant".
! After parsing, there should be a runtime error call for each of the
! failure cases, but they should all be optimized away during OMP
! lowering.
! { dg-final { scan-tree-dump-times "__builtin_GOMP_error" 11 "original" } }
! { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "gimple" } }

122
gcc/testsuite/gfortran.dg/gomp/metadirective-no-score.f90 Normal file
View file

@ -0,0 +1,122 @@
! { dg-do compile { target x86_64-*-* } }
! { dg-additional-options "-foffload=disable" }
! This test is expected to fail with compile-time errors:
! "A trait-score cannot be specified in traits from the construct,
! device or target_device trait-selector-sets."
subroutine f1 (n, a, s)
implicit none
integer :: n
double precision :: a(*)
double precision :: s
integer :: i
!$omp metadirective &
!$omp& when (device={kind (score(5) : host)} &
!$omp& : parallel do)
! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 }
do i = 1, n
a(i) = a(i) * s;
end do
end subroutine
subroutine f2 (n, a, s)
implicit none
integer :: n
double precision :: a(*)
double precision :: s
integer :: i
!$omp metadirective &
!$omp& when (device={kind (host), arch (score(6) : x86_64), isa (avx512f)} &
!$omp& : parallel do)
! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 }
do i = 1, n
a(i) = a(i) * s;
end do
end subroutine
subroutine f3 (n, a, s)
implicit none
integer :: n
double precision :: a(*)
double precision :: s
integer :: i
!$omp metadirective &
!$omp& when (device={kind (host), arch (score(6) : x86_64), &
!$omp& isa (score(7): avx512f)} &
!$omp& : parallel do)
! { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-3 }
do i = 1, n
a(i) = a(i) * s;
end do
end subroutine
subroutine f4 (n, a, s)
implicit none
integer :: n
double precision :: a(*)
double precision :: s
integer :: i
integer, parameter :: omp_initial_device = -1
!$omp metadirective &
!$omp& when (target_device={device_num (score(42) : omp_initial_device), &
!$omp& kind (host)} &
!$omp& : parallel do)
! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 }
do i = 1, n
a(i) = a(i) * s;
end do
end subroutine
subroutine f5 (n, a, s)
implicit none
integer :: n
double precision :: a(*)
double precision :: s
integer :: i
integer, parameter :: omp_initial_device = -1
!$omp metadirective &
!$omp& when (target_device={device_num(omp_initial_device), &
!$omp& kind (score(5) : host)} &
!$omp& : parallel do)
! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 }
do i = 1, n
a(i) = a(i) * s;
end do
end subroutine
subroutine f6 (n, a, s)
implicit none
integer :: n
double precision :: a(*)
double precision :: s
integer :: i
integer, parameter :: omp_initial_device = -1
!$omp metadirective &
!$omp& when (target_device={device_num(omp_initial_device), kind (host), &
!$omp& arch (score(6) : x86_64), isa (avx512f)} &
!$omp& : parallel do)
! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 }
do i = 1, n
a(i) = a(i) * s;
end do
end subroutine
subroutine f7 (n, a, s)
implicit none
integer :: n
double precision :: a(*)
double precision :: s
integer :: i
integer, parameter :: omp_initial_device = -1
!$omp metadirective &
!$omp& when (target_device={device_num(omp_initial_device), kind (host), &
!$omp& arch (score(6) : x86_64), &
!$omp& isa (score(7): avx512f)} &
!$omp& : parallel do)
! { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 }
do i = 1, n
a(i) = a(i) * s;
end do
end subroutine

29
gcc/testsuite/gfortran.dg/gomp/pure-1.f90
View file

@ -110,3 +110,32 @@ pure integer function func_tile(n)
end do
func_tile = r
end
pure logical function func_metadirective()
implicit none
!$omp metadirective
func_metadirective = .false.
end
! not 'parallel' not pure -> invalid in 5.2; + in general invalid in 5.1
pure logical function func_metadirective_2 ()
implicit none
integer :: i, n
n = 0
!$omp metadirective when (device={arch("nvptx")} : parallel do) ! { dg-error "OpenMP directive at .1. is not pure and thus may not appear in a PURE procedure" }
do i = 1, 5
n = n + i
end do
end
! unroll is supposed to be pure, so this case is OK
pure logical function func_metadirective_3()
implicit none
integer :: i, n
n = 0
!$omp metadirective when(device={arch("nvptx")} : unroll full)
do i = 1, 5
n = n + i
end do
end

8
gcc/testsuite/gfortran.dg/gomp/pure-2.f90
View file

@ -26,14 +26,6 @@ logical function func_interchange(n)
end do
end
!pure logical function func_metadirective()
logical function func_metadirective()
implicit none
!$omp metadirective ! { dg-error "Unclassifiable OpenMP directive" }
func_metadirective = .false.
end
!pure logical function func_reverse(n)
logical function func_reverse(n)
implicit none

61
libgomp/testsuite/libgomp.fortran/metadirective-1.f90 Normal file
View file

@ -0,0 +1,61 @@
! { dg-do run }
program test
implicit none
integer, parameter :: N = 100
integer :: x(N), y(N), z(N)
integer :: i
do i = 1, N
x(i) = i;
y(i) = -i;
end do
call f (x, y, z)
do i = 1, N
if (z(i) .ne. x(i) * y(i)) stop 1
end do
! -----
do i = 1, N
x(i) = i;
y(i) = -i;
end do
call g (x, y, z)
do i = 1, N
if (z(i) .ne. x(i) * y(i)) stop 1
end do
contains
subroutine f (x, y, z)
integer :: x(N), y(N), z(N)
!$omp target map (to: x, y) map(from: z)
block
!$omp metadirective &
!$omp& when(device={arch("nvptx")}: teams loop) &
!$omp& default(parallel loop)
do i = 1, N
z(i) = x(i) * y(i)
enddo
end block
end subroutine
subroutine g (x, y, z)
integer :: x(N), y(N), z(N)
!$omp target map (to: x, y) map(from: z)
block
!$omp metadirective &
!$omp& when(device={arch("nvptx")}: teams loop) &
!$omp& default(parallel loop)
do i = 1, N
z(i) = x(i) * y(i)
enddo
end block
!$omp end target
end subroutine
end program

38
libgomp/testsuite/libgomp.fortran/metadirective-2.f90 Normal file
View file

@ -0,0 +1,38 @@
! { dg-do run }
program test
implicit none
integer, parameter :: N = 100
real, parameter :: PI_CONST = 2.0*acos(0.0)
real, parameter :: E_CONST = exp(1.0)
real, parameter :: EPSILON = 0.001
integer :: i
real :: a(N)
!$omp target map(from: a)
call f (a, PI_CONST)
!$omp end target
do i = 1, N
if (abs (a(i) - (PI_CONST * i)) .gt. EPSILON) stop 1
end do
call f (a, E_CONST)
do i = 1, N
if (abs (a(i) - (E_CONST * i)) .gt. EPSILON) stop 2
end do
contains
subroutine f (a, x)
integer :: i
real :: a(N), x
!$omp declare target
!$omp metadirective &
!$omp& when (construct={target}: distribute parallel do ) &
!$omp& default(parallel do simd)
do i = 1, N
a(i) = x * i
end do
end subroutine
end program

29
libgomp/testsuite/libgomp.fortran/metadirective-3.f90 Normal file
View file

@ -0,0 +1,29 @@
! { dg-do run }
program test
implicit none
integer, parameter :: N = 100
integer :: a(N)
integer :: res
if (f (a, .false.)) stop 1
if (.not. f (a, .true.)) stop 2
contains
logical function f (a, flag)
integer :: a(N)
logical :: flag
logical :: res = .false.
integer :: i
f = .false.
!$omp metadirective &
!$omp& when (user={condition(.not. flag)}: &
!$omp& target teams distribute parallel do &
!$omp& map(from: a(1:N)) private(res)) &
!$omp& default(parallel do)
do i = 1, N
a(i) = i
f = .true.
end do
end function
end program

46
libgomp/testsuite/libgomp.fortran/metadirective-4.f90 Normal file
View file

@ -0,0 +1,46 @@
! { dg-do run }
program test
use omp_lib
implicit none
integer, parameter :: N = 100
integer :: a(N)
logical :: is_parallel, is_static
! is_static is always set if run_parallel is false.
call f (a, .false., .false., is_parallel, is_static)
if (is_parallel .or. .not. is_static) stop 1
call f (a, .false., .true., is_parallel, is_static)
if (is_parallel .or. .not. is_static) stop 2
call f (a, .true., .false., is_parallel, is_static)
if (.not. is_parallel .or. is_static) stop 3
call f (a, .true., .true., is_parallel, is_static)
if (.not. is_parallel .or. .not. is_static) stop 4
contains
subroutine f (a, run_parallel, run_static, is_parallel, is_static)
integer :: a(N)
logical, intent(in) :: run_parallel, run_static
logical, intent(out) :: is_parallel, is_static
integer :: i
is_parallel = .false.
is_static = .false.
!$omp begin metadirective when(user={condition(run_parallel)}: parallel)
if (omp_in_parallel ()) is_parallel = .true.
!$omp metadirective &
!$omp& when(construct={parallel}, user={condition(.not. run_static)}: &
!$omp& do schedule(guided) private(is_static)) &
!$omp& when(construct={parallel}: do schedule(static))
do i = 1, N
a(i) = i
is_static = .true.
end do
!$omp end metadirective
end subroutine
end program

44
libgomp/testsuite/libgomp.fortran/metadirective-5.f90 Normal file
View file

@ -0,0 +1,44 @@
! { dg-do run }
program main
use omp_lib
implicit none
integer, parameter :: N = 100
integer :: a(N)
integer :: on_device_count = 0
integer :: i
do i = 1, N
a(i) = i
end do
do i = 0, omp_get_num_devices ()
on_device_count = on_device_count + f (a, i)
end do
if (on_device_count .ne. omp_get_num_devices ()) stop 1
do i = 1, N
if (a(i) .ne. 2 * i) stop 2;
end do
contains
integer function f (a, num)
integer, intent(inout) :: a(N)
integer, intent(in) :: num
integer :: on_device
integer :: i
on_device = 0
!$omp metadirective &
!$omp& when (target_device={device_num(num), kind("gpu")}: &
!$omp& target parallel do map(to: a(1:N)), map(from: on_device)) &
!$omp& default (parallel do private(on_device))
do i = 1, N
a(i) = a(i) + i
on_device = 1
end do
f = on_device;
end function
end program

58
libgomp/testsuite/libgomp.fortran/metadirective-6.f90 Normal file
View file

@ -0,0 +1,58 @@
! { dg-do compile }
program test
implicit none
integer, parameter :: N = 100
integer :: x(N), y(N), z(N)
integer :: i
contains
subroutine f (x, y, z)
integer :: x(N), y(N), z(N)
!$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
block
!$omp metadirective &
!$omp& when(device={arch("nvptx")}: teams loop) &
!$omp& default(parallel loop) ! { dg-error "\\(1\\)" }
! FIXME: The line above should be the same error as above but some fails here with -fno-diagnostics-show-caret
! Seems as if some gcc/testsuite/ fix is missing for libgomp/testsuite
do i = 1, N
z(i) = x(i) * y(i)
enddo
z(N) = z(N) + 1 ! <<< invalid
end block
end subroutine
subroutine f2 (x, y, z)
integer :: x(N), y(N), z(N)
!$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
block
integer :: i ! << invalid
!$omp metadirective &
!$omp& when(device={arch("nvptx")}: teams loop) &
!$omp& default(parallel loop)
do i = 1, N
z(i) = x(i) * y(i)
enddo
end block
end subroutine
subroutine g (x, y, z)
integer :: x(N), y(N), z(N)
!$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" }
block
!$omp metadirective & ! <<<< invalid
!$omp& when(device={arch("nvptx")}: flush) &
!$omp& default(nothing)
!$omp teams loop
do i = 1, N
z(i) = x(i) * y(i)
enddo
end block
!$omp end target
end subroutine
end program