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;
&& 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;
}
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;
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;
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)
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:
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,
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;
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
{
locus where;
gfc_namespace *ns;
+ int omp_region;
}
gfc_st_label;
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
};
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;
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 *);
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);
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
{
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);
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);
{"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.
/* 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)
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;
}
}
+/* 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)
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 ();
}
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");
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;
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
{
|| (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
{
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)
{
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
{
}
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 (" }");
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)
}
+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)
{
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)
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)
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;
}
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);
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);
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;
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 ();
return ST_GET_FCN_CHARACTERISTICS;
}
+gfc_statement
+match_omp_directive (void)
+{
+ return decode_omp_directive ();
+}
+
static gfc_statement
decode_gcc_attribute (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(). */
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;
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;
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";
break;
case ST_ENTRY:
+ case ST_OMP_METADIRECTIVE:
+ case ST_OMP_BEGIN_METADIRECTIVE:
case_executable:
case_exec_markers:
add_statement ();
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. */
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)
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
{
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 ();
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;
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:
return st;
}
-
static gfc_statement
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 "
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
parse_executable (gfc_statement st)
{
int close_flag;
+ bool one_stmt_p = false;
in_exec_part = true;
if (st == ST_NONE)
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)
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;
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;
}
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 */
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;
+ }
+}
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
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 */
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;
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:
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");
}
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;
}
{
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;
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;
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);
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));
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:
}
}
+/* 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)
{
&& 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)
}
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;
+}
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 *);
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:
--- /dev/null
+! { 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
--- /dev/null
+! { 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 }
--- /dev/null
+! { 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
+
+
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { 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 } } } }
--- /dev/null
+! { 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" } }
--- /dev/null
+! { 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" } }
--- /dev/null
+! { 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" } }
--- /dev/null
+! { 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 } } }
--- /dev/null
+! { 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
+
--- /dev/null
+! { 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
+
--- /dev/null
+! { 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" } }
--- /dev/null
+! { 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
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
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
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { 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
projects / thirdparty / gcc.git / commitdiff
