@@ -8404,6 +8404,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;
@@ -8411,6 +8412,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;
}
@@ -8556,6 +8579,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;
@@ -2162,6 +2162,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;
@@ -2359,6 +2360,25 @@ show_omp_node (int level, gfc_code *c)
d = d->block;
}
}
+ else if (c->op == EXEC_OMP_METADIRECTIVE)
+ {
+ gfc_omp_metadirective_variant *variant
+ = c->ext.omp_metadirective_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)
@@ -3493,6 +3513,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:
@@ -317,6 +317,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,
@@ -1583,7 +1584,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;
ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3;
ENUM_BITFIELD (gfc_omp_device_type) device_type:2;
@@ -1703,6 +1704,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_metadirective_variant
+{
+ struct gfc_omp_metadirective_variant *next;
+ locus where; /* Where the metadirective clause occurred. */
+
+ gfc_omp_set_selector *selectors;
+ enum gfc_statement stmt;
+ struct gfc_code *code;
+
+} gfc_omp_metadirective_variant;
+#define gfc_get_omp_metadirective_variant() XCNEW (gfc_omp_metadirective_variant)
typedef struct gfc_omp_udr
{
@@ -1751,6 +1763,7 @@ typedef struct gfc_st_label
locus where;
gfc_namespace *ns;
+ int omp_region;
}
gfc_st_label;
@@ -3009,6 +3022,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_METADIRECTIVE,
EXEC_OMP_ERROR, EXEC_OMP_ALLOCATE, EXEC_OMP_ALLOCATORS
};
@@ -3067,6 +3081,7 @@ typedef struct gfc_code
gfc_omp_clauses *omp_clauses;
const char *omp_name;
gfc_omp_namelist *omp_namelist;
+ gfc_omp_metadirective_variant *omp_metadirective_variants;
bool omp_bool;
}
ext; /* Points to additional structures required by statement */
@@ -3661,6 +3676,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_metadirective_variants (gfc_omp_metadirective_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 *);
@@ -3942,6 +3958,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);
@@ -4038,4 +4056,8 @@ bool gfc_is_reallocatable_lhs (gfc_expr *);
void finish_oacc_declare (gfc_namespace *, gfc_symbol *, bool);
void gfc_adjust_builtins (void);
+/* trans-openmp.c */
+
+bool gfc_skip_omp_metadirective_variant (gfc_omp_metadirective_variant *);
+
#endif /* GCC_GFORTRAN_H */
@@ -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
{
@@ -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);
@@ -178,6 +179,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);
@@ -113,7 +113,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)
@@ -124,17 +125,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;
@@ -339,6 +348,19 @@ gfc_free_omp_udr (gfc_omp_udr *omp_udr)
}
}
+/* Free variants of an !$omp metadirective construct. */
+
+void
+gfc_free_omp_metadirective_variants (gfc_omp_metadirective_variant *variant)
+{
+ while (variant)
+ {
+ gfc_omp_metadirective_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)
@@ -1866,8 +1888,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 ();
@@ -3789,9 +3810,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");
@@ -4479,7 +4498,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;
@@ -5810,14 +5829,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)
{
@@ -5874,7 +5896,7 @@ 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)
{
do
{
@@ -5907,9 +5929,9 @@ 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)
return MATCH_ERROR;
@@ -6010,7 +6032,8 @@ gfc_match_omp_declare_variant (void)
return MATCH_ERROR;
}
- if (gfc_match_omp_context_selector_specification (odv) != MATCH_YES)
+ if (gfc_match_omp_context_selector_specification (&odv->set_selectors)
+ != MATCH_YES)
return MATCH_ERROR;
if (gfc_match (" )") != MATCH_YES)
@@ -6026,6 +6049,157 @@ gfc_match_omp_declare_variant (void)
}
+static match
+match_omp_metadirective (bool begin_p)
+{
+ locus old_loc = gfc_current_locus;
+ gfc_omp_metadirective_variant *variants_head;
+ gfc_omp_metadirective_variant **next_variant = &variants_head;
+ bool default_seen = false;
+
+ /* Parse the context selectors. */
+ for (;;)
+ {
+ bool default_p = false;
+ gfc_omp_set_selector *selectors = NULL;
+ 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_eos () == MATCH_YES)
+ break;
+ 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;
+ }
+
+ if (!default_p)
+ {
+ if (gfc_match_omp_context_selector_specification (&selectors)
+ != 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_metadirective_variant *omv = gfc_get_omp_metadirective_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;
+ }
+
+ if (!gfc_skip_omp_metadirective_variant (omv))
+ {
+ *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_metadirective_variant *omv = gfc_get_omp_metadirective_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_metadirective_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)
{
@@ -10919,6 +11093,19 @@ resolve_omp_do (gfc_code *code)
restructure_intervening_code (&(code->block->next), code, count);
}
+static void
+resolve_omp_metadirective (gfc_code *code, gfc_namespace *ns)
+{
+ gfc_omp_metadirective_variant *variant = code->ext.omp_metadirective_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)
@@ -11462,13 +11649,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_metadirective_variant *mc
+ = c->ext.omp_metadirective_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)
@@ -11595,6 +11801,9 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns)
code->ext.omp_clauses->if_present = false;
resolve_omp_clauses (code, code->ext.omp_clauses, ns);
break;
+ case EXEC_OMP_METADIRECTIVE:
+ resolve_omp_metadirective (code, ns);
+ break;
default:
break;
}
@@ -48,6 +48,10 @@ gfc_state_data *gfc_state_stack;
static bool last_was_use_stmt = false;
bool in_exec_part;
+bool gfc_matching_omp_context_selector;
+bool gfc_in_metadirective_body;
+int gfc_omp_region_count;
+
/* TODO: Re-order functions to kill these forward decls. */
static void check_statement_label (gfc_statement);
static void undo_new_statement (void);
@@ -1041,6 +1045,8 @@ decode_omp_directive (void)
break;
case 'b':
matcho ("barrier", gfc_match_omp_barrier, ST_OMP_BARRIER);
+ matcho ("begin metadirective", gfc_match_omp_begin_metadirective,
+ ST_OMP_BEGIN_METADIRECTIVE);
break;
case 'c':
matcho ("cancellation% point", gfc_match_omp_cancellation_point,
@@ -1085,6 +1091,8 @@ decode_omp_directive (void)
matcho ("end master taskloop", gfc_match_omp_eos_error,
ST_OMP_END_MASTER_TASKLOOP);
matcho ("end master", gfc_match_omp_eos_error, ST_OMP_END_MASTER);
+ matcho ("end metadirective", gfc_match_omp_eos_error,
+ ST_OMP_END_METADIRECTIVE);
matchs ("end ordered", gfc_match_omp_eos_error, ST_OMP_END_ORDERED);
matchs ("end parallel do simd", gfc_match_omp_eos_error,
ST_OMP_END_PARALLEL_DO_SIMD);
@@ -1168,6 +1176,8 @@ decode_omp_directive (void)
matcho ("master taskloop", gfc_match_omp_master_taskloop,
ST_OMP_MASTER_TASKLOOP);
matcho ("master", gfc_match_omp_master, ST_OMP_MASTER);
+ matcho ("metadirective", gfc_match_omp_metadirective,
+ ST_OMP_METADIRECTIVE);
break;
case 'n':
matcho ("nothing", gfc_match_omp_nothing, ST_NONE);
@@ -1296,6 +1306,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 ();
@@ -1413,6 +1427,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)
{
@@ -1935,6 +1955,43 @@ 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
+
/* Block end statements. Errors associated with interchanging these
are detected in gfc_match_end(). */
@@ -2572,6 +2629,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;
@@ -2671,6 +2731,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;
@@ -2815,6 +2878,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";
@@ -3075,6 +3141,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 ();
@@ -5386,6 +5454,140 @@ 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;
+ default:
+ break;
+ }
+ }
+
+ if (omp_structured_p)
+ {
+ switch (omp_st)
+ {
+ case ST_OMP_ASSUME:
+ return ST_OMP_END_ASSUME;
+ 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. */
@@ -5436,94 +5638,16 @@ parse_omp_do (gfc_statement omp_st)
pop_state ();
st = next_statement ();
- 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;
- 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)
@@ -5835,80 +5959,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;
@@ -5954,11 +6013,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:
@@ -6131,6 +6192,88 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
return st;
}
+static gfc_statement
+parse_omp_metadirective_body (gfc_statement omp_st)
+{
+ gfc_omp_metadirective_variant *variant
+ = new_st.ext.omp_metadirective_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_metadirective_body;
+ gfc_in_metadirective_body = true;
+
+ gfc_omp_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);
+ /* TODO: Does st == ST_IMPLIED_ENDDO need special handling? */
+ 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 an OMP METADIRECTIVE block at %C",
+ gfc_ascii_statement (st));
+ reject_statement ();
+ st = next_statement ();
+ }
+ finish:
+
+ gfc_in_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
@@ -6141,6 +6284,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)
@@ -6148,6 +6292,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)
@@ -6257,68 +6407,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_omp_do:
st = parse_omp_do (st);
if (st == ST_IMPLIED_ENDDO)
return st;
@@ -6332,6 +6427,19 @@ parse_executable (gfc_statement st)
st = parse_omp_oacc_atomic (true);
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)
+ {
+ st = next_statement ();
+ return st;
+ }
+ /* FALLTHRU */
+
default:
return st;
}
@@ -7097,6 +7205,10 @@ gfc_parse_file (void)
gfc_statement_label = NULL;
+ gfc_omp_region_count = 0;
+ gfc_in_metadirective_body = false;
+ gfc_matching_omp_context_selector = false;
+
if (setjmp (eof_buf))
return false; /* Come here on unexpected EOF */
@@ -7409,3 +7521,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;
+ }
+}
@@ -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_metadirective_body;
+extern int gfc_omp_region_count;
+
match gfc_match_prefix (gfc_typespec *);
bool is_oacc (gfc_state_data *);
#endif /* GFC_PARSE_H */
@@ -12261,6 +12261,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_metadirective_variant *variant
+ = code->ext.omp_metadirective_variants;
+ variant; variant = variant->next)
+ gfc_resolve_code (variant->code, ns);
else if (code->block)
{
omp_workshare_save = -1;
@@ -12793,6 +12798,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:
@@ -299,6 +299,10 @@ gfc_free_statement (gfc_code *p)
case EXEC_OMP_TASKYIELD:
break;
+ case EXEC_OMP_METADIRECTIVE:
+ gfc_free_omp_metadirective_variants (p->ext.omp_metadirective_variants);
+ break;
+
default:
gfc_internal_error ("gfc_free_statement(): Bad statement");
}
@@ -2630,10 +2630,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;
}
@@ -2683,6 +2686,7 @@ gfc_get_st_label (int labelno)
{
gfc_st_label *lp;
gfc_namespace *ns;
+ int omp_region = gfc_in_metadirective_body ? gfc_omp_region_count : 0;
if (gfc_current_state () == COMP_DERIVED)
ns = gfc_current_block ()->f2k_derived;
@@ -2699,10 +2703,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;
@@ -2714,6 +2724,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);
@@ -327,7 +327,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));
@@ -8228,6 +8228,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:
@@ -8319,6 +8321,104 @@ gfc_trans_omp_declare_simd (gfc_namespace *ns)
}
}
+/* Translate the context selector list GFC_SELECTORS, using WHERE as the
+ locus for error messages. CONV_EXPR_P is true for normal translation
+ and false for early conversion during parsing. */
+
+static tree
+gfc_trans_omp_set_selector (gfc_omp_set_selector *gfc_selectors, locus where,
+ bool conv_expr_p)
+{
+ 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_EXPR:
+ {
+ tree expr = NULL_TREE;
+ if (conv_expr_p)
+ {
+ 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)
{
@@ -8394,89 +8494,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_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, true);
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)
@@ -8577,3 +8596,57 @@ 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_metadirective_variant *variant = code->ext.omp_metadirective_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, true);
+ ctx = omp_check_context_selector (gfc_get_location (&variant->where),
+ ctx, true);
+ if (ctx == error_mark_node)
+ return error_mark_node;
+
+ 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;
+}
+
+/* Hook to check during parsing whether a metadirective is both resolvable
+ and non-matching; if so, it can be eliminated immediately. */
+bool
+gfc_skip_omp_metadirective_variant (gfc_omp_metadirective_variant *variant)
+{
+ tree selector = gfc_trans_omp_set_selector (variant->selectors,
+ variant->where, false);
+ return omp_context_selector_matches (selector, true, true) == 0;
+}
@@ -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 *);
@@ -2611,6 +2611,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:
new file mode 100644
@@ -0,0 +1,55 @@
+! { 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) & ! { dg-error "expected 'when', 'otherwise', or 'default' at .1." }
+ !$omp& where (device={arch("nvptx")}: parallel loop)
+ 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
new file mode 100644
@@ -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 }
new file mode 100644
@@ -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
+
+
new file mode 100644
@@ -0,0 +1,62 @@
+! { 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 empty metadirective.
+ !$omp metadirective
+end program
new file mode 100644
@@ -0,0 +1,34 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-gimple" }
+! { dg-additional-options "-fdump-tree-optimized" }
+
+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
+
+! The metadirective should be resolved after Gimplification.
+
+! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } }
+! { dg-final { scan-tree-dump-times "when \\(device =.*arch.*nvptx.*\\):" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } }
+! { dg-final { scan-tree-dump-times "default:" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } }
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } }
new file mode 100644
@@ -0,0 +1,39 @@
+! { 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
+
+ ! TODO: This does not execute a version of f with the default clause
+ ! active as might be expected.
+ call f (a, 2.71828) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" }
+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 "default:" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } }
+
+! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } }
new file mode 100644
@@ -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" } }
new file mode 100644
@@ -0,0 +1,31 @@
+! { 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
+
+! The outer metadirective should be resolved at parse time, but is
+! currently resolved during Gimplification.
+
+! The inner metadirective should be resolved during Gimplificiation.
+
+! { 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" } }
new file mode 100644
@@ -0,0 +1,36 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ integer, parameter :: N = 256
+contains
+ subroutine f (a, num)
+ 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
+end program
+
+! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"gpu\"\\\[0\\\], &\"amdgcn\"\\\[0\\\], &\"gfx906\"\\\[0\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } }
+! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(.+, &\"cpu\"\\\[0\\\], &\"x86_64\"\\\[0\\\], 0B\\)" "gimple" } }
+! { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(-1, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } }
new file mode 100644
@@ -0,0 +1,22 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+program test
+ integer :: i
+ integer, parameter :: N = 100
+ integer :: sum = 0
+
+ ! The compiler should never consider a situation where both metadirectives
+ ! match. If it does, then the nested metadirective would be an error
+ ! as it is not a loop-nest as per the OpenMP specification.
+
+ !$omp metadirective when (implementation={vendor("ibm")}: &
+ !$omp& target teams distribute)
+ !$omp metadirective when (implementation={vendor("gnu")}: parallel do)
+ do i = 1, N
+ sum = sum + i
+ end do
+end program
+
+! { dg-final { scan-tree-dump-not "when \\(implementation = .*vendor.*ibm.*\\):" "original" } }
+! { dg-final { scan-tree-dump-times "when \\(implementation = .*vendor.*gnu.*\\):" 1 "original" } }
new file mode 100644
@@ -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 an OMP METADIRECTIVE block at .1." }
+
+
+end program
+
new file mode 100644
@@ -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" } }
new file mode 100644
@@ -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
@@ -86,3 +86,10 @@ pure integer function func_simd(n)
end do
func_simd = r
end
+
+!pure logical function func_metadirective()
+logical function func_metadirective()
+ implicit none
+ !$omp metadirective
+ func_metadirective = .false.
+end
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -0,0 +1,40 @@
+! { dg-do run }
+
+program test
+ implicit none
+ integer, parameter :: N = 100
+ real, parameter :: PI_CONST = 3.14159
+ real, parameter :: E_CONST = 2.71828
+ 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
+
+ ! TODO: This does not execute a version of f with the default clause
+ ! active as might be expected.
+ call f (a, E_CONST) ! { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" }
+
+ 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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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