+2014-06-18 Jakub Jelinek <jakub@redhat.com>
+
+ * gimplify.c (omp_notice_variable): If n is non-NULL
+ and no flags change in ORT_TARGET region, don't jump to
+ do_outer.
+ (struct gimplify_adjust_omp_clauses_data): New type.
+ (gimplify_adjust_omp_clauses_1): Adjust for data being
+ a struct gimplify_adjust_omp_clauses_data pointer instead
+ of tree *. Pass pre_p as a new argument to
+ lang_hooks.decls.omp_finish_clause hook.
+ (gimplify_adjust_omp_clauses): Add pre_p argument, adjust
+ splay_tree_foreach to pass both list_p and pre_p.
+ (gimplify_omp_parallel, gimplify_omp_task, gimplify_omp_for,
+ gimplify_omp_workshare, gimplify_omp_target_update): Adjust
+ gimplify_adjust_omp_clauses callers.
+ * langhooks.c (lhd_omp_finish_clause): New function.
+ * langhooks-def.h (lhd_omp_finish_clause): New prototype.
+ (LANG_HOOKS_OMP_FINISH_CLAUSE): Define to lhd_omp_finish_clause.
+ * langhooks.h (struct lang_hooks_for_decls): Add a new
+ gimple_seq * argument to omp_finish_clause hook.
+ * omp-low.c (scan_sharing_clauses): Call scan_omp_op on
+ non-DECL_P OMP_CLAUSE_DECL if ctx->outer.
+ (scan_omp_parallel, lower_omp_for): When adding
+ _LOOPTEMP_ clause var, add it to outer ctx's decl_map
+ as identity.
+ * tree-core.h (OMP_CLAUSE_MAP_TO_PSET): New map kind.
+ * tree-nested.c (convert_nonlocal_omp_clauses,
+ convert_local_omp_clauses): Handle various OpenMP 4.0 clauses.
+ * tree-pretty-print.c (dump_omp_clause): Handle
+ OMP_CLAUSE_MAP_TO_PSET.
+
2014-06-17 Andrew MacLeod <amacleod@redhat.com>
* tree-dfa.h (get_addr_base_and_unit_offset_1): Move from here.
+2014-06-18 Jakub Jelinek <jakub@redhat.com>
+
+ * cp-gimplify.c (cxx_omp_finish_clause): Add a gimple_seq *
+ argument.
+ * cp-tree.h (cxx_omp_finish_clause): Adjust prototype.
+
2014-06-17 Jason Merrill <jason@redhat.com>
PR c++/60605
/* Finalize an implicitly determined clause. */
void
-cxx_omp_finish_clause (tree c)
+cxx_omp_finish_clause (tree c, gimple_seq *)
{
tree decl, inner_type;
bool make_shared = false;
extern tree cxx_omp_clause_copy_ctor (tree, tree, tree);
extern tree cxx_omp_clause_assign_op (tree, tree, tree);
extern tree cxx_omp_clause_dtor (tree, tree);
-extern void cxx_omp_finish_clause (tree);
+extern void cxx_omp_finish_clause (tree, gimple_seq *);
extern bool cxx_omp_privatize_by_reference (const_tree);
/* in name-lookup.c */
+2014-06-18 Jakub Jelinek <jakub@redhat.com>
+
+ * cpp.c (cpp_define_builtins): Change _OPENMP macro to
+ 201307.
+ * dump-parse-tree.c (show_omp_namelist): Add list_type
+ argument. Adjust for rop being u.reduction_op now,
+ handle depend_op or map_op.
+ (show_omp_node): Adjust callers. Print some new
+ OpenMP 4.0 clauses, adjust for OMP_LIST_DEPEND_{IN,OUT}
+ becoming a single OMP_LIST_DEPEND.
+ * f95-lang.c (gfc_handle_omp_declare_target_attribute): New
+ function.
+ (gfc_attribute_table): New variable.
+ (LANG_HOOKS_OMP_FINISH_CLAUSE, LANG_HOOKS_ATTRIBUTE_TABLE): Redefine.
+ * frontend-passes.c (gfc_code_walker): Handle new OpenMP target
+ EXEC_OMP_* codes and new clauses.
+ * gfortran.h (gfc_statement): Add ST_OMP_TARGET, ST_OMP_END_TARGET,
+ ST_OMP_TARGET_DATA, ST_OMP_END_TARGET_DATA, ST_OMP_TARGET_UPDATE,
+ ST_OMP_DECLARE_TARGET, ST_OMP_TEAMS, ST_OMP_END_TEAMS,
+ ST_OMP_DISTRIBUTE, ST_OMP_END_DISTRIBUTE, ST_OMP_DISTRIBUTE_SIMD,
+ ST_OMP_END_DISTRIBUTE_SIMD, ST_OMP_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_END_DISTRIBUTE_PARALLEL_DO, ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
+ ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD, ST_OMP_TARGET_TEAMS,
+ ST_OMP_END_TARGET_TEAMS, ST_OMP_TEAMS_DISTRIBUTE,
+ ST_OMP_END_TEAMS_DISTRIBUTE, ST_OMP_TEAMS_DISTRIBUTE_SIMD,
+ ST_OMP_END_TEAMS_DISTRIBUTE_SIMD, ST_OMP_TARGET_TEAMS_DISTRIBUTE,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD,
+ ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD and
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD.
+ (symbol_attribute): Add omp_declare_target field.
+ (gfc_omp_depend_op, gfc_omp_map_op): New enums.
+ (gfc_omp_namelist): Replace rop field with union
+ containing reduction_op, depend_op and map_op.
+ (OMP_LIST_DEPEND_IN, OMP_LIST_DEPEND_OUT): Remove.
+ (OMP_LIST_DEPEND, OMP_LIST_MAP, OMP_LIST_TO, OMP_LIST_FROM): New.
+ (gfc_omp_clauses): Add num_teams, device, thread_limit,
+ dist_sched_kind, dist_chunk_size fields.
+ (gfc_common_head): Add omp_declare_target field.
+ (gfc_exec_op): Add EXEC_OMP_TARGET, EXEC_OMP_TARGET_DATA,
+ EXEC_OMP_TEAMS, EXEC_OMP_DISTRIBUTE, EXEC_OMP_DISTRIBUTE_SIMD,
+ EXEC_OMP_DISTRIBUTE_PARALLEL_DO, EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
+ EXEC_OMP_TARGET_TEAMS, EXEC_OMP_TEAMS_DISTRIBUTE,
+ EXEC_OMP_TEAMS_DISTRIBUTE_SIMD, EXEC_OMP_TARGET_TEAMS_DISTRIBUTE,
+ EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+ EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD and
+ EXEC_OMP_TARGET_UPDATE.
+ (gfc_add_omp_declare_target): New prototype.
+ * match.h (gfc_match_omp_declare_target, gfc_match_omp_distribute,
+ gfc_match_omp_distribute_parallel_do,
+ gfc_match_omp_distribute_parallel_do_simd,
+ gfc_match_omp_distribute_simd, gfc_match_omp_target,
+ gfc_match_omp_target_data, gfc_match_omp_target_teams,
+ gfc_match_omp_target_teams_distribute,
+ gfc_match_omp_target_teams_distribute_parallel_do,
+ gfc_match_omp_target_teams_distribute_parallel_do_simd,
+ gfc_match_omp_target_teams_distribute_simd,
+ gfc_match_omp_target_update, gfc_match_omp_teams,
+ gfc_match_omp_teams_distribute,
+ gfc_match_omp_teams_distribute_parallel_do,
+ gfc_match_omp_teams_distribute_parallel_do_simd,
+ gfc_match_omp_teams_distribute_simd): New prototypes.
+ * module.c (ab_attribute): Add AB_OMP_DECLARE_TARGET.
+ (attr_bits): Likewise.
+ (mio_symbol_attribute): Handle omp_declare_target attribute.
+ (gfc_free_omp_clauses): Free num_teams, device, thread_limit
+ and dist_chunk_size expressions.
+ (OMP_CLAUSE_PRIVATE, OMP_CLAUSE_FIRSTPRIVATE, OMP_CLAUSE_LASTPRIVATE,
+ OMP_CLAUSE_COPYPRIVATE, OMP_CLAUSE_SHARED, OMP_CLAUSE_COPYIN,
+ OMP_CLAUSE_REDUCTION, OMP_CLAUSE_IF, OMP_CLAUSE_NUM_THREADS,
+ OMP_CLAUSE_SCHEDULE, OMP_CLAUSE_DEFAULT, OMP_CLAUSE_ORDERED,
+ OMP_CLAUSE_COLLAPSE, OMP_CLAUSE_UNTIED, OMP_CLAUSE_FINAL,
+ OMP_CLAUSE_MERGEABLE, OMP_CLAUSE_ALIGNED, OMP_CLAUSE_DEPEND,
+ OMP_CLAUSE_INBRANCH, OMP_CLAUSE_LINEAR, OMP_CLAUSE_NOTINBRANCH,
+ OMP_CLAUSE_PROC_BIND, OMP_CLAUSE_SAFELEN, OMP_CLAUSE_SIMDLEN,
+ OMP_CLAUSE_UNIFORM): Use 1U instead of 1.
+ (OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP, OMP_CLAUSE_TO, OMP_CLAUSE_FROM,
+ OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_THREAD_LIMIT,
+ OMP_CLAUSE_DIST_SCHEDULE): Define.
+ (gfc_match_omp_clauses): Change mask parameter to unsigned int.
+ Adjust for rop becoming u.reduction_op. Disallow inbranch with
+ notinbranch. For depend clause, always create OMP_LIST_DEPEND
+ and fill in u.depend_op. Handle num_teams, device, map,
+ to, from, thread_limit and dist_schedule clauses.
+ (OMP_DECLARE_SIMD_CLAUSES): Or in OMP_CLAUSE_INBRANCH and
+ OMP_CLAUSE_NOTINBRANCH.
+ (OMP_TARGET_CLAUSES, OMP_TARGET_DATA_CLAUSES,
+ OMP_TARGET_UPDATE_CLAUSES, OMP_TEAMS_CLAUSES,
+ OMP_DISTRIBUTE_CLAUSES): Define.
+ (match_omp): New function.
+ (gfc_match_omp_do, gfc_match_omp_do_simd, gfc_match_omp_parallel,
+ gfc_match_omp_parallel_do, gfc_match_omp_parallel_do_simd,
+ gfc_match_omp_parallel_sections, gfc_match_omp_parallel_workshare,
+ gfc_match_omp_sections, gfc_match_omp_simd, gfc_match_omp_single,
+ gfc_match_omp_task): Rewritten using match_omp.
+ (gfc_match_omp_threadprivate, gfc_match_omp_declare_reduction):
+ Diagnose if the directives are followed by unexpected junk.
+ (gfc_match_omp_distribute, gfc_match_omp_distribute_parallel_do,
+ gfc_match_omp_distribute_parallel_do_simd,
+ gfc_match_omp_distrbute_simd, gfc_match_omp_declare_target,
+ gfc_match_omp_target, gfc_match_omp_target_data,
+ gfc_match_omp_target_teams, gfc_match_omp_target_teams_distribute,
+ gfc_match_omp_target_teams_distribute_parallel_do,
+ gfc_match_omp_target_teams_distribute_parallel_do_simd,
+ gfc_match_omp_target_teams_distrbute_simd, gfc_match_omp_target_update,
+ gfc_match_omp_teams, gfc_match_omp_teams_distribute,
+ gfc_match_omp_teams_distribute_parallel_do,
+ gfc_match_omp_teams_distribute_parallel_do_simd,
+ gfc_match_omp_teams_distrbute_simd): New functions.
+ * openmp.c (resolve_omp_clauses): Adjust for
+ OMP_LIST_DEPEND_{IN,OUT} being changed to OMP_LIST_DEPEND. Handle
+ OMP_LIST_MAP, OMP_LIST_FROM, OMP_LIST_TO, num_teams, device,
+ dist_chunk_size and thread_limit.
+ (gfc_resolve_omp_parallel_blocks): Only put sharing clauses into
+ ctx.sharing_clauses. Call gfc_resolve_omp_do_blocks for various
+ new EXEC_OMP_* codes.
+ (resolve_omp_do): Handle various new EXEC_OMP_* codes.
+ (gfc_resolve_omp_directive): Likewise.
+ (gfc_resolve_omp_declare_simd): Add missing space to diagnostics.
+ * parse.c (decode_omp_directive): Handle parsing of OpenMP 4.0
+ offloading related directives.
+ (case_executable): Add ST_OMP_TARGET_UPDATE.
+ (case_exec_markers): Add ST_OMP_TARGET*, ST_OMP_TEAMS*,
+ ST_OMP_DISTRIBUTE*.
+ (case_decl): Add ST_OMP_DECLARE_TARGET.
+ (gfc_ascii_statement): Handle new ST_OMP_* codes.
+ (parse_omp_do): Handle various new ST_OMP_* codes.
+ (parse_executable): Likewise.
+ * resolve.c (gfc_resolve_blocks): Handle various new EXEC_OMP_*
+ codes.
+ (resolve_code): Likewise.
+ (resolve_symbol): Change that !$OMP DECLARE TARGET variables
+ are saved.
+ * st.c (gfc_free_statement): Handle various new EXEC_OMP_* codes.
+ * symbol.c (check_conflict): Check omp_declare_target conflicts.
+ (gfc_add_omp_declare_target): New function.
+ (gfc_copy_attr): Copy omp_declare_target.
+ * trans.c (trans_code): Handle various new EXEC_OMP_* codes.
+ * trans-common.c (build_common_decl): Add "omp declare target"
+ attribute if needed.
+ * trans-decl.c (add_attributes_to_decl): Likewise.
+ * trans.h (gfc_omp_finish_clause): New prototype.
+ * trans-openmp.c (gfc_omp_finish_clause): New function.
+ (gfc_trans_omp_reduction_list): Adjust for rop being renamed
+ to u.reduction_op.
+ (gfc_trans_omp_clauses): Adjust for OMP_LIST_DEPEND_{IN,OUT}
+ change to OMP_LIST_DEPEND and fix up depend handling.
+ Handle OMP_LIST_MAP, OMP_LIST_TO, OMP_LIST_FROM, num_teams,
+ thread_limit, device, dist_chunk_size and dist_sched_kind.
+ (gfc_trans_omp_do): Handle EXEC_OMP_DISTRIBUTE.
+ (GFC_OMP_SPLIT_DISTRIBUTE, GFC_OMP_SPLIT_TEAMS,
+ GFC_OMP_SPLIT_TARGET, GFC_OMP_SPLIT_NUM, GFC_OMP_MASK_DISTRIBUTE,
+ GFC_OMP_MASK_TEAMS, GFC_OMP_MASK_TARGET, GFC_OMP_MASK_NUM): New.
+ (gfc_split_omp_clauses): Handle splitting of clauses for new
+ EXEC_OMP_* codes.
+ (gfc_trans_omp_do_simd): Add pblock argument, adjust for being
+ callable for combined constructs.
+ (gfc_trans_omp_parallel_do, gfc_trans_omp_parallel_do_simd): Likewise.
+ (gfc_trans_omp_distribute, gfc_trans_omp_teams,
+ gfc_trans_omp_target, gfc_trans_omp_target_data,
+ gfc_trans_omp_target_update): New functions.
+ (gfc_trans_omp_directive): Adjust gfc_trans_omp_* callers, handle
+ new EXEC_OMP_* codes.
+
2014-06-18 Tobias Burnus <burnus@net-b.de>
PR fortran/61126
cpp_define (pfile, "_LANGUAGE_FORTRAN=1");
if (gfc_option.gfc_flag_openmp)
- cpp_define (pfile, "_OPENMP=201107");
+ cpp_define (pfile, "_OPENMP=201307");
/* The defines below are necessary for the TARGET_* macros.
}
static void
-show_omp_namelist (gfc_omp_namelist *n)
+show_omp_namelist (int list_type, gfc_omp_namelist *n)
{
for (; n; n = n->next)
{
- switch (n->rop)
- {
- case OMP_REDUCTION_PLUS:
- case OMP_REDUCTION_TIMES:
- case OMP_REDUCTION_MINUS:
- case OMP_REDUCTION_AND:
- case OMP_REDUCTION_OR:
- case OMP_REDUCTION_EQV:
- case OMP_REDUCTION_NEQV:
- fprintf (dumpfile, "%s:", gfc_op2string ((gfc_intrinsic_op) n->rop));
- break;
- case OMP_REDUCTION_MAX: fputs ("max:", dumpfile); break;
- case OMP_REDUCTION_MIN: fputs ("min:", dumpfile); break;
- case OMP_REDUCTION_IAND: fputs ("iand:", dumpfile); break;
- case OMP_REDUCTION_IOR: fputs ("ior:", dumpfile); break;
- case OMP_REDUCTION_IEOR: fputs ("ieor:", dumpfile); break;
- case OMP_REDUCTION_USER:
- if (n->udr)
- fprintf (dumpfile, "%s:", n->udr->name);
- break;
- default: break;
- }
+ if (list_type == OMP_LIST_REDUCTION)
+ switch (n->u.reduction_op)
+ {
+ case OMP_REDUCTION_PLUS:
+ case OMP_REDUCTION_TIMES:
+ case OMP_REDUCTION_MINUS:
+ case OMP_REDUCTION_AND:
+ case OMP_REDUCTION_OR:
+ case OMP_REDUCTION_EQV:
+ case OMP_REDUCTION_NEQV:
+ fprintf (dumpfile, "%s:",
+ gfc_op2string ((gfc_intrinsic_op) n->u.reduction_op));
+ break;
+ case OMP_REDUCTION_MAX: fputs ("max:", dumpfile); break;
+ case OMP_REDUCTION_MIN: fputs ("min:", dumpfile); break;
+ case OMP_REDUCTION_IAND: fputs ("iand:", dumpfile); break;
+ case OMP_REDUCTION_IOR: fputs ("ior:", dumpfile); break;
+ case OMP_REDUCTION_IEOR: fputs ("ieor:", dumpfile); break;
+ case OMP_REDUCTION_USER:
+ if (n->udr)
+ fprintf (dumpfile, "%s:", n->udr->name);
+ break;
+ default: break;
+ }
+ else if (list_type == OMP_LIST_DEPEND)
+ switch (n->u.depend_op)
+ {
+ case OMP_DEPEND_IN: fputs ("in:", dumpfile); break;
+ case OMP_DEPEND_OUT: fputs ("out:", dumpfile); break;
+ case OMP_DEPEND_INOUT: fputs ("inout:", dumpfile); break;
+ default: break;
+ }
+ else if (list_type == OMP_LIST_MAP)
+ switch (n->u.map_op)
+ {
+ case OMP_MAP_ALLOC: fputs ("alloc:", dumpfile); break;
+ case OMP_MAP_TO: fputs ("to:", dumpfile); break;
+ case OMP_MAP_FROM: fputs ("from:", dumpfile); break;
+ case OMP_MAP_TOFROM: fputs ("tofrom:", dumpfile); break;
+ default: break;
+ }
fprintf (dumpfile, "%s", n->sym->name);
if (n->expr)
{
if (c->ext.omp_namelist)
{
fputs (" (", dumpfile);
- show_omp_namelist (c->ext.omp_namelist);
+ show_omp_namelist (OMP_LIST_NUM, c->ext.omp_namelist);
fputc (')', dumpfile);
}
return;
case OMP_LIST_ALIGNED: type = "ALIGNED"; break;
case OMP_LIST_LINEAR: type = "LINEAR"; break;
case OMP_LIST_REDUCTION: type = "REDUCTION"; break;
- case OMP_LIST_DEPEND_IN:
- fprintf (dumpfile, " DEPEND(IN:");
- break;
- case OMP_LIST_DEPEND_OUT:
- fprintf (dumpfile, " DEPEND(OUT:");
- break;
+ case OMP_LIST_DEPEND: type = "DEPEND"; break;
default:
gcc_unreachable ();
}
- if (type)
- fprintf (dumpfile, " %s(", type);
- show_omp_namelist (omp_clauses->lists[list_type]);
+ fprintf (dumpfile, " %s(", type);
+ show_omp_namelist (list_type, omp_clauses->lists[list_type]);
fputc (')', dumpfile);
}
if (omp_clauses->safelen_expr)
}
fprintf (dumpfile, " PROC_BIND(%s)", type);
}
+ if (omp_clauses->num_teams)
+ {
+ fputs (" NUM_TEAMS(", dumpfile);
+ show_expr (omp_clauses->num_teams);
+ fputc (')', dumpfile);
+ }
+ if (omp_clauses->device)
+ {
+ fputs (" DEVICE(", dumpfile);
+ show_expr (omp_clauses->device);
+ fputc (')', dumpfile);
+ }
+ if (omp_clauses->thread_limit)
+ {
+ fputs (" THREAD_LIMIT(", dumpfile);
+ show_expr (omp_clauses->thread_limit);
+ fputc (')', dumpfile);
+ }
+ if (omp_clauses->dist_sched_kind != OMP_SCHED_NONE)
+ {
+ fprintf (dumpfile, " DIST_SCHEDULE (static");
+ if (omp_clauses->dist_chunk_size)
+ {
+ fputc (',', dumpfile);
+ show_expr (omp_clauses->dist_chunk_size);
+ }
+ fputc (')', dumpfile);
+ }
}
fputc ('\n', dumpfile);
if (c->op == EXEC_OMP_SECTIONS || c->op == EXEC_OMP_PARALLEL_SECTIONS)
if (omp_clauses->lists[OMP_LIST_COPYPRIVATE])
{
fputs (" COPYPRIVATE(", dumpfile);
- show_omp_namelist (omp_clauses->lists[OMP_LIST_COPYPRIVATE]);
+ show_omp_namelist (OMP_LIST_COPYPRIVATE,
+ omp_clauses->lists[OMP_LIST_COPYPRIVATE]);
fputc (')', dumpfile);
}
else if (omp_clauses->nowait)
static void gfc_init_ts (void);
static tree gfc_builtin_function (tree);
+/* Handle an "omp declare target" attribute; arguments as in
+ struct attribute_spec.handler. */
+static tree
+gfc_handle_omp_declare_target_attribute (tree *, tree, tree, int, bool *)
+{
+ return NULL_TREE;
+}
+
+/* Table of valid Fortran attributes. */
+static const struct attribute_spec gfc_attribute_table[] =
+{
+ /* { name, min_len, max_len, decl_req, type_req, fn_type_req, handler,
+ affects_type_identity } */
+ { "omp declare target", 0, 0, true, false, false,
+ gfc_handle_omp_declare_target_attribute, false },
+ { NULL, 0, 0, false, false, false, NULL, false }
+};
+
#undef LANG_HOOKS_NAME
#undef LANG_HOOKS_INIT
#undef LANG_HOOKS_FINISH
#undef LANG_HOOKS_OMP_CLAUSE_COPY_CTOR
#undef LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP
#undef LANG_HOOKS_OMP_CLAUSE_DTOR
+#undef LANG_HOOKS_OMP_FINISH_CLAUSE
#undef LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR
#undef LANG_HOOKS_OMP_PRIVATE_DEBUG_CLAUSE
#undef LANG_HOOKS_OMP_PRIVATE_OUTER_REF
#undef LANG_HOOKS_BUILTIN_FUNCTION
#undef LANG_HOOKS_BUILTIN_FUNCTION
#undef LANG_HOOKS_GET_ARRAY_DESCR_INFO
+#undef LANG_HOOKS_ATTRIBUTE_TABLE
/* Define lang hooks. */
#define LANG_HOOKS_NAME "GNU Fortran"
#define LANG_HOOKS_OMP_CLAUSE_COPY_CTOR gfc_omp_clause_copy_ctor
#define LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP gfc_omp_clause_assign_op
#define LANG_HOOKS_OMP_CLAUSE_DTOR gfc_omp_clause_dtor
+#define LANG_HOOKS_OMP_FINISH_CLAUSE gfc_omp_finish_clause
#define LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR gfc_omp_disregard_value_expr
#define LANG_HOOKS_OMP_PRIVATE_DEBUG_CLAUSE gfc_omp_private_debug_clause
#define LANG_HOOKS_OMP_PRIVATE_OUTER_REF gfc_omp_private_outer_ref
#define LANG_HOOKS_OMP_FIRSTPRIVATIZE_TYPE_SIZES \
gfc_omp_firstprivatize_type_sizes
-#define LANG_HOOKS_BUILTIN_FUNCTION gfc_builtin_function
-#define LANG_HOOKS_GET_ARRAY_DESCR_INFO gfc_get_array_descr_info
+#define LANG_HOOKS_BUILTIN_FUNCTION gfc_builtin_function
+#define LANG_HOOKS_GET_ARRAY_DESCR_INFO gfc_get_array_descr_info
+#define LANG_HOOKS_ATTRIBUTE_TABLE gfc_attribute_table
struct lang_hooks lang_hooks = LANG_HOOKS_INITIALIZER;
in_omp_workshare = true;
/* Fall through */
-
+
+ case EXEC_OMP_DISTRIBUTE:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_DISTRIBUTE_SIMD:
case EXEC_OMP_DO:
case EXEC_OMP_DO_SIMD:
case EXEC_OMP_SECTIONS:
case EXEC_OMP_SINGLE:
case EXEC_OMP_END_SINGLE:
case EXEC_OMP_SIMD:
+ case EXEC_OMP_TARGET:
+ case EXEC_OMP_TARGET_DATA:
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TARGET_UPDATE:
case EXEC_OMP_TASK:
+ case EXEC_OMP_TEAMS:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
/* Come to this label only from the
EXEC_OMP_PARALLEL_* cases above. */
if (co->ext.omp_clauses)
{
+ gfc_omp_namelist *n;
+ static int list_types[]
+ = { OMP_LIST_ALIGNED, OMP_LIST_LINEAR, OMP_LIST_DEPEND,
+ OMP_LIST_MAP, OMP_LIST_TO, OMP_LIST_FROM };
+ size_t idx;
WALK_SUBEXPR (co->ext.omp_clauses->if_expr);
WALK_SUBEXPR (co->ext.omp_clauses->final_expr);
WALK_SUBEXPR (co->ext.omp_clauses->num_threads);
WALK_SUBEXPR (co->ext.omp_clauses->chunk_size);
WALK_SUBEXPR (co->ext.omp_clauses->safelen_expr);
WALK_SUBEXPR (co->ext.omp_clauses->simdlen_expr);
+ WALK_SUBEXPR (co->ext.omp_clauses->num_teams);
+ WALK_SUBEXPR (co->ext.omp_clauses->device);
+ WALK_SUBEXPR (co->ext.omp_clauses->thread_limit);
+ WALK_SUBEXPR (co->ext.omp_clauses->dist_chunk_size);
+ for (idx = 0;
+ idx < sizeof (list_types) / sizeof (list_types[0]);
+ idx++)
+ for (n = co->ext.omp_clauses->lists[list_types[idx]];
+ n; n = n->next)
+ WALK_SUBEXPR (n->expr);
}
- {
- gfc_omp_namelist *n;
- for (n = co->ext.omp_clauses->lists[OMP_LIST_ALIGNED];
- n; n = n->next)
- WALK_SUBEXPR (n->expr);
- for (n = co->ext.omp_clauses->lists[OMP_LIST_LINEAR];
- n; n = n->next)
- WALK_SUBEXPR (n->expr);
- for (n = co->ext.omp_clauses->lists[OMP_LIST_DEPEND_IN];
- n; n = n->next)
- WALK_SUBEXPR (n->expr);
- for (n = co->ext.omp_clauses->lists[OMP_LIST_DEPEND_OUT];
- n; n = n->next)
- WALK_SUBEXPR (n->expr);
- }
break;
default:
break;
ST_OMP_TASKGROUP, ST_OMP_END_TASKGROUP, ST_OMP_SIMD, ST_OMP_END_SIMD,
ST_OMP_DO_SIMD, ST_OMP_END_DO_SIMD, ST_OMP_PARALLEL_DO_SIMD,
ST_OMP_END_PARALLEL_DO_SIMD, ST_OMP_DECLARE_SIMD, ST_OMP_DECLARE_REDUCTION,
+ ST_OMP_TARGET, ST_OMP_END_TARGET, ST_OMP_TARGET_DATA, ST_OMP_END_TARGET_DATA,
+ ST_OMP_TARGET_UPDATE, ST_OMP_DECLARE_TARGET,
+ ST_OMP_TEAMS, ST_OMP_END_TEAMS, ST_OMP_DISTRIBUTE, ST_OMP_END_DISTRIBUTE,
+ ST_OMP_DISTRIBUTE_SIMD, ST_OMP_END_DISTRIBUTE_SIMD,
+ ST_OMP_DISTRIBUTE_PARALLEL_DO, ST_OMP_END_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD, ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD,
+ ST_OMP_TARGET_TEAMS, ST_OMP_END_TARGET_TEAMS, ST_OMP_TEAMS_DISTRIBUTE,
+ ST_OMP_END_TEAMS_DISTRIBUTE, ST_OMP_TEAMS_DISTRIBUTE_SIMD,
+ ST_OMP_END_TEAMS_DISTRIBUTE_SIMD, ST_OMP_TARGET_TEAMS_DISTRIBUTE,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE, ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD, ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
ST_PROCEDURE, ST_GENERIC, ST_CRITICAL, ST_END_CRITICAL,
ST_GET_FCN_CHARACTERISTICS, ST_LOCK, ST_UNLOCK, ST_NONE
}
!$OMP DECLARE REDUCTION. */
unsigned omp_udr_artificial_var:1;
+ /* Mentioned in OMP DECLARE TARGET. */
+ unsigned omp_declare_target:1;
+
/* Attributes set by compiler extensions (!GCC$ ATTRIBUTES). */
unsigned ext_attr:EXT_ATTR_NUM;
}
gfc_omp_reduction_op;
+typedef enum
+{
+ OMP_DEPEND_IN,
+ OMP_DEPEND_OUT,
+ OMP_DEPEND_INOUT
+}
+gfc_omp_depend_op;
+
+typedef enum
+{
+ OMP_MAP_ALLOC,
+ OMP_MAP_TO,
+ OMP_MAP_FROM,
+ OMP_MAP_TOFROM
+}
+gfc_omp_map_op;
+
/* For use in OpenMP clauses in case we need extra information
(aligned clause alignment, linear clause step, etc.). */
{
struct gfc_symbol *sym;
struct gfc_expr *expr;
- gfc_omp_reduction_op rop;
+ union
+ {
+ gfc_omp_reduction_op reduction_op;
+ gfc_omp_depend_op depend_op;
+ gfc_omp_map_op map_op;
+ } u;
struct gfc_omp_udr *udr;
struct gfc_omp_namelist *next;
}
OMP_LIST_UNIFORM,
OMP_LIST_ALIGNED,
OMP_LIST_LINEAR,
- OMP_LIST_DEPEND_IN,
- OMP_LIST_DEPEND_OUT,
+ OMP_LIST_DEPEND,
+ OMP_LIST_MAP,
+ OMP_LIST_TO,
+ OMP_LIST_FROM,
OMP_LIST_REDUCTION,
OMP_LIST_NUM
};
enum gfc_omp_proc_bind_kind proc_bind;
struct gfc_expr *safelen_expr;
struct gfc_expr *simdlen_expr;
+ struct gfc_expr *num_teams;
+ struct gfc_expr *device;
+ struct gfc_expr *thread_limit;
+ enum gfc_omp_sched_kind dist_sched_kind;
+ struct gfc_expr *dist_chunk_size;
}
gfc_omp_clauses;
typedef struct gfc_common_head
{
locus where;
- char use_assoc, saved, threadprivate;
+ char use_assoc, saved, threadprivate, omp_declare_target;
char name[GFC_MAX_SYMBOL_LEN + 1];
struct gfc_symbol *head;
const char* binding_label;
EXEC_OMP_END_SINGLE, EXEC_OMP_TASK, EXEC_OMP_TASKWAIT,
EXEC_OMP_TASKYIELD, EXEC_OMP_CANCEL, EXEC_OMP_CANCELLATION_POINT,
EXEC_OMP_TASKGROUP, EXEC_OMP_SIMD, EXEC_OMP_DO_SIMD,
- EXEC_OMP_PARALLEL_DO_SIMD
+ EXEC_OMP_PARALLEL_DO_SIMD, EXEC_OMP_TARGET, EXEC_OMP_TARGET_DATA,
+ EXEC_OMP_TEAMS, EXEC_OMP_DISTRIBUTE, EXEC_OMP_DISTRIBUTE_SIMD,
+ EXEC_OMP_DISTRIBUTE_PARALLEL_DO, EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
+ EXEC_OMP_TARGET_TEAMS, EXEC_OMP_TEAMS_DISTRIBUTE,
+ EXEC_OMP_TEAMS_DISTRIBUTE_SIMD, EXEC_OMP_TARGET_TEAMS_DISTRIBUTE,
+ EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+ EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ EXEC_OMP_TARGET_UPDATE
}
gfc_exec_op;
bool gfc_add_result (symbol_attribute *, const char *, locus *);
bool gfc_add_save (symbol_attribute *, save_state, const char *, locus *);
bool gfc_add_threadprivate (symbol_attribute *, const char *, locus *);
+bool gfc_add_omp_declare_target (symbol_attribute *, const char *, locus *);
bool gfc_add_saved_common (symbol_attribute *, locus *);
bool gfc_add_target (symbol_attribute *, locus *);
bool gfc_add_dummy (symbol_attribute *, const char *, locus *);
match gfc_match_omp_critical (void);
match gfc_match_omp_declare_reduction (void);
match gfc_match_omp_declare_simd (void);
+match gfc_match_omp_declare_target (void);
+match gfc_match_omp_distribute (void);
+match gfc_match_omp_distribute_parallel_do (void);
+match gfc_match_omp_distribute_parallel_do_simd (void);
+match gfc_match_omp_distribute_simd (void);
match gfc_match_omp_do (void);
match gfc_match_omp_do_simd (void);
match gfc_match_omp_flush (void);
match gfc_match_omp_sections (void);
match gfc_match_omp_simd (void);
match gfc_match_omp_single (void);
+match gfc_match_omp_target (void);
+match gfc_match_omp_target_data (void);
+match gfc_match_omp_target_teams (void);
+match gfc_match_omp_target_teams_distribute (void);
+match gfc_match_omp_target_teams_distribute_parallel_do (void);
+match gfc_match_omp_target_teams_distribute_parallel_do_simd (void);
+match gfc_match_omp_target_teams_distribute_simd (void);
+match gfc_match_omp_target_update (void);
match gfc_match_omp_task (void);
match gfc_match_omp_taskgroup (void);
match gfc_match_omp_taskwait (void);
match gfc_match_omp_taskyield (void);
+match gfc_match_omp_teams (void);
+match gfc_match_omp_teams_distribute (void);
+match gfc_match_omp_teams_distribute_parallel_do (void);
+match gfc_match_omp_teams_distribute_parallel_do_simd (void);
+match gfc_match_omp_teams_distribute_simd (void);
match gfc_match_omp_threadprivate (void);
match gfc_match_omp_workshare (void);
match gfc_match_omp_end_nowait (void);
AB_IS_BIND_C, AB_IS_C_INTEROP, AB_IS_ISO_C, AB_ABSTRACT, AB_ZERO_COMP,
AB_IS_CLASS, AB_PROCEDURE, AB_PROC_POINTER, AB_ASYNCHRONOUS, AB_CODIMENSION,
AB_COARRAY_COMP, AB_VTYPE, AB_VTAB, AB_CONTIGUOUS, AB_CLASS_POINTER,
- AB_IMPLICIT_PURE, AB_ARTIFICIAL, AB_UNLIMITED_POLY
+ AB_IMPLICIT_PURE, AB_ARTIFICIAL, AB_UNLIMITED_POLY, AB_OMP_DECLARE_TARGET
}
ab_attribute;
minit ("CLASS_POINTER", AB_CLASS_POINTER),
minit ("IMPLICIT_PURE", AB_IMPLICIT_PURE),
minit ("UNLIMITED_POLY", AB_UNLIMITED_POLY),
+ minit ("OMP_DECLARE_TARGET", AB_OMP_DECLARE_TARGET),
minit (NULL, -1)
};
MIO_NAME (ab_attribute) (AB_VTYPE, attr_bits);
if (attr->vtab)
MIO_NAME (ab_attribute) (AB_VTAB, attr_bits);
+ if (attr->omp_declare_target)
+ MIO_NAME (ab_attribute) (AB_OMP_DECLARE_TARGET, attr_bits);
mio_rparen ();
case AB_VTAB:
attr->vtab = 1;
break;
+ case AB_OMP_DECLARE_TARGET:
+ attr->omp_declare_target = 1;
+ break;
}
}
}
gfc_free_expr (c->chunk_size);
gfc_free_expr (c->safelen_expr);
gfc_free_expr (c->simdlen_expr);
+ gfc_free_expr (c->num_teams);
+ gfc_free_expr (c->device);
+ gfc_free_expr (c->thread_limit);
+ gfc_free_expr (c->dist_chunk_size);
for (i = 0; i < OMP_LIST_NUM; i++)
gfc_free_omp_namelist (c->lists[i]);
free (c);
return MATCH_ERROR;
}
-#define OMP_CLAUSE_PRIVATE (1 << 0)
-#define OMP_CLAUSE_FIRSTPRIVATE (1 << 1)
-#define OMP_CLAUSE_LASTPRIVATE (1 << 2)
-#define OMP_CLAUSE_COPYPRIVATE (1 << 3)
-#define OMP_CLAUSE_SHARED (1 << 4)
-#define OMP_CLAUSE_COPYIN (1 << 5)
-#define OMP_CLAUSE_REDUCTION (1 << 6)
-#define OMP_CLAUSE_IF (1 << 7)
-#define OMP_CLAUSE_NUM_THREADS (1 << 8)
-#define OMP_CLAUSE_SCHEDULE (1 << 9)
-#define OMP_CLAUSE_DEFAULT (1 << 10)
-#define OMP_CLAUSE_ORDERED (1 << 11)
-#define OMP_CLAUSE_COLLAPSE (1 << 12)
-#define OMP_CLAUSE_UNTIED (1 << 13)
-#define OMP_CLAUSE_FINAL (1 << 14)
-#define OMP_CLAUSE_MERGEABLE (1 << 15)
-#define OMP_CLAUSE_ALIGNED (1 << 16)
-#define OMP_CLAUSE_DEPEND (1 << 17)
-#define OMP_CLAUSE_INBRANCH (1 << 18)
-#define OMP_CLAUSE_LINEAR (1 << 19)
-#define OMP_CLAUSE_NOTINBRANCH (1 << 20)
-#define OMP_CLAUSE_PROC_BIND (1 << 21)
-#define OMP_CLAUSE_SAFELEN (1 << 22)
-#define OMP_CLAUSE_SIMDLEN (1 << 23)
-#define OMP_CLAUSE_UNIFORM (1 << 24)
+#define OMP_CLAUSE_PRIVATE (1U << 0)
+#define OMP_CLAUSE_FIRSTPRIVATE (1U << 1)
+#define OMP_CLAUSE_LASTPRIVATE (1U << 2)
+#define OMP_CLAUSE_COPYPRIVATE (1U << 3)
+#define OMP_CLAUSE_SHARED (1U << 4)
+#define OMP_CLAUSE_COPYIN (1U << 5)
+#define OMP_CLAUSE_REDUCTION (1U << 6)
+#define OMP_CLAUSE_IF (1U << 7)
+#define OMP_CLAUSE_NUM_THREADS (1U << 8)
+#define OMP_CLAUSE_SCHEDULE (1U << 9)
+#define OMP_CLAUSE_DEFAULT (1U << 10)
+#define OMP_CLAUSE_ORDERED (1U << 11)
+#define OMP_CLAUSE_COLLAPSE (1U << 12)
+#define OMP_CLAUSE_UNTIED (1U << 13)
+#define OMP_CLAUSE_FINAL (1U << 14)
+#define OMP_CLAUSE_MERGEABLE (1U << 15)
+#define OMP_CLAUSE_ALIGNED (1U << 16)
+#define OMP_CLAUSE_DEPEND (1U << 17)
+#define OMP_CLAUSE_INBRANCH (1U << 18)
+#define OMP_CLAUSE_LINEAR (1U << 19)
+#define OMP_CLAUSE_NOTINBRANCH (1U << 20)
+#define OMP_CLAUSE_PROC_BIND (1U << 21)
+#define OMP_CLAUSE_SAFELEN (1U << 22)
+#define OMP_CLAUSE_SIMDLEN (1U << 23)
+#define OMP_CLAUSE_UNIFORM (1U << 24)
+#define OMP_CLAUSE_DEVICE (1U << 25)
+#define OMP_CLAUSE_MAP (1U << 26)
+#define OMP_CLAUSE_TO (1U << 27)
+#define OMP_CLAUSE_FROM (1U << 28)
+#define OMP_CLAUSE_NUM_TEAMS (1U << 29)
+#define OMP_CLAUSE_THREAD_LIMIT (1U << 30)
+#define OMP_CLAUSE_DIST_SCHEDULE (1U << 31)
/* Match OpenMP directive clauses. MASK is a bitmask of
clauses that are allowed for a particular directive. */
static match
-gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true,
- bool needs_space = true)
+gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned int mask,
+ bool first = true, bool needs_space = true)
{
gfc_omp_clauses *c = gfc_get_omp_clauses ();
locus old_loc;
else
for (n = *head; n; n = n->next)
{
- n->rop = rop;
+ n->u.reduction_op = rop;
n->udr = udr;
}
continue;
continue;
}
}
- if ((mask & OMP_CLAUSE_INBRANCH) && !c->inbranch
+ if ((mask & OMP_CLAUSE_INBRANCH) && !c->inbranch && !c->notinbranch
&& gfc_match ("inbranch") == MATCH_YES)
{
c->inbranch = needs_space = true;
continue;
}
- if ((mask & OMP_CLAUSE_NOTINBRANCH) && !c->notinbranch
+ if ((mask & OMP_CLAUSE_NOTINBRANCH) && !c->notinbranch && !c->inbranch
&& gfc_match ("notinbranch") == MATCH_YES)
{
c->notinbranch = needs_space = true;
continue;
}
if ((mask & OMP_CLAUSE_DEPEND)
- && gfc_match_omp_variable_list ("depend ( in : ",
- &c->lists[OMP_LIST_DEPEND_IN], false,
- NULL, NULL, true)
- == MATCH_YES)
+ && gfc_match ("depend ( ") == MATCH_YES)
+ {
+ match m = MATCH_YES;
+ gfc_omp_depend_op depend_op = OMP_DEPEND_OUT;
+ if (gfc_match ("inout") == MATCH_YES)
+ depend_op = OMP_DEPEND_INOUT;
+ else if (gfc_match ("in") == MATCH_YES)
+ depend_op = OMP_DEPEND_IN;
+ else if (gfc_match ("out") == MATCH_YES)
+ depend_op = OMP_DEPEND_OUT;
+ else
+ m = MATCH_NO;
+ head = NULL;
+ if (m == MATCH_YES
+ && gfc_match_omp_variable_list (" : ",
+ &c->lists[OMP_LIST_DEPEND],
+ false, NULL, &head, true)
+ == MATCH_YES)
+ {
+ gfc_omp_namelist *n;
+ for (n = *head; n; n = n->next)
+ n->u.depend_op = depend_op;
+ continue;
+ }
+ else
+ gfc_current_locus = old_loc;
+ }
+ if ((mask & OMP_CLAUSE_DIST_SCHEDULE)
+ && c->dist_sched_kind == OMP_SCHED_NONE
+ && gfc_match ("dist_schedule ( static") == MATCH_YES)
+ {
+ match m = MATCH_NO;
+ c->dist_sched_kind = OMP_SCHED_STATIC;
+ m = gfc_match (" , %e )", &c->dist_chunk_size);
+ if (m != MATCH_YES)
+ m = gfc_match_char (')');
+ if (m != MATCH_YES)
+ {
+ c->dist_sched_kind = OMP_SCHED_NONE;
+ gfc_current_locus = old_loc;
+ }
+ else
+ continue;
+ }
+ if ((mask & OMP_CLAUSE_NUM_TEAMS) && c->num_teams == NULL
+ && gfc_match ("num_teams ( %e )", &c->num_teams) == MATCH_YES)
continue;
- if ((mask & OMP_CLAUSE_DEPEND)
- && gfc_match_omp_variable_list ("depend ( out : ",
- &c->lists[OMP_LIST_DEPEND_OUT], false,
- NULL, NULL, true)
+ if ((mask & OMP_CLAUSE_DEVICE) && c->device == NULL
+ && gfc_match ("device ( %e )", &c->device) == MATCH_YES)
+ continue;
+ if ((mask & OMP_CLAUSE_THREAD_LIMIT) && c->thread_limit == NULL
+ && gfc_match ("thread_limit ( %e )", &c->thread_limit) == MATCH_YES)
+ continue;
+ if ((mask & OMP_CLAUSE_MAP)
+ && gfc_match ("map ( ") == MATCH_YES)
+ {
+ gfc_omp_map_op map_op = OMP_MAP_TOFROM;
+ if (gfc_match ("alloc : ") == MATCH_YES)
+ map_op = OMP_MAP_ALLOC;
+ else if (gfc_match ("tofrom : ") == MATCH_YES)
+ map_op = OMP_MAP_TOFROM;
+ else if (gfc_match ("to : ") == MATCH_YES)
+ map_op = OMP_MAP_TO;
+ else if (gfc_match ("from : ") == MATCH_YES)
+ map_op = OMP_MAP_FROM;
+ head = NULL;
+ if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP],
+ false, NULL, &head, true)
+ == MATCH_YES)
+ {
+ gfc_omp_namelist *n;
+ for (n = *head; n; n = n->next)
+ n->u.map_op = map_op;
+ continue;
+ }
+ else
+ gfc_current_locus = old_loc;
+ }
+ if ((mask & OMP_CLAUSE_TO)
+ && gfc_match_omp_variable_list ("to (",
+ &c->lists[OMP_LIST_TO], false,
+ NULL, &head, true)
== MATCH_YES)
continue;
- if ((mask & OMP_CLAUSE_DEPEND)
- && gfc_match_omp_variable_list ("depend ( inout : ",
- &c->lists[OMP_LIST_DEPEND_OUT], false,
- NULL, NULL, true)
+ if ((mask & OMP_CLAUSE_FROM)
+ && gfc_match_omp_variable_list ("from (",
+ &c->lists[OMP_LIST_FROM], false,
+ NULL, &head, true)
== MATCH_YES)
continue;
| OMP_CLAUSE_NUM_THREADS | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PROC_BIND)
#define OMP_DECLARE_SIMD_CLAUSES \
(OMP_CLAUSE_SIMDLEN | OMP_CLAUSE_LINEAR | OMP_CLAUSE_UNIFORM \
- | OMP_CLAUSE_ALIGNED)
+ | OMP_CLAUSE_ALIGNED | OMP_CLAUSE_INBRANCH | OMP_CLAUSE_NOTINBRANCH)
#define OMP_DO_CLAUSES \
(OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_REDUCTION \
(OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_SHARED \
| OMP_CLAUSE_IF | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED \
| OMP_CLAUSE_FINAL | OMP_CLAUSE_MERGEABLE | OMP_CLAUSE_DEPEND)
-
-match
-gfc_match_omp_parallel (void)
-{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_PARALLEL_CLAUSES) != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_PARALLEL;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
-}
+#define OMP_TARGET_CLAUSES \
+ (OMP_CLAUSE_DEVICE | OMP_CLAUSE_MAP | OMP_CLAUSE_IF)
+#define OMP_TARGET_DATA_CLAUSES \
+ (OMP_CLAUSE_DEVICE | OMP_CLAUSE_MAP | OMP_CLAUSE_IF)
+#define OMP_TARGET_UPDATE_CLAUSES \
+ (OMP_CLAUSE_DEVICE | OMP_CLAUSE_IF | OMP_CLAUSE_TO | OMP_CLAUSE_FROM)
+#define OMP_TEAMS_CLAUSES \
+ (OMP_CLAUSE_NUM_TEAMS | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_DEFAULT \
+ | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_SHARED \
+ | OMP_CLAUSE_REDUCTION)
+#define OMP_DISTRIBUTE_CLAUSES \
+ (OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_COLLAPSE \
+ | OMP_CLAUSE_DIST_SCHEDULE)
-match
-gfc_match_omp_task (void)
+static match
+match_omp (gfc_exec_op op, unsigned int mask)
{
gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_TASK_CLAUSES) != MATCH_YES)
+ if (gfc_match_omp_clauses (&c, mask) != MATCH_YES)
return MATCH_ERROR;
- new_st.op = EXEC_OMP_TASK;
+ new_st.op = op;
new_st.ext.omp_clauses = c;
return MATCH_YES;
}
match
-gfc_match_omp_taskwait (void)
+gfc_match_omp_critical (void)
{
+ char n[GFC_MAX_SYMBOL_LEN+1];
+
+ if (gfc_match (" ( %n )", n) != MATCH_YES)
+ n[0] = '\0';
if (gfc_match_omp_eos () != MATCH_YES)
{
- gfc_error ("Unexpected junk after TASKWAIT clause at %C");
+ gfc_error ("Unexpected junk after $OMP CRITICAL statement at %C");
return MATCH_ERROR;
}
- new_st.op = EXEC_OMP_TASKWAIT;
- new_st.ext.omp_clauses = NULL;
+ new_st.op = EXEC_OMP_CRITICAL;
+ new_st.ext.omp_name = n[0] ? xstrdup (n) : NULL;
return MATCH_YES;
}
match
-gfc_match_omp_taskyield (void)
+gfc_match_omp_distribute (void)
{
- if (gfc_match_omp_eos () != MATCH_YES)
- {
- gfc_error ("Unexpected junk after TASKYIELD clause at %C");
- return MATCH_ERROR;
- }
- new_st.op = EXEC_OMP_TASKYIELD;
- new_st.ext.omp_clauses = NULL;
- return MATCH_YES;
+ return match_omp (EXEC_OMP_DISTRIBUTE, OMP_DISTRIBUTE_CLAUSES);
}
match
-gfc_match_omp_critical (void)
+gfc_match_omp_distribute_parallel_do (void)
{
- char n[GFC_MAX_SYMBOL_LEN+1];
+ return match_omp (EXEC_OMP_DISTRIBUTE_PARALLEL_DO,
+ OMP_DISTRIBUTE_CLAUSES | OMP_PARALLEL_CLAUSES
+ | OMP_DO_CLAUSES);
+}
- if (gfc_match (" ( %n )", n) != MATCH_YES)
- n[0] = '\0';
- if (gfc_match_omp_eos () != MATCH_YES)
- {
- gfc_error ("Unexpected junk after $OMP CRITICAL statement at %C");
- return MATCH_ERROR;
- }
- new_st.op = EXEC_OMP_CRITICAL;
- new_st.ext.omp_name = n[0] ? xstrdup (n) : NULL;
- return MATCH_YES;
+
+match
+gfc_match_omp_distribute_parallel_do_simd (void)
+{
+ return match_omp (EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
+ (OMP_DISTRIBUTE_CLAUSES | OMP_PARALLEL_CLAUSES
+ | OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
+ & ~OMP_CLAUSE_ORDERED);
+}
+
+
+match
+gfc_match_omp_distribute_simd (void)
+{
+ return match_omp (EXEC_OMP_DISTRIBUTE_SIMD,
+ OMP_DISTRIBUTE_CLAUSES | OMP_SIMD_CLAUSES);
}
match
gfc_match_omp_do (void)
{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_DO_CLAUSES) != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_DO;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
+ return match_omp (EXEC_OMP_DO, OMP_DO_CLAUSES);
}
match
gfc_match_omp_do_simd (void)
{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, ((OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
- & ~OMP_CLAUSE_ORDERED))
- != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_DO_SIMD;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
+ return match_omp (EXEC_OMP_DO_SIMD, ((OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
+ & ~OMP_CLAUSE_ORDERED));
}
match
-gfc_match_omp_simd (void)
-{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_SIMD_CLAUSES) != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_SIMD;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
-}
-
-
-match
gfc_match_omp_declare_simd (void)
{
locus where = gfc_current_locus;
if (end_loc_set)
{
gfc_current_locus = end_loc;
+ if (gfc_match_omp_eos () != MATCH_YES)
+ {
+ gfc_error ("Unexpected junk after !$OMP DECLARE REDUCTION at %C");
+ gfc_current_locus = where;
+ return MATCH_ERROR;
+ }
+
return MATCH_YES;
}
gfc_clear_error ();
match
+gfc_match_omp_declare_target (void)
+{
+ locus old_loc;
+ char n[GFC_MAX_SYMBOL_LEN+1];
+ gfc_symbol *sym;
+ match m;
+ gfc_symtree *st;
+
+ old_loc = gfc_current_locus;
+
+ m = gfc_match (" (");
+
+ if (gfc_current_ns->proc_name
+ && gfc_current_ns->proc_name->attr.if_source == IFSRC_IFBODY
+ && m == MATCH_YES)
+ {
+ gfc_error ("Only the !$OMP DECLARE TARGET form without "
+ "list is allowed in interface block at %C");
+ goto cleanup;
+ }
+
+ if (m == MATCH_NO
+ && gfc_current_ns->proc_name
+ && gfc_match_omp_eos () == MATCH_YES)
+ {
+ if (!gfc_add_omp_declare_target (&gfc_current_ns->proc_name->attr,
+ gfc_current_ns->proc_name->name,
+ &old_loc))
+ goto cleanup;
+ return MATCH_YES;
+ }
+
+ if (m != MATCH_YES)
+ return m;
+
+ for (;;)
+ {
+ m = gfc_match_symbol (&sym, 0);
+ switch (m)
+ {
+ case MATCH_YES:
+ if (sym->attr.in_common)
+ gfc_error_now ("OMP DECLARE TARGET on a variable at %C is an "
+ "element of a COMMON block");
+ else if (!gfc_add_omp_declare_target (&sym->attr, sym->name,
+ &sym->declared_at))
+ goto cleanup;
+ goto next_item;
+ case MATCH_NO:
+ break;
+ case MATCH_ERROR:
+ goto cleanup;
+ }
+
+ m = gfc_match (" / %n /", n);
+ if (m == MATCH_ERROR)
+ goto cleanup;
+ if (m == MATCH_NO || n[0] == '\0')
+ goto syntax;
+
+ st = gfc_find_symtree (gfc_current_ns->common_root, n);
+ if (st == NULL)
+ {
+ gfc_error ("COMMON block /%s/ not found at %C", n);
+ goto cleanup;
+ }
+ st->n.common->omp_declare_target = 1;
+ for (sym = st->n.common->head; sym; sym = sym->common_next)
+ if (!gfc_add_omp_declare_target (&sym->attr, sym->name,
+ &sym->declared_at))
+ goto cleanup;
+
+ next_item:
+ if (gfc_match_char (')') == MATCH_YES)
+ break;
+ if (gfc_match_char (',') != MATCH_YES)
+ goto syntax;
+ }
+
+ if (gfc_match_omp_eos () != MATCH_YES)
+ {
+ gfc_error ("Unexpected junk after !$OMP DECLARE TARGET at %C");
+ goto cleanup;
+ }
+ return MATCH_YES;
+
+syntax:
+ gfc_error ("Syntax error in !$OMP DECLARE TARGET list at %C");
+
+cleanup:
+ gfc_current_locus = old_loc;
+ return MATCH_ERROR;
+}
+
+
+match
gfc_match_omp_threadprivate (void)
{
locus old_loc;
goto syntax;
}
+ if (gfc_match_omp_eos () != MATCH_YES)
+ {
+ gfc_error ("Unexpected junk after OMP THREADPRIVATE at %C");
+ goto cleanup;
+ }
+
return MATCH_YES;
syntax:
match
+gfc_match_omp_parallel (void)
+{
+ return match_omp (EXEC_OMP_PARALLEL, OMP_PARALLEL_CLAUSES);
+}
+
+
+match
gfc_match_omp_parallel_do (void)
{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES)
- != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_PARALLEL_DO;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
+ return match_omp (EXEC_OMP_PARALLEL_DO,
+ OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES);
}
match
gfc_match_omp_parallel_do_simd (void)
{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, (OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES
- | OMP_SIMD_CLAUSES) & ~OMP_CLAUSE_ORDERED)
- != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_PARALLEL_DO_SIMD;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
+ return match_omp (EXEC_OMP_PARALLEL_DO_SIMD,
+ (OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
+ & ~OMP_CLAUSE_ORDERED);
}
match
gfc_match_omp_parallel_sections (void)
{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_PARALLEL_CLAUSES | OMP_SECTIONS_CLAUSES)
- != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_PARALLEL_SECTIONS;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
+ return match_omp (EXEC_OMP_PARALLEL_SECTIONS,
+ OMP_PARALLEL_CLAUSES | OMP_SECTIONS_CLAUSES);
}
match
gfc_match_omp_parallel_workshare (void)
{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_PARALLEL_CLAUSES) != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_PARALLEL_WORKSHARE;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
+ return match_omp (EXEC_OMP_PARALLEL_WORKSHARE, OMP_PARALLEL_CLAUSES);
}
match
gfc_match_omp_sections (void)
{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_SECTIONS_CLAUSES) != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_SECTIONS;
- new_st.ext.omp_clauses = c;
- return MATCH_YES;
+ return match_omp (EXEC_OMP_SECTIONS, OMP_SECTIONS_CLAUSES);
+}
+
+
+match
+gfc_match_omp_simd (void)
+{
+ return match_omp (EXEC_OMP_SIMD, OMP_SIMD_CLAUSES);
}
match
gfc_match_omp_single (void)
{
- gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE)
- != MATCH_YES)
- return MATCH_ERROR;
- new_st.op = EXEC_OMP_SINGLE;
- new_st.ext.omp_clauses = c;
+ return match_omp (EXEC_OMP_SINGLE,
+ OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE);
+}
+
+
+match
+gfc_match_omp_task (void)
+{
+ return match_omp (EXEC_OMP_TASK, OMP_TASK_CLAUSES);
+}
+
+
+match
+gfc_match_omp_taskwait (void)
+{
+ if (gfc_match_omp_eos () != MATCH_YES)
+ {
+ gfc_error ("Unexpected junk after TASKWAIT clause at %C");
+ return MATCH_ERROR;
+ }
+ new_st.op = EXEC_OMP_TASKWAIT;
+ new_st.ext.omp_clauses = NULL;
return MATCH_YES;
}
match
+gfc_match_omp_taskyield (void)
+{
+ if (gfc_match_omp_eos () != MATCH_YES)
+ {
+ gfc_error ("Unexpected junk after TASKYIELD clause at %C");
+ return MATCH_ERROR;
+ }
+ new_st.op = EXEC_OMP_TASKYIELD;
+ new_st.ext.omp_clauses = NULL;
+ return MATCH_YES;
+}
+
+
+match
+gfc_match_omp_target (void)
+{
+ return match_omp (EXEC_OMP_TARGET, OMP_TARGET_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_data (void)
+{
+ return match_omp (EXEC_OMP_TARGET_DATA, OMP_TARGET_DATA_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_teams (void)
+{
+ return match_omp (EXEC_OMP_TARGET_TEAMS,
+ OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_teams_distribute (void)
+{
+ return match_omp (EXEC_OMP_TARGET_TEAMS_DISTRIBUTE,
+ OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES
+ | OMP_DISTRIBUTE_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_teams_distribute_parallel_do (void)
+{
+ return match_omp (EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES
+ | OMP_DISTRIBUTE_CLAUSES | OMP_PARALLEL_CLAUSES
+ | OMP_DO_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_teams_distribute_parallel_do_simd (void)
+{
+ return match_omp (EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ (OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES
+ | OMP_DISTRIBUTE_CLAUSES | OMP_PARALLEL_CLAUSES
+ | OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
+ & ~OMP_CLAUSE_ORDERED);
+}
+
+
+match
+gfc_match_omp_target_teams_distribute_simd (void)
+{
+ return match_omp (EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+ OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES
+ | OMP_DISTRIBUTE_CLAUSES | OMP_SIMD_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_update (void)
+{
+ return match_omp (EXEC_OMP_TARGET_UPDATE, OMP_TARGET_UPDATE_CLAUSES);
+}
+
+
+match
+gfc_match_omp_teams (void)
+{
+ return match_omp (EXEC_OMP_TEAMS, OMP_TEAMS_CLAUSES);
+}
+
+
+match
+gfc_match_omp_teams_distribute (void)
+{
+ return match_omp (EXEC_OMP_TEAMS_DISTRIBUTE,
+ OMP_TEAMS_CLAUSES | OMP_DISTRIBUTE_CLAUSES);
+}
+
+
+match
+gfc_match_omp_teams_distribute_parallel_do (void)
+{
+ return match_omp (EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+ OMP_TEAMS_CLAUSES | OMP_DISTRIBUTE_CLAUSES
+ | OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES);
+}
+
+
+match
+gfc_match_omp_teams_distribute_parallel_do_simd (void)
+{
+ return match_omp (EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+ (OMP_TEAMS_CLAUSES | OMP_DISTRIBUTE_CLAUSES
+ | OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES
+ | OMP_SIMD_CLAUSES) & ~OMP_CLAUSE_ORDERED);
+}
+
+
+match
+gfc_match_omp_teams_distribute_simd (void)
+{
+ return match_omp (EXEC_OMP_TEAMS_DISTRIBUTE_SIMD,
+ OMP_TEAMS_CLAUSES | OMP_DISTRIBUTE_CLAUSES
+ | OMP_SIMD_CLAUSES);
+}
+
+
+match
gfc_match_omp_workshare (void)
{
if (gfc_match_omp_eos () != MATCH_YES)
int list;
static const char *clause_names[]
= { "PRIVATE", "FIRSTPRIVATE", "LASTPRIVATE", "COPYPRIVATE", "SHARED",
- "COPYIN", "UNIFORM", "ALIGNED", "LINEAR", "DEPEND", "DEPEND",
- "REDUCTION" };
+ "COPYIN", "UNIFORM", "ALIGNED", "LINEAR", "DEPEND", "MAP",
+ "TO", "FROM", "REDUCTION" };
if (omp_clauses == NULL)
return;
if (list != OMP_LIST_FIRSTPRIVATE
&& list != OMP_LIST_LASTPRIVATE
&& list != OMP_LIST_ALIGNED
- && list != OMP_LIST_DEPEND_IN
- && list != OMP_LIST_DEPEND_OUT)
+ && list != OMP_LIST_DEPEND
+ && list != OMP_LIST_MAP
+ && list != OMP_LIST_FROM
+ && list != OMP_LIST_TO)
for (n = omp_clauses->lists[list]; n; n = n->next)
{
if (n->sym->mark)
n->sym->mark = 1;
}
+ for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
+ n->sym->mark = 0;
+ for (n = omp_clauses->lists[OMP_LIST_FROM]; n; n = n->next)
+ if (n->expr == NULL)
+ n->sym->mark = 1;
+ for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
+ {
+ if (n->expr == NULL && n->sym->mark)
+ gfc_error ("Symbol '%s' present on both FROM and TO clauses at %L",
+ n->sym->name, where);
+ else
+ n->sym->mark = 1;
+ }
+
for (list = 0; list < OMP_LIST_NUM; list++)
if ((n = omp_clauses->lists[list]) != NULL)
{
}
}
break;
- case OMP_LIST_DEPEND_IN:
- case OMP_LIST_DEPEND_OUT:
+ case OMP_LIST_DEPEND:
+ case OMP_LIST_MAP:
+ case OMP_LIST_TO:
+ case OMP_LIST_FROM:
for (; n != NULL; n = n->next)
if (n->expr)
{
|| n->expr->ref == NULL
|| n->expr->ref->next
|| n->expr->ref->type != REF_ARRAY)
- gfc_error ("'%s' in DEPEND clause at %L is not a proper "
- "array section", n->sym->name, where);
+ gfc_error ("'%s' in %s clause at %L is not a proper "
+ "array section", n->sym->name, name, where);
else if (n->expr->ref->u.ar.codimen)
- gfc_error ("Coarrays not supported in DEPEND clause at %L",
- where);
+ gfc_error ("Coarrays not supported in %s clause at %L",
+ name, where);
else
{
int i;
if (ar->stride[i])
{
gfc_error ("Stride should not be specified for "
- "array section in DEPEND clause at %L",
- where);
+ "array section in %s clause at %L",
+ name, where);
break;
}
else if (ar->dimen_type[i] != DIMEN_ELEMENT
&& ar->dimen_type[i] != DIMEN_RANGE)
{
- gfc_error ("'%s' in DEPEND clause at %L is not a "
+ gfc_error ("'%s' in %s clause at %L is not a "
"proper array section",
- n->sym->name, where);
+ n->sym->name, name, where);
break;
}
- else if (ar->start[i]
+ else if (list == OMP_LIST_DEPEND
+ && ar->start[i]
&& ar->start[i]->expr_type == EXPR_CONSTANT
&& ar->end[i]
&& ar->end[i]->expr_type == EXPR_CONSTANT
}
}
}
+ if (list != OMP_LIST_DEPEND)
+ for (n = omp_clauses->lists[list]; n != NULL; n = n->next)
+ {
+ n->sym->attr.referenced = 1;
+ if (n->sym->attr.threadprivate)
+ gfc_error ("THREADPRIVATE object '%s' in %s clause at %L",
+ n->sym->name, name, where);
+ if (n->sym->attr.cray_pointee)
+ gfc_error ("Cray pointee '%s' in %s clause at %L",
+ n->sym->name, name, where);
+ }
break;
default:
for (; n != NULL; n = n->next)
switch (list)
{
case OMP_LIST_REDUCTION:
- switch (n->rop)
+ switch (n->u.reduction_op)
{
case OMP_REDUCTION_PLUS:
case OMP_REDUCTION_TIMES:
if (n->udr == NULL)
{
if (udr_name == NULL)
- switch (n->rop)
+ switch (n->u.reduction_op)
{
case OMP_REDUCTION_PLUS:
case OMP_REDUCTION_TIMES:
case OMP_REDUCTION_EQV:
case OMP_REDUCTION_NEQV:
udr_name = gfc_op2string ((gfc_intrinsic_op)
- n->rop);
+ n->u.reduction_op);
break;
case OMP_REDUCTION_MAX:
udr_name = "max";
gfc_typename (&n->sym->ts), where);
}
else
- n->rop = OMP_REDUCTION_USER;
+ n->u.reduction_op = OMP_REDUCTION_USER;
}
break;
case OMP_LIST_LINEAR:
gfc_error ("SIMDLEN clause at %L requires a scalar "
"INTEGER expression", &expr->where);
}
+ if (omp_clauses->num_teams)
+ {
+ gfc_expr *expr = omp_clauses->num_teams;
+ if (!gfc_resolve_expr (expr)
+ || expr->ts.type != BT_INTEGER || expr->rank != 0)
+ gfc_error ("NUM_TEAMS clause at %L requires a scalar "
+ "INTEGER expression", &expr->where);
+ }
+ if (omp_clauses->device)
+ {
+ gfc_expr *expr = omp_clauses->device;
+ if (!gfc_resolve_expr (expr)
+ || expr->ts.type != BT_INTEGER || expr->rank != 0)
+ gfc_error ("DEVICE clause at %L requires a scalar "
+ "INTEGER expression", &expr->where);
+ }
+ if (omp_clauses->dist_chunk_size)
+ {
+ gfc_expr *expr = omp_clauses->dist_chunk_size;
+ if (!gfc_resolve_expr (expr)
+ || expr->ts.type != BT_INTEGER || expr->rank != 0)
+ gfc_error ("DIST_SCHEDULE clause's chunk_size at %L requires "
+ "a scalar INTEGER expression", &expr->where);
+ }
+ if (omp_clauses->thread_limit)
+ {
+ gfc_expr *expr = omp_clauses->thread_limit;
+ if (!gfc_resolve_expr (expr)
+ || expr->ts.type != BT_INTEGER || expr->rank != 0)
+ gfc_error ("THREAD_LIMIT clause at %L requires a scalar "
+ "INTEGER expression", &expr->where);
+ }
}
omp_current_ctx = &ctx;
for (list = 0; list < OMP_LIST_NUM; list++)
- for (n = omp_clauses->lists[list]; n; n = n->next)
- pointer_set_insert (ctx.sharing_clauses, n->sym);
+ switch (list)
+ {
+ case OMP_LIST_SHARED:
+ case OMP_LIST_PRIVATE:
+ case OMP_LIST_FIRSTPRIVATE:
+ case OMP_LIST_LASTPRIVATE:
+ case OMP_LIST_REDUCTION:
+ case OMP_LIST_LINEAR:
+ for (n = omp_clauses->lists[list]; n; n = n->next)
+ pointer_set_insert (ctx.sharing_clauses, n->sym);
+ break;
+ default:
+ break;
+ }
- if (code->op == EXEC_OMP_PARALLEL_DO
- || code->op == EXEC_OMP_PARALLEL_DO_SIMD)
- gfc_resolve_omp_do_blocks (code, ns);
- else
- gfc_resolve_blocks (code->block, ns);
+ switch (code->op)
+ {
+ case EXEC_OMP_PARALLEL_DO:
+ case EXEC_OMP_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+ gfc_resolve_omp_do_blocks (code, ns);
+ break;
+ default:
+ gfc_resolve_blocks (code->block, ns);
+ }
omp_current_ctx = ctx.previous;
pointer_set_destroy (ctx.sharing_clauses);
switch (code->op)
{
+ case EXEC_OMP_DISTRIBUTE: name = "!$OMP DISTRIBUTE"; break;
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ name = "!$OMP DISTRIBUTE PARALLEL DO";
+ break;
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ name = "!$OMP DISTRIBUTE PARALLEL DO SIMD";
+ is_simd = true;
+ break;
+ case EXEC_OMP_DISTRIBUTE_SIMD:
+ name = "!$OMP DISTRIBUTE SIMD";
+ is_simd = true;
+ break;
case EXEC_OMP_DO: name = "!$OMP DO"; break;
case EXEC_OMP_DO_SIMD: name = "!$OMP DO SIMD"; is_simd = true; break;
case EXEC_OMP_PARALLEL_DO: name = "!$OMP PARALLEL DO"; break;
case EXEC_OMP_PARALLEL_DO_SIMD:
name = "!$OMP PARALLEL DO SIMD";
- is_simd = true; break;
+ is_simd = true;
+ break;
case EXEC_OMP_SIMD: name = "!$OMP SIMD"; is_simd = true; break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ name = "!$OMP TARGET TEAMS_DISTRIBUTE";
+ break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ name = "!$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO";
+ break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ name = "!$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD";
+ is_simd = true;
+ break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ name = "!$OMP TARGET TEAMS DISTRIBUTE SIMD";
+ is_simd = true;
+ break;
+ case EXEC_OMP_TEAMS_DISTRIBUTE: name = "!$OMP TEAMS_DISTRIBUTE"; break;
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ name = "!$OMP TEAMS DISTRIBUTE PARALLEL DO";
+ break;
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ name = "!$OMP TEAMS DISTRIBUTE PARALLEL DO SIMD";
+ is_simd = true;
+ break;
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+ name = "!$OMP TEAMS DISTRIBUTE SIMD";
+ is_simd = true;
+ break;
default: gcc_unreachable ();
}
switch (code->op)
{
+ case EXEC_OMP_DISTRIBUTE:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_DISTRIBUTE_SIMD:
case EXEC_OMP_DO:
case EXEC_OMP_DO_SIMD:
case EXEC_OMP_PARALLEL_DO:
case EXEC_OMP_PARALLEL_DO_SIMD:
case EXEC_OMP_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
resolve_omp_do (code);
break;
case EXEC_OMP_CANCEL:
case EXEC_OMP_PARALLEL_SECTIONS:
case EXEC_OMP_SECTIONS:
case EXEC_OMP_SINGLE:
+ case EXEC_OMP_TARGET:
+ case EXEC_OMP_TARGET_DATA:
+ case EXEC_OMP_TARGET_TEAMS:
case EXEC_OMP_TASK:
+ case EXEC_OMP_TEAMS:
case EXEC_OMP_WORKSHARE:
if (code->ext.omp_clauses)
resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL);
break;
+ case EXEC_OMP_TARGET_UPDATE:
+ if (code->ext.omp_clauses)
+ resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL);
+ if (code->ext.omp_clauses == NULL
+ || (code->ext.omp_clauses->lists[OMP_LIST_TO] == NULL
+ && code->ext.omp_clauses->lists[OMP_LIST_FROM] == NULL))
+ gfc_error ("OMP TARGET UPDATE at %L requires at least one TO or "
+ "FROM clause", &code->loc);
+ break;
case EXEC_OMP_ATOMIC:
resolve_omp_atomic (code);
break;
for (ods = ns->omp_declare_simd; ods; ods = ods->next)
{
if (ods->proc_name != ns->proc_name)
- gfc_error ("!$OMP DECLARE SIMD should refer to containing procedure"
+ gfc_error ("!$OMP DECLARE SIMD should refer to containing procedure "
"'%s' at %L", ns->proc_name->name, &ods->where);
if (ods->clauses)
resolve_omp_clauses (NULL, &ods->where, ods->clauses, ns);
ST_OMP_DECLARE_REDUCTION);
matchs ("declare simd", gfc_match_omp_declare_simd,
ST_OMP_DECLARE_SIMD);
+ matcho ("declare target", gfc_match_omp_declare_target,
+ ST_OMP_DECLARE_TARGET);
+ matchs ("distribute parallel do simd",
+ gfc_match_omp_distribute_parallel_do_simd,
+ ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD);
+ matcho ("distribute parallel do", gfc_match_omp_distribute_parallel_do,
+ ST_OMP_DISTRIBUTE_PARALLEL_DO);
+ matchs ("distribute simd", gfc_match_omp_distribute_simd,
+ ST_OMP_DISTRIBUTE_SIMD);
+ matcho ("distribute", gfc_match_omp_distribute, ST_OMP_DISTRIBUTE);
matchs ("do simd", gfc_match_omp_do_simd, ST_OMP_DO_SIMD);
matcho ("do", gfc_match_omp_do, ST_OMP_DO);
break;
case 'e':
matcho ("end atomic", gfc_match_omp_eos, ST_OMP_END_ATOMIC);
matcho ("end critical", gfc_match_omp_critical, ST_OMP_END_CRITICAL);
+ matchs ("end distribute parallel do simd", gfc_match_omp_eos,
+ ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD);
+ matcho ("end distribute parallel do", gfc_match_omp_eos,
+ ST_OMP_END_DISTRIBUTE_PARALLEL_DO);
+ matchs ("end distribute simd", gfc_match_omp_eos,
+ ST_OMP_END_DISTRIBUTE_SIMD);
+ matcho ("end distribute", gfc_match_omp_eos, ST_OMP_END_DISTRIBUTE);
matchs ("end do simd", gfc_match_omp_end_nowait, ST_OMP_END_DO_SIMD);
matcho ("end do", gfc_match_omp_end_nowait, ST_OMP_END_DO);
matchs ("end simd", gfc_match_omp_eos, ST_OMP_END_SIMD);
matcho ("end parallel", gfc_match_omp_eos, ST_OMP_END_PARALLEL);
matcho ("end sections", gfc_match_omp_end_nowait, ST_OMP_END_SECTIONS);
matcho ("end single", gfc_match_omp_end_single, ST_OMP_END_SINGLE);
+ matcho ("end target data", gfc_match_omp_eos, ST_OMP_END_TARGET_DATA);
+ matchs ("end target teams distribute parallel do simd",
+ gfc_match_omp_eos,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
+ matcho ("end target teams distribute parallel do", gfc_match_omp_eos,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO);
+ matchs ("end target teams distribute simd", gfc_match_omp_eos,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD);
+ matcho ("end target teams distribute", gfc_match_omp_eos,
+ ST_OMP_END_TARGET_TEAMS_DISTRIBUTE);
+ matcho ("end target teams", gfc_match_omp_eos, ST_OMP_END_TARGET_TEAMS);
+ matcho ("end target", gfc_match_omp_eos, ST_OMP_END_TARGET);
matcho ("end taskgroup", gfc_match_omp_eos, ST_OMP_END_TASKGROUP);
matcho ("end task", gfc_match_omp_eos, ST_OMP_END_TASK);
+ matchs ("end teams distribute parallel do simd", gfc_match_omp_eos,
+ ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
+ matcho ("end teams distribute parallel do", gfc_match_omp_eos,
+ ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO);
+ matchs ("end teams distribute simd", gfc_match_omp_eos,
+ ST_OMP_END_TEAMS_DISTRIBUTE_SIMD);
+ matcho ("end teams distribute", gfc_match_omp_eos,
+ ST_OMP_END_TEAMS_DISTRIBUTE);
+ matcho ("end teams", gfc_match_omp_eos, ST_OMP_END_TEAMS);
matcho ("end workshare", gfc_match_omp_end_nowait,
ST_OMP_END_WORKSHARE);
break;
matcho ("single", gfc_match_omp_single, ST_OMP_SINGLE);
break;
case 't':
+ matcho ("target data", gfc_match_omp_target_data, ST_OMP_TARGET_DATA);
+ matchs ("target teams distribute parallel do simd",
+ gfc_match_omp_target_teams_distribute_parallel_do_simd,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
+ matcho ("target teams distribute parallel do",
+ gfc_match_omp_target_teams_distribute_parallel_do,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO);
+ matchs ("target teams distribute simd",
+ gfc_match_omp_target_teams_distribute_simd,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD);
+ matcho ("target teams distribute", gfc_match_omp_target_teams_distribute,
+ ST_OMP_TARGET_TEAMS_DISTRIBUTE);
+ matcho ("target teams", gfc_match_omp_target_teams, ST_OMP_TARGET_TEAMS);
+ matcho ("target update", gfc_match_omp_target_update,
+ ST_OMP_TARGET_UPDATE);
+ matcho ("target", gfc_match_omp_target, ST_OMP_TARGET);
matcho ("taskgroup", gfc_match_omp_taskgroup, ST_OMP_TASKGROUP);
matcho ("taskwait", gfc_match_omp_taskwait, ST_OMP_TASKWAIT);
matcho ("taskyield", gfc_match_omp_taskyield, ST_OMP_TASKYIELD);
matcho ("task", gfc_match_omp_task, ST_OMP_TASK);
+ matchs ("teams distribute parallel do simd",
+ gfc_match_omp_teams_distribute_parallel_do_simd,
+ ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
+ matcho ("teams distribute parallel do",
+ gfc_match_omp_teams_distribute_parallel_do,
+ ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO);
+ matchs ("teams distribute simd", gfc_match_omp_teams_distribute_simd,
+ ST_OMP_TEAMS_DISTRIBUTE_SIMD);
+ matcho ("teams distribute", gfc_match_omp_teams_distribute,
+ ST_OMP_TEAMS_DISTRIBUTE);
+ matcho ("teams", gfc_match_omp_teams, ST_OMP_TEAMS);
matcho ("threadprivate", gfc_match_omp_threadprivate,
ST_OMP_THREADPRIVATE);
break;
case ST_LABEL_ASSIGNMENT: case ST_FLUSH: case ST_OMP_FLUSH: \
case ST_OMP_BARRIER: case ST_OMP_TASKWAIT: case ST_OMP_TASKYIELD: \
case ST_OMP_CANCEL: case ST_OMP_CANCELLATION_POINT: \
- case ST_ERROR_STOP: case ST_SYNC_ALL: case ST_SYNC_IMAGES: \
- case ST_SYNC_MEMORY: case ST_LOCK: case ST_UNLOCK
+ case ST_OMP_TARGET_UPDATE: case ST_ERROR_STOP: case ST_SYNC_ALL: \
+ case ST_SYNC_IMAGES: case ST_SYNC_MEMORY: case ST_LOCK: case ST_UNLOCK
/* Statements that mark other executable statements. */
case ST_OMP_DO: case ST_OMP_PARALLEL_DO: case ST_OMP_ATOMIC: \
case ST_OMP_WORKSHARE: case ST_OMP_PARALLEL_WORKSHARE: \
case ST_OMP_TASK: case ST_OMP_TASKGROUP: case ST_OMP_SIMD: \
- case ST_OMP_DO_SIMD: case ST_OMP_PARALLEL_DO_SIMD: case ST_CRITICAL
+ case ST_OMP_DO_SIMD: case ST_OMP_PARALLEL_DO_SIMD: case ST_OMP_TARGET: \
+ case ST_OMP_TARGET_DATA: case ST_OMP_TARGET_TEAMS: \
+ case ST_OMP_TARGET_TEAMS_DISTRIBUTE: \
+ case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: \
+ case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: \
+ case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: \
+ case ST_OMP_TEAMS: case ST_OMP_TEAMS_DISTRIBUTE: \
+ case ST_OMP_TEAMS_DISTRIBUTE_SIMD: \
+ case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: \
+ case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: case ST_OMP_DISTRIBUTE: \
+ case ST_OMP_DISTRIBUTE_SIMD: case ST_OMP_DISTRIBUTE_PARALLEL_DO: \
+ case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: \
+ case ST_CRITICAL
/* Declaration statements */
#define case_decl case ST_ATTR_DECL: case ST_COMMON: case ST_DATA_DECL: \
case ST_EQUIVALENCE: case ST_NAMELIST: case ST_STATEMENT_FUNCTION: \
case ST_TYPE: case ST_INTERFACE: case ST_OMP_THREADPRIVATE: \
- case ST_PROCEDURE: case ST_OMP_DECLARE_SIMD: case ST_OMP_DECLARE_REDUCTION
+ case ST_PROCEDURE: case ST_OMP_DECLARE_SIMD: case ST_OMP_DECLARE_REDUCTION: \
+ case ST_OMP_DECLARE_TARGET
/* Block end statements. Errors associated with interchanging these
are detected in gfc_match_end(). */
case ST_OMP_DECLARE_SIMD:
p = "!$OMP DECLARE SIMD";
break;
+ case ST_OMP_DECLARE_TARGET:
+ p = "!$OMP DECLARE TARGET";
+ break;
+ case ST_OMP_DISTRIBUTE:
+ p = "!$OMP DISTRIBUTE";
+ break;
+ case ST_OMP_DISTRIBUTE_PARALLEL_DO:
+ p = "!$OMP DISTRIBUTE PARALLEL DO";
+ break;
+ case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ p = "!$OMP DISTRIBUTE PARALLEL DO SIMD";
+ break;
+ case ST_OMP_DISTRIBUTE_SIMD:
+ p = "!$OMP DISTRIBUTE SIMD";
+ break;
case ST_OMP_DO:
p = "!$OMP DO";
break;
case ST_OMP_END_CRITICAL:
p = "!$OMP END CRITICAL";
break;
+ case ST_OMP_END_DISTRIBUTE:
+ p = "!$OMP END DISTRIBUTE";
+ break;
+ case ST_OMP_END_DISTRIBUTE_PARALLEL_DO:
+ p = "!$OMP END DISTRIBUTE PARALLEL DO";
+ break;
+ case ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD:
+ p = "!$OMP END DISTRIBUTE PARALLEL DO SIMD";
+ break;
+ case ST_OMP_END_DISTRIBUTE_SIMD:
+ p = "!$OMP END DISTRIBUTE SIMD";
+ break;
case ST_OMP_END_DO:
p = "!$OMP END DO";
break;
case ST_OMP_END_TASK:
p = "!$OMP END TASK";
break;
+ case ST_OMP_END_TARGET:
+ p = "!$OMP END TARGET";
+ break;
+ case ST_OMP_END_TARGET_DATA:
+ p = "!$OMP END TARGET DATA";
+ break;
+ case ST_OMP_END_TARGET_TEAMS:
+ p = "!$OMP END TARGET TEAMS";
+ break;
+ case ST_OMP_END_TARGET_TEAMS_DISTRIBUTE:
+ p = "!$OMP END TARGET TEAMS DISTRIBUTE";
+ break;
+ case ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ p = "!$OMP END TARGET TEAMS DISTRIBUTE PARALLEL DO";
+ break;
+ case ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ p = "!$OMP END TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD";
+ break;
+ case ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ p = "!$OMP END TARGET TEAMS DISTRIBUTE SIMD";
+ break;
case ST_OMP_END_TASKGROUP:
p = "!$OMP END TASKGROUP";
break;
+ case ST_OMP_END_TEAMS:
+ p = "!$OMP END TEAMS";
+ break;
+ case ST_OMP_END_TEAMS_DISTRIBUTE:
+ p = "!$OMP END TEAMS DISTRIBUTE";
+ break;
+ case ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ p = "!$OMP END TEAMS DISTRIBUTE PARALLEL DO";
+ break;
+ case ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ p = "!$OMP END TEAMS DISTRIBUTE PARALLEL DO SIMD";
+ break;
+ case ST_OMP_END_TEAMS_DISTRIBUTE_SIMD:
+ p = "!$OMP END TEAMS DISTRIBUTE SIMD";
+ break;
case ST_OMP_END_WORKSHARE:
p = "!$OMP END WORKSHARE";
break;
case ST_OMP_SINGLE:
p = "!$OMP SINGLE";
break;
+ case ST_OMP_TARGET:
+ p = "!$OMP TARGET";
+ break;
+ case ST_OMP_TARGET_DATA:
+ p = "!$OMP TARGET DATA";
+ break;
+ case ST_OMP_TARGET_TEAMS:
+ p = "!$OMP TARGET TEAMS";
+ break;
+ case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
+ p = "!$OMP TARGET TEAMS DISTRIBUTE";
+ break;
+ case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ p = "!$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO";
+ break;
+ case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ p = "!$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD";
+ break;
+ case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ p = "!$OMP TARGET TEAMS DISTRIBUTE SIMD";
+ break;
+ case ST_OMP_TARGET_UPDATE:
+ p = "!$OMP TARGET UPDATE";
+ break;
case ST_OMP_TASK:
p = "!$OMP TASK";
break;
case ST_OMP_TASKYIELD:
p = "!$OMP TASKYIELD";
break;
+ case ST_OMP_TEAMS:
+ p = "!$OMP TEAMS";
+ break;
+ case ST_OMP_TEAMS_DISTRIBUTE:
+ p = "!$OMP TEAMS DISTRIBUTE";
+ break;
+ case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ p = "!$OMP TEAMS DISTRIBUTE PARALLEL DO";
+ break;
+ case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ p = "!$OMP TEAMS DISTRIBUTE PARALLEL DO SIMD";
+ break;
+ case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
+ p = "!$OMP TEAMS DISTRIBUTE SIMD";
+ break;
case ST_OMP_THREADPRIVATE:
p = "!$OMP THREADPRIVATE";
break;
gfc_statement omp_end_st = ST_OMP_END_DO;
switch (omp_st)
{
- case ST_OMP_SIMD: omp_end_st = ST_OMP_END_SIMD; break;
+ 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_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_SIMD: omp_end_st = ST_OMP_END_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_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;
default: gcc_unreachable ();
}
if (st == omp_end_st)
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_TEAMS:
+ omp_end_st = ST_OMP_END_TARGET_TEAMS;
+ 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_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_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_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_WORKSHARE:
omp_end_st = ST_OMP_END_WORKSHARE;
break;
case ST_OMP_CRITICAL:
case ST_OMP_MASTER:
case ST_OMP_SINGLE:
+ case ST_OMP_TARGET:
+ case ST_OMP_TARGET_DATA:
+ case ST_OMP_TARGET_TEAMS:
+ case ST_OMP_TEAMS:
case ST_OMP_TASK:
case ST_OMP_TASKGROUP:
parse_omp_structured_block (st, false);
parse_omp_structured_block (st, true);
break;
+ 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_PARALLEL_DO:
case ST_OMP_PARALLEL_DO_SIMD:
case ST_OMP_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_TEAMS_DISTRIBUTE:
+ case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
st = parse_omp_do (st);
if (st == ST_IMPLIED_ENDDO)
return st;
case EXEC_OMP_ATOMIC:
case EXEC_OMP_CRITICAL:
+ case EXEC_OMP_DISTRIBUTE:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_DISTRIBUTE_SIMD:
case EXEC_OMP_DO:
case EXEC_OMP_DO_SIMD:
case EXEC_OMP_MASTER:
case EXEC_OMP_SECTIONS:
case EXEC_OMP_SIMD:
case EXEC_OMP_SINGLE:
+ case EXEC_OMP_TARGET:
+ case EXEC_OMP_TARGET_DATA:
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TARGET_UPDATE:
case EXEC_OMP_TASK:
case EXEC_OMP_TASKGROUP:
case EXEC_OMP_TASKWAIT:
case EXEC_OMP_TASKYIELD:
+ case EXEC_OMP_TEAMS:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
case EXEC_OMP_WORKSHARE:
break;
case EXEC_OMP_PARALLEL_DO:
case EXEC_OMP_PARALLEL_DO_SIMD:
case EXEC_OMP_PARALLEL_SECTIONS:
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
case EXEC_OMP_TASK:
+ case EXEC_OMP_TEAMS:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
omp_workshare_save = omp_workshare_flag;
omp_workshare_flag = 0;
gfc_resolve_omp_parallel_blocks (code, ns);
break;
+ case EXEC_OMP_DISTRIBUTE:
+ case EXEC_OMP_DISTRIBUTE_SIMD:
case EXEC_OMP_DO:
case EXEC_OMP_DO_SIMD:
case EXEC_OMP_SIMD:
case EXEC_OMP_CANCELLATION_POINT:
case EXEC_OMP_CRITICAL:
case EXEC_OMP_FLUSH:
+ case EXEC_OMP_DISTRIBUTE:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_DISTRIBUTE_SIMD:
case EXEC_OMP_DO:
case EXEC_OMP_DO_SIMD:
case EXEC_OMP_MASTER:
case EXEC_OMP_SECTIONS:
case EXEC_OMP_SIMD:
case EXEC_OMP_SINGLE:
+ case EXEC_OMP_TARGET:
+ case EXEC_OMP_TARGET_DATA:
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TARGET_UPDATE:
+ case EXEC_OMP_TASK:
case EXEC_OMP_TASKGROUP:
case EXEC_OMP_TASKWAIT:
case EXEC_OMP_TASKYIELD:
+ case EXEC_OMP_TEAMS:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
case EXEC_OMP_WORKSHARE:
gfc_resolve_omp_directive (code, ns);
break;
case EXEC_OMP_PARALLEL_DO_SIMD:
case EXEC_OMP_PARALLEL_SECTIONS:
case EXEC_OMP_PARALLEL_WORKSHARE:
- case EXEC_OMP_TASK:
omp_workshare_save = omp_workshare_flag;
omp_workshare_flag = 0;
gfc_resolve_omp_directive (code, ns);
|| sym->ns->proc_name->attr.flavor != FL_MODULE)))
gfc_error ("Threadprivate at %L isn't SAVEd", &sym->declared_at);
+ /* Check omp declare target restrictions. */
+ if (sym->attr.omp_declare_target
+ && sym->attr.flavor == FL_VARIABLE
+ && !sym->attr.save
+ && !sym->ns->save_all
+ && (!sym->attr.in_common
+ && sym->module == NULL
+ && (sym->ns->proc_name == NULL
+ || sym->ns->proc_name->attr.flavor != FL_MODULE)))
+ gfc_error ("!$OMP DECLARE TARGET variable '%s' at %L isn't SAVEd",
+ sym->name, &sym->declared_at);
+
/* If we have come this far we can apply default-initializers, as
described in 14.7.5, to those variables that have not already
been assigned one. */
case EXEC_OMP_CANCEL:
case EXEC_OMP_CANCELLATION_POINT:
+ case EXEC_OMP_DISTRIBUTE:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_DISTRIBUTE_SIMD:
case EXEC_OMP_DO:
case EXEC_OMP_DO_SIMD:
case EXEC_OMP_END_SINGLE:
case EXEC_OMP_SECTIONS:
case EXEC_OMP_SIMD:
case EXEC_OMP_SINGLE:
+ case EXEC_OMP_TARGET:
+ case EXEC_OMP_TARGET_DATA:
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TARGET_UPDATE:
case EXEC_OMP_TASK:
+ case EXEC_OMP_TEAMS:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
case EXEC_OMP_WORKSHARE:
case EXEC_OMP_PARALLEL_WORKSHARE:
gfc_free_omp_clauses (p->ext.omp_clauses);
*asynchronous = "ASYNCHRONOUS", *codimension = "CODIMENSION",
*contiguous = "CONTIGUOUS", *generic = "GENERIC";
static const char *threadprivate = "THREADPRIVATE";
+ static const char *omp_declare_target = "OMP DECLARE TARGET";
const char *a1, *a2;
int standard;
conf (dummy, entry);
conf (dummy, intrinsic);
conf (dummy, threadprivate);
+ conf (dummy, omp_declare_target);
conf (pointer, target);
conf (pointer, intrinsic);
conf (pointer, elemental);
conf (in_equivalence, entry);
conf (in_equivalence, allocatable);
conf (in_equivalence, threadprivate);
+ conf (in_equivalence, omp_declare_target);
conf (dummy, result);
conf (entry, result);
conf (cray_pointee, in_common);
conf (cray_pointee, in_equivalence);
conf (cray_pointee, threadprivate);
+ conf (cray_pointee, omp_declare_target);
conf (data, dummy);
conf (data, function);
conf (proc_pointer, abstract)
+ conf (entry, omp_declare_target)
+
a1 = gfc_code2string (flavors, attr->flavor);
if (attr->in_namelist
conf2 (function);
conf2 (subroutine);
conf2 (threadprivate);
+ conf2 (omp_declare_target);
if (attr->access == ACCESS_PUBLIC || attr->access == ACCESS_PRIVATE)
{
conf2 (subroutine);
conf2 (threadprivate);
conf2 (result);
+ conf2 (omp_declare_target);
if (attr->intent != INTENT_UNKNOWN)
{
bool
+gfc_add_omp_declare_target (symbol_attribute *attr, const char *name,
+ locus *where)
+{
+
+ if (check_used (attr, name, where))
+ return false;
+
+ if (attr->omp_declare_target)
+ return true;
+
+ attr->omp_declare_target = 1;
+ return check_conflict (attr, name, where);
+}
+
+
+bool
gfc_add_target (symbol_attribute *attr, locus *where)
{
if (src->threadprivate
&& !gfc_add_threadprivate (dest, NULL, where))
goto fail;
+ if (src->omp_declare_target
+ && !gfc_add_omp_declare_target (dest, NULL, where))
+ goto fail;
if (src->target && !gfc_add_target (dest, where))
goto fail;
if (src->dummy && !gfc_add_dummy (dest, NULL, where))
if (com->threadprivate)
set_decl_tls_model (decl, decl_default_tls_model (decl));
+ if (com->omp_declare_target)
+ DECL_ATTRIBUTES (decl)
+ = tree_cons (get_identifier ("omp declare target"),
+ NULL_TREE, DECL_ATTRIBUTES (decl));
+
/* Place the back end declaration for this common block in
GLOBAL_BINDING_LEVEL. */
gfc_map_of_all_commons[identifier] = pushdecl_top_level (decl);
list = chainon (list, attr);
}
+ if (sym_attr.omp_declare_target)
+ list = tree_cons (get_identifier ("omp declare target"),
+ NULL_TREE, list);
+
return list;
}
}
+void
+gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
+{
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ return;
+
+ tree decl = OMP_CLAUSE_DECL (c);
+ tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
+ if (POINTER_TYPE_P (TREE_TYPE (decl)))
+ {
+ if (!gfc_omp_privatize_by_reference (decl)
+ && !GFC_DECL_GET_SCALAR_POINTER (decl)
+ && !GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
+ && !GFC_DECL_CRAY_POINTEE (decl)
+ && !GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
+ return;
+ c4 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (c4) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_DECL (c4) = decl;
+ OMP_CLAUSE_SIZE (c4) = size_int (0);
+ decl = build_fold_indirect_ref (decl);
+ OMP_CLAUSE_DECL (c) = decl;
+ }
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ {
+ stmtblock_t block;
+ gfc_start_block (&block);
+ tree type = TREE_TYPE (decl);
+ tree ptr = gfc_conv_descriptor_data_get (decl);
+ ptr = fold_convert (build_pointer_type (char_type_node), ptr);
+ ptr = build_fold_indirect_ref (ptr);
+ OMP_CLAUSE_DECL (c) = ptr;
+ c2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_TO_PSET;
+ OMP_CLAUSE_DECL (c2) = decl;
+ OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (type);
+ c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (c3) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_DECL (c3) = gfc_conv_descriptor_data_get (decl);
+ OMP_CLAUSE_SIZE (c3) = size_int (0);
+ tree size = create_tmp_var (gfc_array_index_type, NULL);
+ tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ if (GFC_TYPE_ARRAY_AKIND (type) == GFC_ARRAY_POINTER
+ || GFC_TYPE_ARRAY_AKIND (type) == GFC_ARRAY_POINTER_CONT)
+ {
+ stmtblock_t cond_block;
+ tree tem, then_b, else_b, zero, cond;
+
+ gfc_init_block (&cond_block);
+ tem = gfc_full_array_size (&cond_block, decl,
+ GFC_TYPE_ARRAY_RANK (type));
+ gfc_add_modify (&cond_block, size, tem);
+ gfc_add_modify (&cond_block, size,
+ fold_build2 (MULT_EXPR, gfc_array_index_type,
+ size, elemsz));
+ then_b = gfc_finish_block (&cond_block);
+ gfc_init_block (&cond_block);
+ zero = build_int_cst (gfc_array_index_type, 0);
+ gfc_add_modify (&cond_block, size, zero);
+ else_b = gfc_finish_block (&cond_block);
+ tem = gfc_conv_descriptor_data_get (decl);
+ tem = fold_convert (pvoid_type_node, tem);
+ cond = fold_build2_loc (input_location, NE_EXPR,
+ boolean_type_node, tem, null_pointer_node);
+ gfc_add_expr_to_block (&block, build3_loc (input_location, COND_EXPR,
+ void_type_node, cond,
+ then_b, else_b));
+ }
+ else
+ {
+ gfc_add_modify (&block, size,
+ gfc_full_array_size (&block, decl,
+ GFC_TYPE_ARRAY_RANK (type)));
+ gfc_add_modify (&block, size,
+ fold_build2 (MULT_EXPR, gfc_array_index_type,
+ size, elemsz));
+ }
+ OMP_CLAUSE_SIZE (c) = size;
+ tree stmt = gfc_finish_block (&block);
+ gimplify_and_add (stmt, pre_p);
+ }
+ tree last = c;
+ if (c2)
+ {
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (last);
+ OMP_CLAUSE_CHAIN (last) = c2;
+ last = c2;
+ }
+ if (c3)
+ {
+ OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (last);
+ OMP_CLAUSE_CHAIN (last) = c3;
+ last = c3;
+ }
+ if (c4)
+ {
+ OMP_CLAUSE_CHAIN (c4) = OMP_CLAUSE_CHAIN (last);
+ OMP_CLAUSE_CHAIN (last) = c4;
+ last = c4;
+ }
+}
+
+
/* Return true if DECL's DECL_VALUE_EXPR (if any) should be
disregarded in OpenMP construct, because it is going to be
remapped during OpenMP lowering. SHARED is true if DECL
tree node = build_omp_clause (where.lb->location,
OMP_CLAUSE_REDUCTION);
OMP_CLAUSE_DECL (node) = t;
- switch (namelist->rop)
+ switch (namelist->u.reduction_op)
{
case OMP_REDUCTION_PLUS:
OMP_CLAUSE_REDUCTION_CODE (node) = PLUS_EXPR;
gcc_unreachable ();
}
if (namelist->sym->attr.dimension
- || namelist->rop == OMP_REDUCTION_USER
+ || namelist->u.reduction_op == OMP_REDUCTION_USER
|| namelist->sym->attr.allocatable)
gfc_trans_omp_array_reduction_or_udr (node, namelist, where);
list = gfc_trans_add_clause (node, list);
}
}
break;
- case OMP_LIST_DEPEND_IN:
- case OMP_LIST_DEPEND_OUT:
+ case OMP_LIST_DEPEND:
for (; n != NULL; n = n->next)
{
if (!n->sym->attr.referenced)
tree node = build_omp_clause (input_location, OMP_CLAUSE_DEPEND);
if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
{
- OMP_CLAUSE_DECL (node) = gfc_get_symbol_decl (n->sym);
- if (DECL_P (OMP_CLAUSE_DECL (node)))
- TREE_ADDRESSABLE (OMP_CLAUSE_DECL (node)) = 1;
+ tree decl = gfc_get_symbol_decl (n->sym);
+ if (gfc_omp_privatize_by_reference (decl))
+ decl = build_fold_indirect_ref (decl);
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ {
+ decl = gfc_conv_descriptor_data_get (decl);
+ decl = fold_convert (build_pointer_type (char_type_node),
+ decl);
+ decl = build_fold_indirect_ref (decl);
+ }
+ else if (DECL_P (decl))
+ TREE_ADDRESSABLE (decl) = 1;
+ OMP_CLAUSE_DECL (node) = decl;
}
else
{
}
gfc_add_block_to_block (block, &se.pre);
gfc_add_block_to_block (block, &se.post);
- OMP_CLAUSE_DECL (node)
- = fold_build1_loc (input_location, INDIRECT_REF,
- TREE_TYPE (TREE_TYPE (ptr)), ptr);
+ ptr = fold_convert (build_pointer_type (char_type_node),
+ ptr);
+ OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+ }
+ switch (n->u.depend_op)
+ {
+ case OMP_DEPEND_IN:
+ OMP_CLAUSE_DEPEND_KIND (node) = OMP_CLAUSE_DEPEND_IN;
+ break;
+ case OMP_DEPEND_OUT:
+ OMP_CLAUSE_DEPEND_KIND (node) = OMP_CLAUSE_DEPEND_OUT;
+ break;
+ case OMP_DEPEND_INOUT:
+ OMP_CLAUSE_DEPEND_KIND (node) = OMP_CLAUSE_DEPEND_INOUT;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+ }
+ break;
+ case OMP_LIST_MAP:
+ for (; n != NULL; n = n->next)
+ {
+ if (!n->sym->attr.referenced)
+ continue;
+
+ tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ tree node2 = NULL_TREE;
+ tree node3 = NULL_TREE;
+ tree node4 = NULL_TREE;
+ tree decl = gfc_get_symbol_decl (n->sym);
+ if (DECL_P (decl))
+ TREE_ADDRESSABLE (decl) = 1;
+ if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
+ {
+ if (POINTER_TYPE_P (TREE_TYPE (decl)))
+ {
+ node4 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (node4) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_DECL (node4) = decl;
+ OMP_CLAUSE_SIZE (node4) = size_int (0);
+ decl = build_fold_indirect_ref (decl);
+ }
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ {
+ tree type = TREE_TYPE (decl);
+ tree ptr = gfc_conv_descriptor_data_get (decl);
+ ptr = fold_convert (build_pointer_type (char_type_node),
+ ptr);
+ ptr = build_fold_indirect_ref (ptr);
+ OMP_CLAUSE_DECL (node) = ptr;
+ node2 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (node2) = OMP_CLAUSE_MAP_TO_PSET;
+ OMP_CLAUSE_DECL (node2) = decl;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ node3 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (node3) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_DECL (node3)
+ = gfc_conv_descriptor_data_get (decl);
+ OMP_CLAUSE_SIZE (node3) = size_int (0);
+ if (n->sym->attr.pointer)
+ {
+ stmtblock_t cond_block;
+ tree size
+ = gfc_create_var (gfc_array_index_type, NULL);
+ tree tem, then_b, else_b, zero, cond;
+
+ gfc_init_block (&cond_block);
+ tem
+ = gfc_full_array_size (&cond_block, decl,
+ GFC_TYPE_ARRAY_RANK (type));
+ gfc_add_modify (&cond_block, size, tem);
+ then_b = gfc_finish_block (&cond_block);
+ gfc_init_block (&cond_block);
+ zero = build_int_cst (gfc_array_index_type, 0);
+ gfc_add_modify (&cond_block, size, zero);
+ else_b = gfc_finish_block (&cond_block);
+ tem = gfc_conv_descriptor_data_get (decl);
+ tem = fold_convert (pvoid_type_node, tem);
+ cond = fold_build2_loc (input_location, NE_EXPR,
+ boolean_type_node,
+ tem, null_pointer_node);
+ gfc_add_expr_to_block (block,
+ build3_loc (input_location,
+ COND_EXPR,
+ void_type_node,
+ cond, then_b,
+ else_b));
+ OMP_CLAUSE_SIZE (node) = size;
+ }
+ else
+ OMP_CLAUSE_SIZE (node)
+ = gfc_full_array_size (block, decl,
+ GFC_TYPE_ARRAY_RANK (type));
+ tree elemsz
+ = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ OMP_CLAUSE_SIZE (node)
+ = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ OMP_CLAUSE_SIZE (node), elemsz);
+ }
+ else
+ OMP_CLAUSE_DECL (node) = decl;
+ }
+ else
+ {
+ tree ptr, ptr2;
+ gfc_init_se (&se, NULL);
+ if (n->expr->ref->u.ar.type == AR_ELEMENT)
+ {
+ gfc_conv_expr_reference (&se, n->expr);
+ gfc_add_block_to_block (block, &se.pre);
+ ptr = se.expr;
+ OMP_CLAUSE_SIZE (node)
+ = TYPE_SIZE_UNIT (TREE_TYPE (ptr));
+ }
+ else
+ {
+ gfc_conv_expr_descriptor (&se, n->expr);
+ ptr = gfc_conv_array_data (se.expr);
+ tree type = TREE_TYPE (se.expr);
+ gfc_add_block_to_block (block, &se.pre);
+ OMP_CLAUSE_SIZE (node)
+ = gfc_full_array_size (block, se.expr,
+ GFC_TYPE_ARRAY_RANK (type));
+ tree elemsz
+ = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ OMP_CLAUSE_SIZE (node)
+ = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ OMP_CLAUSE_SIZE (node), elemsz);
+ }
+ gfc_add_block_to_block (block, &se.post);
+ ptr = fold_convert (build_pointer_type (char_type_node),
+ ptr);
+ OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+
+ if (POINTER_TYPE_P (TREE_TYPE (decl))
+ && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
+ {
+ node4 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (node4) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_DECL (node4) = decl;
+ OMP_CLAUSE_SIZE (node4) = size_int (0);
+ decl = build_fold_indirect_ref (decl);
+ }
+ ptr = fold_convert (sizetype, ptr);
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ {
+ tree type = TREE_TYPE (decl);
+ ptr2 = gfc_conv_descriptor_data_get (decl);
+ node2 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (node2) = OMP_CLAUSE_MAP_TO_PSET;
+ OMP_CLAUSE_DECL (node2) = decl;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ node3 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (node3) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_DECL (node3)
+ = gfc_conv_descriptor_data_get (decl);
+ }
+ else
+ {
+ if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ ptr2 = build_fold_addr_expr (decl);
+ else
+ {
+ gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
+ ptr2 = decl;
+ }
+ node3 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (node3) = OMP_CLAUSE_MAP_POINTER;
+ OMP_CLAUSE_DECL (node3) = decl;
+ }
+ ptr2 = fold_convert (sizetype, ptr2);
+ OMP_CLAUSE_SIZE (node3)
+ = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+ }
+ switch (n->u.map_op)
+ {
+ case OMP_MAP_ALLOC:
+ OMP_CLAUSE_MAP_KIND (node) = OMP_CLAUSE_MAP_ALLOC;
+ break;
+ case OMP_MAP_TO:
+ OMP_CLAUSE_MAP_KIND (node) = OMP_CLAUSE_MAP_TO;
+ break;
+ case OMP_MAP_FROM:
+ OMP_CLAUSE_MAP_KIND (node) = OMP_CLAUSE_MAP_FROM;
+ break;
+ case OMP_MAP_TOFROM:
+ OMP_CLAUSE_MAP_KIND (node) = OMP_CLAUSE_MAP_TOFROM;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+ if (node2)
+ omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
+ if (node3)
+ omp_clauses = gfc_trans_add_clause (node3, omp_clauses);
+ if (node4)
+ omp_clauses = gfc_trans_add_clause (node4, omp_clauses);
+ }
+ break;
+ case OMP_LIST_TO:
+ case OMP_LIST_FROM:
+ for (; n != NULL; n = n->next)
+ {
+ if (!n->sym->attr.referenced)
+ continue;
+
+ tree node = build_omp_clause (input_location,
+ list == OMP_LIST_TO
+ ? OMP_CLAUSE_TO : OMP_CLAUSE_FROM);
+ if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
+ {
+ tree decl = gfc_get_symbol_decl (n->sym);
+ if (gfc_omp_privatize_by_reference (decl))
+ decl = build_fold_indirect_ref (decl);
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ {
+ tree type = TREE_TYPE (decl);
+ tree ptr = gfc_conv_descriptor_data_get (decl);
+ ptr = fold_convert (build_pointer_type (char_type_node),
+ ptr);
+ ptr = build_fold_indirect_ref (ptr);
+ OMP_CLAUSE_DECL (node) = ptr;
+ OMP_CLAUSE_SIZE (node)
+ = gfc_full_array_size (block, decl,
+ GFC_TYPE_ARRAY_RANK (type));
+ tree elemsz
+ = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ OMP_CLAUSE_SIZE (node)
+ = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ OMP_CLAUSE_SIZE (node), elemsz);
+ }
+ else
+ OMP_CLAUSE_DECL (node) = decl;
+ }
+ else
+ {
+ tree ptr;
+ gfc_init_se (&se, NULL);
+ if (n->expr->ref->u.ar.type == AR_ELEMENT)
+ {
+ gfc_conv_expr_reference (&se, n->expr);
+ ptr = se.expr;
+ gfc_add_block_to_block (block, &se.pre);
+ OMP_CLAUSE_SIZE (node)
+ = TYPE_SIZE_UNIT (TREE_TYPE (ptr));
+ }
+ else
+ {
+ gfc_conv_expr_descriptor (&se, n->expr);
+ ptr = gfc_conv_array_data (se.expr);
+ tree type = TREE_TYPE (se.expr);
+ gfc_add_block_to_block (block, &se.pre);
+ OMP_CLAUSE_SIZE (node)
+ = gfc_full_array_size (block, se.expr,
+ GFC_TYPE_ARRAY_RANK (type));
+ tree elemsz
+ = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ OMP_CLAUSE_SIZE (node)
+ = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ OMP_CLAUSE_SIZE (node), elemsz);
+ }
+ gfc_add_block_to_block (block, &se.post);
+ ptr = fold_convert (build_pointer_type (char_type_node),
+ ptr);
+ OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
}
- OMP_CLAUSE_DEPEND_KIND (node)
- = ((list == OMP_LIST_DEPEND_IN)
- ? OMP_CLAUSE_DEPEND_IN : OMP_CLAUSE_DEPEND_OUT);
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
}
break;
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
- return omp_clauses;
+ if (clauses->num_teams)
+ {
+ tree num_teams;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->num_teams);
+ gfc_add_block_to_block (block, &se.pre);
+ num_teams = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_NUM_TEAMS);
+ OMP_CLAUSE_NUM_TEAMS_EXPR (c) = num_teams;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
+ if (clauses->device)
+ {
+ tree device;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->device);
+ gfc_add_block_to_block (block, &se.pre);
+ device = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_DEVICE);
+ OMP_CLAUSE_DEVICE_ID (c) = device;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
+ if (clauses->thread_limit)
+ {
+ tree thread_limit;
+
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->thread_limit);
+ gfc_add_block_to_block (block, &se.pre);
+ thread_limit = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_THREAD_LIMIT);
+ OMP_CLAUSE_THREAD_LIMIT_EXPR (c) = thread_limit;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
+ chunk_size = NULL_TREE;
+ if (clauses->dist_chunk_size)
+ {
+ gfc_init_se (&se, NULL);
+ gfc_conv_expr (&se, clauses->dist_chunk_size);
+ gfc_add_block_to_block (block, &se.pre);
+ chunk_size = gfc_evaluate_now (se.expr, block);
+ gfc_add_block_to_block (block, &se.post);
+ }
+
+ if (clauses->dist_sched_kind != OMP_SCHED_NONE)
+ {
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_DIST_SCHEDULE);
+ OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (c) = chunk_size;
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+
+ return nreverse (omp_clauses);
}
/* Like gfc_trans_code, but force creation of a BIND_EXPR around it. */
if (clauses)
{
- gfc_omp_namelist *n;
- for (n = clauses->lists[(op == EXEC_OMP_SIMD && collapse == 1)
- ? OMP_LIST_LINEAR : OMP_LIST_LASTPRIVATE];
- n != NULL; n = n->next)
- if (code->ext.iterator->var->symtree->n.sym == n->sym)
- break;
+ gfc_omp_namelist *n = NULL;
+ if (op != EXEC_OMP_DISTRIBUTE)
+ for (n = clauses->lists[(op == EXEC_OMP_SIMD && collapse == 1)
+ ? OMP_LIST_LINEAR : OMP_LIST_LASTPRIVATE];
+ n != NULL; n = n->next)
+ if (code->ext.iterator->var->symtree->n.sym == n->sym)
+ break;
if (n != NULL)
dovar_found = 1;
else if (n == NULL && op != EXEC_OMP_SIMD)
}
/* End of loop body. */
- stmt = make_node (op == EXEC_OMP_SIMD ? OMP_SIMD : OMP_FOR);
+ switch (op)
+ {
+ case EXEC_OMP_SIMD: stmt = make_node (OMP_SIMD); break;
+ case EXEC_OMP_DO: stmt = make_node (OMP_FOR); break;
+ case EXEC_OMP_DISTRIBUTE: stmt = make_node (OMP_DISTRIBUTE); break;
+ default: gcc_unreachable ();
+ }
TREE_TYPE (stmt) = void_type_node;
OMP_FOR_BODY (stmt) = gfc_finish_block (&body);
GFC_OMP_SPLIT_SIMD,
GFC_OMP_SPLIT_DO,
GFC_OMP_SPLIT_PARALLEL,
+ GFC_OMP_SPLIT_DISTRIBUTE,
+ GFC_OMP_SPLIT_TEAMS,
+ GFC_OMP_SPLIT_TARGET,
GFC_OMP_SPLIT_NUM
};
{
GFC_OMP_MASK_SIMD = (1 << GFC_OMP_SPLIT_SIMD),
GFC_OMP_MASK_DO = (1 << GFC_OMP_SPLIT_DO),
- GFC_OMP_MASK_PARALLEL = (1 << GFC_OMP_SPLIT_PARALLEL)
+ GFC_OMP_MASK_PARALLEL = (1 << GFC_OMP_SPLIT_PARALLEL),
+ GFC_OMP_MASK_DISTRIBUTE = (1 << GFC_OMP_SPLIT_DISTRIBUTE),
+ GFC_OMP_MASK_TEAMS = (1 << GFC_OMP_SPLIT_TEAMS),
+ GFC_OMP_MASK_TARGET = (1 << GFC_OMP_SPLIT_TARGET)
};
static void
memset (clausesa, 0, GFC_OMP_SPLIT_NUM * sizeof (gfc_omp_clauses));
switch (code->op)
{
+ case EXEC_OMP_DISTRIBUTE:
+ innermost = GFC_OMP_SPLIT_DISTRIBUTE;
+ break;
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ mask = GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO;
+ innermost = GFC_OMP_SPLIT_DO;
+ break;
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ mask = GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_PARALLEL
+ | GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
+ innermost = GFC_OMP_SPLIT_SIMD;
+ break;
+ case EXEC_OMP_DISTRIBUTE_SIMD:
+ mask = GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_SIMD;
+ innermost = GFC_OMP_SPLIT_SIMD;
+ break;
+ case EXEC_OMP_DO:
+ innermost = GFC_OMP_SPLIT_DO;
+ break;
case EXEC_OMP_DO_SIMD:
mask = GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
innermost = GFC_OMP_SPLIT_SIMD;
break;
+ case EXEC_OMP_PARALLEL:
+ innermost = GFC_OMP_SPLIT_PARALLEL;
+ break;
case EXEC_OMP_PARALLEL_DO:
mask = GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO;
innermost = GFC_OMP_SPLIT_DO;
mask = GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
innermost = GFC_OMP_SPLIT_SIMD;
break;
+ case EXEC_OMP_SIMD:
+ innermost = GFC_OMP_SPLIT_SIMD;
+ break;
+ case EXEC_OMP_TARGET:
+ innermost = GFC_OMP_SPLIT_TARGET;
+ break;
+ case EXEC_OMP_TARGET_TEAMS:
+ mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS;
+ innermost = GFC_OMP_SPLIT_TEAMS;
+ break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS
+ | GFC_OMP_MASK_DISTRIBUTE;
+ innermost = GFC_OMP_SPLIT_DISTRIBUTE;
+ break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE
+ | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO;
+ innermost = GFC_OMP_SPLIT_DO;
+ break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE
+ | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
+ innermost = GFC_OMP_SPLIT_SIMD;
+ break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS
+ | GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_SIMD;
+ innermost = GFC_OMP_SPLIT_SIMD;
+ break;
+ case EXEC_OMP_TEAMS:
+ innermost = GFC_OMP_SPLIT_TEAMS;
+ break;
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ mask = GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE;
+ innermost = GFC_OMP_SPLIT_DISTRIBUTE;
+ break;
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ mask = GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE
+ | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO;
+ innermost = GFC_OMP_SPLIT_DO;
+ break;
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ mask = GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE
+ | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
+ innermost = GFC_OMP_SPLIT_SIMD;
+ break;
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+ mask = GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_SIMD;
+ innermost = GFC_OMP_SPLIT_SIMD;
+ break;
default:
gcc_unreachable ();
}
+ if (mask == 0)
+ {
+ clausesa[innermost] = *code->ext.omp_clauses;
+ return;
+ }
if (code->ext.omp_clauses != NULL)
{
+ if (mask & GFC_OMP_MASK_TARGET)
+ {
+ /* First the clauses that are unique to some constructs. */
+ clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_MAP]
+ = code->ext.omp_clauses->lists[OMP_LIST_MAP];
+ clausesa[GFC_OMP_SPLIT_TARGET].device
+ = code->ext.omp_clauses->device;
+ }
+ if (mask & GFC_OMP_MASK_TEAMS)
+ {
+ /* First the clauses that are unique to some constructs. */
+ clausesa[GFC_OMP_SPLIT_TEAMS].num_teams
+ = code->ext.omp_clauses->num_teams;
+ clausesa[GFC_OMP_SPLIT_TEAMS].thread_limit
+ = code->ext.omp_clauses->thread_limit;
+ /* Shared and default clauses are allowed on parallel and teams. */
+ clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_SHARED]
+ = code->ext.omp_clauses->lists[OMP_LIST_SHARED];
+ clausesa[GFC_OMP_SPLIT_TEAMS].default_sharing
+ = code->ext.omp_clauses->default_sharing;
+ }
+ if (mask & GFC_OMP_MASK_DISTRIBUTE)
+ {
+ /* First the clauses that are unique to some constructs. */
+ clausesa[GFC_OMP_SPLIT_DISTRIBUTE].dist_sched_kind
+ = code->ext.omp_clauses->dist_sched_kind;
+ clausesa[GFC_OMP_SPLIT_DISTRIBUTE].dist_chunk_size
+ = code->ext.omp_clauses->dist_chunk_size;
+ /* Duplicate collapse. */
+ clausesa[GFC_OMP_SPLIT_DISTRIBUTE].collapse
+ = code->ext.omp_clauses->collapse;
+ }
if (mask & GFC_OMP_MASK_PARALLEL)
{
/* First the clauses that are unique to some constructs. */
= code->ext.omp_clauses->lists[OMP_LIST_SHARED];
clausesa[GFC_OMP_SPLIT_PARALLEL].default_sharing
= code->ext.omp_clauses->default_sharing;
- /* FIXME: This is currently being discussed. */
- clausesa[GFC_OMP_SPLIT_PARALLEL].if_expr
- = code->ext.omp_clauses->if_expr;
}
if (mask & GFC_OMP_MASK_DO)
{
/* Firstprivate clause is supported on all constructs but
target and simd. Put it on the outermost of those and
duplicate on parallel. */
+ if (mask & GFC_OMP_MASK_TEAMS)
+ clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_FIRSTPRIVATE]
+ = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE];
+ else if (mask & GFC_OMP_MASK_DISTRIBUTE)
+ clausesa[GFC_OMP_SPLIT_DISTRIBUTE].lists[OMP_LIST_FIRSTPRIVATE]
+ = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE];
if (mask & GFC_OMP_MASK_PARALLEL)
clausesa[GFC_OMP_SPLIT_PARALLEL].lists[OMP_LIST_FIRSTPRIVATE]
= code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE];
/* Reduction is allowed on simd, do, parallel and teams.
Duplicate it on all of them, but omit on do if
parallel is present. */
+ if (mask & GFC_OMP_MASK_TEAMS)
+ clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_REDUCTION]
+ = code->ext.omp_clauses->lists[OMP_LIST_REDUCTION];
if (mask & GFC_OMP_MASK_PARALLEL)
clausesa[GFC_OMP_SPLIT_PARALLEL].lists[OMP_LIST_REDUCTION]
= code->ext.omp_clauses->lists[OMP_LIST_REDUCTION];
if (mask & GFC_OMP_MASK_SIMD)
clausesa[GFC_OMP_SPLIT_SIMD].lists[OMP_LIST_REDUCTION]
= code->ext.omp_clauses->lists[OMP_LIST_REDUCTION];
+ /* FIXME: This is currently being discussed. */
+ if (mask & GFC_OMP_MASK_PARALLEL)
+ clausesa[GFC_OMP_SPLIT_PARALLEL].if_expr
+ = code->ext.omp_clauses->if_expr;
+ else
+ clausesa[GFC_OMP_SPLIT_TARGET].if_expr
+ = code->ext.omp_clauses->if_expr;
}
if ((mask & (GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO))
== (GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO))
}
static tree
-gfc_trans_omp_do_simd (gfc_code *code, gfc_omp_clauses *clausesa,
- tree omp_clauses)
+gfc_trans_omp_do_simd (gfc_code *code, stmtblock_t *pblock,
+ gfc_omp_clauses *clausesa, tree omp_clauses)
{
- stmtblock_t block, *pblock = NULL;
+ stmtblock_t block;
gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
tree stmt, body, omp_do_clauses = NULL_TREE;
- gfc_start_block (&block);
+ if (pblock == NULL)
+ gfc_start_block (&block);
+ else
+ gfc_init_block (&block);
if (clausesa == NULL)
{
if (gfc_option.gfc_flag_openmp)
omp_do_clauses
= gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_DO], code->loc);
- pblock = █
- body = gfc_trans_omp_do (code, EXEC_OMP_SIMD, pblock,
+ body = gfc_trans_omp_do (code, EXEC_OMP_SIMD, pblock ? pblock : &block,
&clausesa[GFC_OMP_SPLIT_SIMD], omp_clauses);
- if (TREE_CODE (body) != BIND_EXPR)
- body = build3_v (BIND_EXPR, NULL, body, poplevel (1, 0));
- else
- poplevel (0, 0);
+ if (pblock == NULL)
+ {
+ if (TREE_CODE (body) != BIND_EXPR)
+ body = build3_v (BIND_EXPR, NULL, body, poplevel (1, 0));
+ else
+ poplevel (0, 0);
+ }
+ else if (TREE_CODE (body) != BIND_EXPR)
+ body = build3_v (BIND_EXPR, NULL, body, NULL_TREE);
if (gfc_option.gfc_flag_openmp)
{
stmt = make_node (OMP_FOR);
}
static tree
-gfc_trans_omp_parallel_do (gfc_code *code)
+gfc_trans_omp_parallel_do (gfc_code *code, stmtblock_t *pblock,
+ gfc_omp_clauses *clausesa)
{
- stmtblock_t block, *pblock = NULL;
- gfc_omp_clauses clausesa[GFC_OMP_SPLIT_NUM];
+ stmtblock_t block, *new_pblock = pblock;
+ gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
tree stmt, omp_clauses = NULL_TREE;
- gfc_start_block (&block);
+ if (pblock == NULL)
+ gfc_start_block (&block);
+ else
+ gfc_init_block (&block);
- gfc_split_omp_clauses (code, clausesa);
+ if (clausesa == NULL)
+ {
+ clausesa = clausesa_buf;
+ gfc_split_omp_clauses (code, clausesa);
+ }
omp_clauses
= gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_PARALLEL],
code->loc);
- if (!clausesa[GFC_OMP_SPLIT_DO].ordered
- && clausesa[GFC_OMP_SPLIT_DO].sched_kind != OMP_SCHED_STATIC)
- pblock = █
- else
- pushlevel ();
- stmt = gfc_trans_omp_do (code, EXEC_OMP_DO, pblock,
+ if (pblock == NULL)
+ {
+ if (!clausesa[GFC_OMP_SPLIT_DO].ordered
+ && clausesa[GFC_OMP_SPLIT_DO].sched_kind != OMP_SCHED_STATIC)
+ new_pblock = █
+ else
+ pushlevel ();
+ }
+ stmt = gfc_trans_omp_do (code, EXEC_OMP_DO, new_pblock,
&clausesa[GFC_OMP_SPLIT_DO], omp_clauses);
- if (TREE_CODE (stmt) != BIND_EXPR)
- stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
- else
- poplevel (0, 0);
+ if (pblock == NULL)
+ {
+ if (TREE_CODE (stmt) != BIND_EXPR)
+ stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+ else
+ poplevel (0, 0);
+ }
+ else if (TREE_CODE (stmt) != BIND_EXPR)
+ stmt = build3_v (BIND_EXPR, NULL, stmt, NULL_TREE);
stmt = build2_loc (input_location, OMP_PARALLEL, void_type_node, stmt,
omp_clauses);
OMP_PARALLEL_COMBINED (stmt) = 1;
}
static tree
-gfc_trans_omp_parallel_do_simd (gfc_code *code)
+gfc_trans_omp_parallel_do_simd (gfc_code *code, stmtblock_t *pblock,
+ gfc_omp_clauses *clausesa)
{
stmtblock_t block;
- gfc_omp_clauses clausesa[GFC_OMP_SPLIT_NUM];
+ gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
tree stmt, omp_clauses = NULL_TREE;
- gfc_start_block (&block);
+ if (pblock == NULL)
+ gfc_start_block (&block);
+ else
+ gfc_init_block (&block);
- gfc_split_omp_clauses (code, clausesa);
+ if (clausesa == NULL)
+ {
+ clausesa = clausesa_buf;
+ gfc_split_omp_clauses (code, clausesa);
+ }
if (gfc_option.gfc_flag_openmp)
omp_clauses
= gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_PARALLEL],
code->loc);
- pushlevel ();
- stmt = gfc_trans_omp_do_simd (code, clausesa, omp_clauses);
- if (TREE_CODE (stmt) != BIND_EXPR)
- stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
- else
- poplevel (0, 0);
+ if (pblock == NULL)
+ pushlevel ();
+ stmt = gfc_trans_omp_do_simd (code, pblock, clausesa, omp_clauses);
+ if (pblock == NULL)
+ {
+ if (TREE_CODE (stmt) != BIND_EXPR)
+ stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+ else
+ poplevel (0, 0);
+ }
+ else if (TREE_CODE (stmt) != BIND_EXPR)
+ stmt = build3_v (BIND_EXPR, NULL, stmt, NULL_TREE);
if (gfc_option.gfc_flag_openmp)
{
stmt = build2_loc (input_location, OMP_PARALLEL, void_type_node, stmt,
}
static tree
+gfc_trans_omp_distribute (gfc_code *code, gfc_omp_clauses *clausesa)
+{
+ stmtblock_t block;
+ gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
+ tree stmt, omp_clauses = NULL_TREE;
+
+ gfc_start_block (&block);
+ if (clausesa == NULL)
+ {
+ clausesa = clausesa_buf;
+ gfc_split_omp_clauses (code, clausesa);
+ }
+ if (gfc_option.gfc_flag_openmp)
+ omp_clauses
+ = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_DISTRIBUTE],
+ code->loc);
+ switch (code->op)
+ {
+ case EXEC_OMP_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ /* This is handled in gfc_trans_omp_do. */
+ gcc_unreachable ();
+ break;
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ stmt = gfc_trans_omp_parallel_do (code, &block, clausesa);
+ if (TREE_CODE (stmt) != BIND_EXPR)
+ stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+ else
+ poplevel (0, 0);
+ break;
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ stmt = gfc_trans_omp_parallel_do_simd (code, &block, clausesa);
+ if (TREE_CODE (stmt) != BIND_EXPR)
+ stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+ else
+ poplevel (0, 0);
+ break;
+ case EXEC_OMP_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+ stmt = gfc_trans_omp_do (code, EXEC_OMP_SIMD, &block,
+ &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE);
+ if (TREE_CODE (stmt) != BIND_EXPR)
+ stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+ else
+ poplevel (0, 0);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ if (gfc_option.gfc_flag_openmp)
+ {
+ tree distribute = make_node (OMP_DISTRIBUTE);
+ TREE_TYPE (distribute) = void_type_node;
+ OMP_FOR_BODY (distribute) = stmt;
+ OMP_FOR_CLAUSES (distribute) = omp_clauses;
+ stmt = distribute;
+ }
+ gfc_add_expr_to_block (&block, stmt);
+ return gfc_finish_block (&block);
+}
+
+static tree
+gfc_trans_omp_teams (gfc_code *code, gfc_omp_clauses *clausesa)
+{
+ stmtblock_t block;
+ gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
+ tree stmt, omp_clauses = NULL_TREE;
+
+ gfc_start_block (&block);
+ if (clausesa == NULL)
+ {
+ clausesa = clausesa_buf;
+ gfc_split_omp_clauses (code, clausesa);
+ }
+ if (gfc_option.gfc_flag_openmp)
+ omp_clauses
+ = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_TEAMS],
+ code->loc);
+ switch (code->op)
+ {
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TEAMS:
+ stmt = gfc_trans_omp_code (code->block->next, true);
+ break;
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ stmt = gfc_trans_omp_do (code, EXEC_OMP_DISTRIBUTE, NULL,
+ &clausesa[GFC_OMP_SPLIT_DISTRIBUTE],
+ NULL);
+ break;
+ default:
+ stmt = gfc_trans_omp_distribute (code, clausesa);
+ break;
+ }
+ stmt = build2_loc (input_location, OMP_TEAMS, void_type_node, stmt,
+ omp_clauses);
+ gfc_add_expr_to_block (&block, stmt);
+ return gfc_finish_block (&block);
+}
+
+static tree
+gfc_trans_omp_target (gfc_code *code)
+{
+ stmtblock_t block;
+ gfc_omp_clauses clausesa[GFC_OMP_SPLIT_NUM];
+ tree stmt, omp_clauses = NULL_TREE;
+
+ gfc_start_block (&block);
+ gfc_split_omp_clauses (code, clausesa);
+ if (gfc_option.gfc_flag_openmp)
+ omp_clauses
+ = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_TARGET],
+ code->loc);
+ if (code->op == EXEC_OMP_TARGET)
+ stmt = gfc_trans_omp_code (code->block->next, true);
+ else
+ stmt = gfc_trans_omp_teams (code, clausesa);
+ if (TREE_CODE (stmt) != BIND_EXPR)
+ stmt = build3_v (BIND_EXPR, NULL, stmt, NULL_TREE);
+ if (gfc_option.gfc_flag_openmp)
+ stmt = build2_loc (input_location, OMP_TARGET, void_type_node, stmt,
+ omp_clauses);
+ gfc_add_expr_to_block (&block, stmt);
+ return gfc_finish_block (&block);
+}
+
+static tree
+gfc_trans_omp_target_data (gfc_code *code)
+{
+ stmtblock_t block;
+ tree stmt, omp_clauses;
+
+ gfc_start_block (&block);
+ omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+ code->loc);
+ stmt = gfc_trans_omp_code (code->block->next, true);
+ stmt = build2_loc (input_location, OMP_TARGET_DATA, void_type_node, stmt,
+ omp_clauses);
+ gfc_add_expr_to_block (&block, stmt);
+ return gfc_finish_block (&block);
+}
+
+static tree
+gfc_trans_omp_target_update (gfc_code *code)
+{
+ stmtblock_t block;
+ tree stmt, omp_clauses;
+
+ gfc_start_block (&block);
+ omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+ code->loc);
+ stmt = build1_loc (input_location, OMP_TARGET_UPDATE, void_type_node,
+ omp_clauses);
+ gfc_add_expr_to_block (&block, stmt);
+ return gfc_finish_block (&block);
+}
+
+static tree
gfc_trans_omp_workshare (gfc_code *code, gfc_omp_clauses *clauses)
{
tree res, tmp, stmt;
return gfc_trans_omp_cancellation_point (code);
case EXEC_OMP_CRITICAL:
return gfc_trans_omp_critical (code);
+ case EXEC_OMP_DISTRIBUTE:
case EXEC_OMP_DO:
case EXEC_OMP_SIMD:
return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses,
NULL);
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_DISTRIBUTE_SIMD:
+ return gfc_trans_omp_distribute (code, NULL);
case EXEC_OMP_DO_SIMD:
- return gfc_trans_omp_do_simd (code, NULL, NULL_TREE);
+ return gfc_trans_omp_do_simd (code, NULL, NULL, NULL_TREE);
case EXEC_OMP_FLUSH:
return gfc_trans_omp_flush ();
case EXEC_OMP_MASTER:
case EXEC_OMP_PARALLEL:
return gfc_trans_omp_parallel (code);
case EXEC_OMP_PARALLEL_DO:
- return gfc_trans_omp_parallel_do (code);
+ return gfc_trans_omp_parallel_do (code, NULL, NULL);
case EXEC_OMP_PARALLEL_DO_SIMD:
- return gfc_trans_omp_parallel_do_simd (code);
+ return gfc_trans_omp_parallel_do_simd (code, NULL, NULL);
case EXEC_OMP_PARALLEL_SECTIONS:
return gfc_trans_omp_parallel_sections (code);
case EXEC_OMP_PARALLEL_WORKSHARE:
return gfc_trans_omp_sections (code, code->ext.omp_clauses);
case EXEC_OMP_SINGLE:
return gfc_trans_omp_single (code, code->ext.omp_clauses);
+ case EXEC_OMP_TARGET:
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ return gfc_trans_omp_target (code);
+ case EXEC_OMP_TARGET_DATA:
+ return gfc_trans_omp_target_data (code);
+ case EXEC_OMP_TARGET_UPDATE:
+ return gfc_trans_omp_target_update (code);
case EXEC_OMP_TASK:
return gfc_trans_omp_task (code);
case EXEC_OMP_TASKGROUP:
return gfc_trans_omp_taskwait ();
case EXEC_OMP_TASKYIELD:
return gfc_trans_omp_taskyield ();
+ case EXEC_OMP_TEAMS:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+ return gfc_trans_omp_teams (code, NULL);
case EXEC_OMP_WORKSHARE:
return gfc_trans_omp_workshare (code, code->ext.omp_clauses);
default:
case EXEC_OMP_CANCEL:
case EXEC_OMP_CANCELLATION_POINT:
case EXEC_OMP_CRITICAL:
+ case EXEC_OMP_DISTRIBUTE:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_DISTRIBUTE_SIMD:
case EXEC_OMP_DO:
case EXEC_OMP_DO_SIMD:
case EXEC_OMP_FLUSH:
case EXEC_OMP_SECTIONS:
case EXEC_OMP_SIMD:
case EXEC_OMP_SINGLE:
+ case EXEC_OMP_TARGET:
+ case EXEC_OMP_TARGET_DATA:
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TARGET_UPDATE:
case EXEC_OMP_TASK:
case EXEC_OMP_TASKGROUP:
case EXEC_OMP_TASKWAIT:
case EXEC_OMP_TASKYIELD:
+ case EXEC_OMP_TEAMS:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
case EXEC_OMP_WORKSHARE:
res = gfc_trans_omp_directive (code);
break;
tree gfc_omp_clause_copy_ctor (tree, tree, tree);
tree gfc_omp_clause_assign_op (tree, tree, tree);
tree gfc_omp_clause_dtor (tree, tree);
+void gfc_omp_finish_clause (tree, gimple_seq *);
bool gfc_omp_disregard_value_expr (tree, bool);
bool gfc_omp_private_debug_clause (tree, bool);
bool gfc_omp_private_outer_ref (tree);
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
if (ctx->region_type == ORT_TARGET)
{
+ ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
if (n == NULL)
{
if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
omp_add_variable (ctx, decl, GOVD_MAP | flags);
}
else
- n->value |= flags;
- ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
+ {
+ /* If nothing changed, there's nothing left to do. */
+ if ((n->value & flags) == flags)
+ return ret;
+ n->value |= flags;
+ }
goto do_outer;
}
gimplify_omp_ctxp = ctx;
}
+struct gimplify_adjust_omp_clauses_data
+{
+ tree *list_p;
+ gimple_seq *pre_p;
+};
+
/* For all variables that were not actually used within the context,
remove PRIVATE, SHARED, and FIRSTPRIVATE clauses. */
static int
gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
{
- tree *list_p = (tree *) data;
+ tree *list_p = ((struct gimplify_adjust_omp_clauses_data *) data)->list_p;
+ gimple_seq *pre_p
+ = ((struct gimplify_adjust_omp_clauses_data *) data)->pre_p;
tree decl = (tree) n->key;
unsigned flags = n->value;
enum omp_clause_code code;
OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (nc) = 1;
OMP_CLAUSE_CHAIN (nc) = *list_p;
OMP_CLAUSE_CHAIN (clause) = nc;
- lang_hooks.decls.omp_finish_clause (nc);
+ struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = ctx->outer_context;
+ lang_hooks.decls.omp_finish_clause (nc, pre_p);
+ gimplify_omp_ctxp = ctx;
}
*list_p = clause;
- lang_hooks.decls.omp_finish_clause (clause);
+ struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = ctx->outer_context;
+ lang_hooks.decls.omp_finish_clause (clause, pre_p);
+ gimplify_omp_ctxp = ctx;
return 0;
}
static void
-gimplify_adjust_omp_clauses (tree *list_p)
+gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
{
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
tree c, decl;
}
/* Add in any implicit data sharing. */
- splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, list_p);
+ struct gimplify_adjust_omp_clauses_data data;
+ data.list_p = list_p;
+ data.pre_p = pre_p;
+ splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
gimplify_omp_ctxp = ctx->outer_context;
delete_omp_context (ctx);
else
pop_gimplify_context (NULL);
- gimplify_adjust_omp_clauses (&OMP_PARALLEL_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_PARALLEL_CLAUSES (expr));
g = gimple_build_omp_parallel (body,
OMP_PARALLEL_CLAUSES (expr),
else
pop_gimplify_context (NULL);
- gimplify_adjust_omp_clauses (&OMP_TASK_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_TASK_CLAUSES (expr));
g = gimple_build_omp_task (body,
OMP_TASK_CLAUSES (expr),
TREE_OPERAND (TREE_OPERAND (t, 1), 0) = var;
}
- gimplify_adjust_omp_clauses (&OMP_FOR_CLAUSES (orig_for_stmt));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_FOR_CLAUSES (orig_for_stmt));
int kind;
switch (TREE_CODE (orig_for_stmt))
}
else
gimplify_and_add (OMP_BODY (expr), &body);
- gimplify_adjust_omp_clauses (&OMP_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_CLAUSES (expr));
switch (TREE_CODE (expr))
{
gimplify_scan_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr), pre_p,
ORT_WORKSHARE);
- gimplify_adjust_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_TARGET_UPDATE_CLAUSES (expr));
stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_UPDATE,
OMP_TARGET_UPDATE_CLAUSES (expr));
extern int lhd_gimplify_expr (tree *, gimple_seq *, gimple_seq *);
extern enum omp_clause_default_kind lhd_omp_predetermined_sharing (tree);
extern tree lhd_omp_assignment (tree, tree, tree);
+extern void lhd_omp_finish_clause (tree, gimple_seq *);
struct gimplify_omp_ctx;
extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
tree);
#define LANG_HOOKS_OMP_CLAUSE_COPY_CTOR lhd_omp_assignment
#define LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP lhd_omp_assignment
#define LANG_HOOKS_OMP_CLAUSE_DTOR hook_tree_tree_tree_null
-#define LANG_HOOKS_OMP_FINISH_CLAUSE hook_void_tree
+#define LANG_HOOKS_OMP_FINISH_CLAUSE lhd_omp_finish_clause
#define LANG_HOOKS_DECLS { \
LANG_HOOKS_GLOBAL_BINDINGS_P, \
return build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, src);
}
+/* Finalize clause C. */
+
+void
+lhd_omp_finish_clause (tree, gimple_seq *)
+{
+}
+
/* Register language specific type size variables as potentially OpenMP
firstprivate variables. */
tree (*omp_clause_dtor) (tree clause, tree decl);
/* Do language specific checking on an implicitly determined clause. */
- void (*omp_finish_clause) (tree clause);
+ void (*omp_finish_clause) (tree clause, gimple_seq *pre_p);
};
/* Language hooks related to LTO serialization. */
}
else
{
+ if (ctx->outer)
+ {
+ scan_omp_op (&OMP_CLAUSE_DECL (c), ctx->outer);
+ decl = OMP_CLAUSE_DECL (c);
+ }
gcc_assert (!splay_tree_lookup (ctx->field_map,
(splay_tree_key) decl));
tree field
tree temp = create_tmp_var (type, NULL);
tree c = build_omp_clause (UNKNOWN_LOCATION,
OMP_CLAUSE__LOOPTEMP_);
+ insert_decl_map (&outer_ctx->cb, temp, temp);
OMP_CLAUSE_DECL (c) = temp;
OMP_CLAUSE_CHAIN (c) = gimple_omp_parallel_clauses (stmt);
gimple_omp_parallel_set_clauses (stmt, c);
return false;
}
break;
+ case GIMPLE_OMP_TARGET:
+ for (; ctx != NULL; ctx = ctx->outer)
+ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+ && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION)
+ {
+ const char *name;
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION: name = "target"; break;
+ case GF_OMP_TARGET_KIND_DATA: name = "target data"; break;
+ case GF_OMP_TARGET_KIND_UPDATE: name = "target update"; break;
+ default: gcc_unreachable ();
+ }
+ warning_at (gimple_location (stmt), 0,
+ "%s construct inside of target region", name);
+ }
+ break;
default:
break;
}
OMP_CLAUSE__LOOPTEMP_);
}
else
- temp = create_tmp_var (type, NULL);
+ {
+ temp = create_tmp_var (type, NULL);
+ insert_decl_map (&ctx->outer->cb, temp, temp);
+ }
*pc = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE__LOOPTEMP_);
OMP_CLAUSE_DECL (*pc) = temp;
pc = &OMP_CLAUSE_CHAIN (*pc);
+2014-06-18 Jakub Jelinek <jakub@redhat.com>
+
+ * gfortran.dg/gomp/declare-simd-1.f90: New test.
+ * gfortran.dg/gomp/depend-1.f90: New test.
+ * gfortran.dg/gomp/target1.f90: New test.
+ * gfortran.dg/gomp/target2.f90: New test.
+ * gfortran.dg/gomp/target3.f90: New test.
+ * gfortran.dg/gomp/udr4.f90: Adjust expected diagnostics.
+ * gfortran.dg/openmp-define-3.f90: Expect _OPENMP 201307 instead of
+ 201107.
+
2014-06-18 Dominique d'Humieres <dominiq@lps.ens.fr>
PR fortran/61126
--- /dev/null
+! { dg-do compile }
+
+subroutine fn1 (x)
+ integer :: x
+!$omp declare simd (fn1) inbranch notinbranch uniform (x) ! { dg-error "Unclassifiable OpenMP directive" }
+end subroutine fn1
+subroutine fn2 (x)
+!$omp declare simd (fn100) ! { dg-error "should refer to containing procedure" }
+end subroutine fn2
--- /dev/null
+! { dg-do compile }
+
+subroutine foo (x)
+ integer :: x(5, *)
+!$omp parallel
+!$omp single
+!$omp task depend(in:x(:,5))
+!$omp end task
+!$omp task depend(in:x(5,:)) ! { dg-error "Rightmost upper bound of assumed size array section|proper array section" }
+!$omp end task
+!$omp end single
+!$omp end parallel
+end
--- /dev/null
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+module target1
+ interface
+ subroutine dosomething (a, n, m)
+ integer :: a (:), n, m
+ !$omp declare target
+ end subroutine dosomething
+ end interface
+contains
+ subroutine foo (n, o, p, q, r, pp)
+ integer :: n, o, p, q, r, s, i, j
+ integer :: a (2:o)
+ integer, pointer :: pp
+ !$omp target data device (n + 1) if (n .ne. 6) map (tofrom: n, r)
+ !$omp target device (n + 1) if (n .ne. 6) map (from: n) map (alloc: a(2:o))
+ call dosomething (a, n, 0)
+ !$omp end target
+ !$omp target teams device (n + 1) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r)
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ !$omp end target teams
+ !$omp target teams distribute device (n + 1) num_teams (n + 4) collapse (2) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ end do
+ !$omp target teams distribute device (n + 1) num_teams (n + 4) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ end do
+ !$omp end target teams distribute
+ !$omp target teams distribute parallel do device (n + 1) num_teams (n + 4) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+ !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+ !$omp & ordered schedule (static, 8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ !$omp ordered
+ p = q
+ !$omp end ordered
+ s = i * 10 + j
+ end do
+ end do
+ !$omp target teams distribute parallel do device (n + 1) num_teams (n + 4) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) num_threads (n + 4) &
+ !$omp & proc_bind (master) lastprivate (s) ordered schedule (static, 8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ !$omp ordered
+ p = q
+ !$omp end ordered
+ s = i * 10
+ end do
+ !$omp end target teams distribute parallel do
+ !$omp target teams distribute parallel do simd device (n + 1) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+ !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+ !$omp & schedule (static, 8) num_teams (n + 4) safelen(8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ a(2+i*10+j) = p + q
+ s = i * 10 + j
+ end do
+ end do
+ !$omp target teams distribute parallel do simd device (n + 1) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) num_threads (n + 4) &
+ !$omp & proc_bind (master) lastprivate (s) schedule (static, 8) &
+ !$omp & num_teams (n + 4) safelen(16) linear(i:1) aligned (pp:4)
+ do i = 1, 10
+ r = r + 1
+ p = q
+ a(1+i) = p + q
+ s = i * 10
+ end do
+ !$omp end target teams distribute parallel do simd
+ !$omp target teams distribute simd device (n + 1) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+ !$omp & lastprivate (s) num_teams (n + 4) safelen(8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ a(2+i*10+j) = p + q
+ s = i * 10 + j
+ end do
+ end do
+ !$omp target teams distribute simd device (n + 1) &
+ !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) lastprivate (s) &
+ !$omp & num_teams (n + 4) safelen(16) linear(i:1) aligned (pp:4)
+ do i = 1, 10
+ r = r + 1
+ p = q
+ a(1+i) = p + q
+ s = i * 10
+ end do
+ !$omp end target teams distribute simd
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams num_teams (n + 4) thread_limit (n * 2) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r)
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ !$omp end teams
+ !$omp end target
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams distribute num_teams (n + 4) collapse (2) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ end do
+ !$omp end target
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams distribute num_teams (n + 4) default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ end do
+ !$omp end teams distribute
+ !$omp end target
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams distribute parallel do num_teams (n + 4) &
+ !$omp & if (n .ne. 6) default(shared) ordered schedule (static, 8) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+ !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ !$omp ordered
+ p = q
+ !$omp end ordered
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams distribute parallel do num_teams (n + 4)if(n.ne.6)default(shared)&
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) num_threads (n + 4) &
+ !$omp & proc_bind (master) lastprivate (s) ordered schedule (static, 8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ !$omp ordered
+ p = q
+ !$omp end ordered
+ s = i * 10
+ end do
+ !$omp end teams distribute parallel do
+ !$omp end target
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams distribute parallel do simd if(n.ne.6)default(shared)&
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+ !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+ !$omp & schedule (static, 8) num_teams (n + 4) safelen(8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ a(2+i*10+j) = p + q
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams distribute parallel do simd if (n .ne. 6)default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) num_threads (n + 4) &
+ !$omp & proc_bind (master) lastprivate (s) schedule (static, 8) &
+ !$omp & num_teams (n + 4) safelen(16) linear(i:1) aligned (pp:4)
+ do i = 1, 10
+ r = r + 1
+ p = q
+ a(1+i) = p + q
+ s = i * 10
+ end do
+ !$omp end teams distribute parallel do simd
+ !$omp end target
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams distribute simd default(shared) safelen(8) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+ !$omp & lastprivate (s) num_teams (n + 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ a(2+i*10+j) = p + q
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target
+ !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+ !$omp teams distribute simd default(shared) aligned (pp:4) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & thread_limit (n * 2) dist_schedule (static, 4) lastprivate (s)
+ do i = 1, 10
+ r = r + 1
+ p = q
+ a(1+i) = p + q
+ s = i * 10
+ end do
+ !$omp end teams distribute simd
+ !$omp end target
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction ( + : r )
+ !$omp distribute collapse (2) firstprivate (q) dist_schedule (static, 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ end do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute firstprivate (q) dist_schedule (static, 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ end do
+ !$omp end distribute
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute parallel do if (n .ne. 6) default(shared) &
+ !$omp & ordered schedule (static, 8) private (p) firstprivate (q) &
+ !$omp & shared(n)reduction(+:r)dist_schedule(static,4)collapse(2)&
+ !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ !$omp ordered
+ p = q
+ !$omp end ordered
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute parallel do if(n.ne.6)default(shared)&
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) num_threads (n + 4) &
+ !$omp & proc_bind (master) lastprivate (s) ordered schedule (static, 8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ !$omp ordered
+ p = q
+ !$omp end ordered
+ s = i * 10
+ end do
+ !$omp end distribute parallel do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute parallel do simd if(n.ne.6)default(shared)&
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) collapse (2) safelen(8) &
+ !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+ !$omp & schedule (static, 8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ a(2+i*10+j) = p + q
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute parallel do simd if (n .ne. 6)default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) num_threads (n + 4) &
+ !$omp & proc_bind (master) lastprivate (s) schedule (static, 8) &
+ !$omp & safelen(16) linear(i:1) aligned (pp:4)
+ do i = 1, 10
+ r = r + 1
+ p = q
+ a(1+i) = p + q
+ s = i * 10
+ end do
+ !$omp end distribute parallel do simd
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute simd safelen(8) lastprivate(s) &
+ !$omp & private (p) firstprivate (q) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) collapse (2)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ a(2+i*10+j) = p + q
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute simd aligned (pp:4) &
+ !$omp & private (p) firstprivate (q) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) lastprivate (s)
+ do i = 1, 10
+ r = r + 1
+ p = q
+ a(1+i) = p + q
+ s = i * 10
+ end do
+ !$omp end distribute simd
+ !$omp end target teams
+ !$omp end target data
+ end subroutine
+ subroutine bar (n, o, p, r, pp)
+ integer :: n, o, p, q, r, s, i, j
+ integer :: a (2:o)
+ integer, pointer :: pp
+ common /blk/ i, j, q
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction ( + : r )
+ !$omp distribute collapse (2) firstprivate (q) dist_schedule (static, 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ end do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute firstprivate (q) dist_schedule (static, 4)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ end do
+ !$omp end distribute
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute parallel do if (n .ne. 6) default(shared) &
+ !$omp & ordered schedule (static, 8) private (p) firstprivate (q) &
+ !$omp & shared(n)reduction(+:r)dist_schedule(static,4)collapse(2)&
+ !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ !$omp ordered
+ p = q
+ !$omp end ordered
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute parallel do if(n.ne.6)default(shared)&
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) num_threads (n + 4) &
+ !$omp & proc_bind (master) lastprivate (s) ordered schedule (static, 8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ call dosomething (a, n, p + q)
+ end do
+ !$omp ordered
+ p = q
+ !$omp end ordered
+ s = i * 10
+ end do
+ !$omp end distribute parallel do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute parallel do simd if(n.ne.6)default(shared)&
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) collapse (2) safelen(8) &
+ !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+ !$omp & schedule (static, 8)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ a(2+i*10+j) = p + q
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute parallel do simd if (n .ne. 6)default(shared) &
+ !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) num_threads (n + 4) &
+ !$omp & proc_bind (master) lastprivate (s) schedule (static, 8) &
+ !$omp & safelen(16) linear(i:1) aligned (pp:4)
+ do i = 1, 10
+ r = r + 1
+ p = q
+ a(1+i) = p + q
+ s = i * 10
+ end do
+ !$omp end distribute parallel do simd
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute simd safelen(8) lastprivate(s) &
+ !$omp & private (p) firstprivate (q) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) collapse (2)
+ do i = 1, 10
+ do j = 1, 10
+ r = r + 1
+ p = q
+ a(2+i*10+j) = p + q
+ s = i * 10 + j
+ end do
+ end do
+ !$omp end target teams
+ !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+ !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+ !$omp & default(shared) shared(n) private (p) reduction(+:r)
+ !$omp distribute simd aligned (pp:4) &
+ !$omp & private (p) firstprivate (q) reduction (+: r) &
+ !$omp & dist_schedule (static, 4) lastprivate (s)
+ do i = 1, 10
+ r = r + 1
+ p = q
+ a(1+i) = p + q
+ s = i * 10
+ end do
+ !$omp end distribute simd
+ !$omp end target teams
+ end subroutine
+end module
--- /dev/null
+! { dg-do compile }
+! { dg-options "-fopenmp -ffree-line-length-160" }
+
+subroutine foo (n, s, t, u, v, w)
+ integer :: n, i, s, t, u, v, w
+ common /bar/ i
+ !$omp simd safelen(s + 1)
+ do i = 1, n
+ end do
+ !$omp do schedule (static, t * 2)
+ do i = 1, n
+ end do
+ !$omp do simd safelen(s + 1) schedule (static, t * 2)
+ do i = 1, n
+ end do
+ !$omp parallel do schedule (static, t * 2) num_threads (u - 1)
+ do i = 1, n
+ end do
+ !$omp parallel do simd safelen(s + 1) schedule (static, t * 2) num_threads (u - 1)
+ do i = 1, n
+ end do
+ !$omp distribute dist_schedule (static, v + 8)
+ do i = 1, n
+ end do
+ !$omp distribute simd dist_schedule (static, v + 8) safelen(s + 1)
+ do i = 1, n
+ end do
+ !$omp distribute parallel do simd dist_schedule (static, v + 8) safelen(s + 1) &
+ !$omp & schedule (static, t * 2) num_threads (u - 1)
+ do i = 1, n
+ end do
+ !$omp distribute parallel do dist_schedule (static, v + 8) num_threads (u - 1) &
+ !$omp & schedule (static, t * 2)
+ do i = 1, n
+ end do
+ !$omp target
+ !$omp teams distribute dist_schedule (static, v + 8) num_teams (w + 8)
+ do i = 1, n
+ end do
+ !$omp end target
+ !$omp target
+ !$omp teams distribute simd dist_schedule (static, v + 8) safelen(s + 1) &
+ !$omp & num_teams (w + 8)
+ do i = 1, n
+ end do
+ !$omp end target
+ !$omp target
+ !$omp teams distribute parallel do simd dist_schedule (static, v + 8) safelen(s + 1) &
+ !$omp & schedule (static, t * 2) num_threads (u - 1) num_teams (w + 8)
+ do i = 1, n
+ end do
+ !$omp end target
+ !$omp target
+ !$omp teams distribute parallel do dist_schedule (static, v + 8) num_threads (u - 1) &
+ !$omp & schedule (static, t * 2) num_teams (w + 8)
+ do i = 1, n
+ end do
+ !$omp end target
+ !$omp target teams distribute dist_schedule (static, v + 8) num_teams (w + 8)
+ do i = 1, n
+ end do
+ !$omp target teams distribute simd dist_schedule (static, v + 8) safelen(s + 1) &
+ !$omp & num_teams (w + 8)
+ do i = 1, n
+ end do
+ !$omp target teams distribute parallel do simd dist_schedule (static, v + 8) safelen(s + 1) &
+ !$omp & schedule (static, t * 2) num_threads (u - 1) num_teams (w + 8)
+ do i = 1, n
+ end do
+ !$omp target teams distribute parallel do dist_schedule (static, v + 8) num_threads (u - 1) &
+ !$omp & schedule (static, t * 2) num_teams (w + 8)
+ do i = 1, n
+ end do
+end subroutine
--- /dev/null
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+subroutine foo (r)
+ integer :: i, r
+ !$omp target
+ !$omp target teams distribute parallel do reduction (+: r) ! { dg-warning "target construct inside of target region" }
+ do i = 1, 10
+ r = r + 1
+ end do
+ !$omp end target
+end subroutine
!$omp declare reduction (foo) ! { dg-error "Unclassifiable OpenMP directive" }
!$omp declare reduction (foo:integer) ! { dg-error "Unclassifiable OpenMP directive" }
!$omp declare reduction (foo:integer:omp_out=omp_out+omp_in) &
-!$omp & initializer(omp_priv=0) initializer(omp_priv=0) ! { dg-error "Unclassifiable statement" }
+!$omp & initializer(omp_priv=0) initializer(omp_priv=0) ! { dg-error "Unexpected junk after" }
end subroutine f3
subroutine f4
implicit integer (o)
# error _OPENMP not defined
#endif
-#if _OPENMP != 201107
+#if _OPENMP != 201307
# error _OPENMP defined to wrong value
#endif
array sections. OMP_CLAUSE_SIZE for these is not the pointer size,
which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias. */
OMP_CLAUSE_MAP_POINTER,
+ /* Also internal, behaves like OMP_CLAUS_MAP_TO, but additionally any
+ OMP_CLAUSE_MAP_POINTER records consecutive after it which have addresses
+ falling into that range will not be ignored if OMP_CLAUSE_MAP_TO_PSET
+ wasn't mapped already. */
+ OMP_CLAUSE_MAP_TO_PSET,
OMP_CLAUSE_MAP_LAST
};
case OMP_CLAUSE_LINEAR:
if (OMP_CLAUSE_LINEAR_GIMPLE_SEQ (clause))
need_stmts = true;
+ wi->val_only = true;
+ wi->is_lhs = false;
+ convert_nonlocal_reference_op (&OMP_CLAUSE_LINEAR_STEP (clause),
+ &dummy, wi);
goto do_decl_clause;
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_DEPEND:
+ case OMP_CLAUSE_DEVICE:
+ case OMP_CLAUSE_NUM_TEAMS:
+ case OMP_CLAUSE_THREAD_LIMIT:
+ case OMP_CLAUSE_SAFELEN:
wi->val_only = true;
wi->is_lhs = false;
convert_nonlocal_reference_op (&OMP_CLAUSE_OPERAND (clause, 0),
- &dummy, wi);
+ &dummy, wi);
+ break;
+
+ case OMP_CLAUSE_DIST_SCHEDULE:
+ if (OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (clause) != NULL)
+ {
+ wi->val_only = true;
+ wi->is_lhs = false;
+ convert_nonlocal_reference_op (&OMP_CLAUSE_OPERAND (clause, 0),
+ &dummy, wi);
+ }
+ break;
+
+ case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ if (OMP_CLAUSE_SIZE (clause))
+ {
+ wi->val_only = true;
+ wi->is_lhs = false;
+ convert_nonlocal_reference_op (&OMP_CLAUSE_SIZE (clause),
+ &dummy, wi);
+ }
+ if (DECL_P (OMP_CLAUSE_DECL (clause)))
+ goto do_decl_clause;
+ wi->val_only = true;
+ wi->is_lhs = false;
+ convert_nonlocal_reference_op (&OMP_CLAUSE_DECL (clause),
+ &dummy, wi);
break;
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_MERGEABLE:
+ case OMP_CLAUSE_PROC_BIND:
break;
default:
case OMP_CLAUSE_LINEAR:
if (OMP_CLAUSE_LINEAR_GIMPLE_SEQ (clause))
need_stmts = true;
+ wi->val_only = true;
+ wi->is_lhs = false;
+ convert_local_reference_op (&OMP_CLAUSE_LINEAR_STEP (clause), &dummy,
+ wi);
goto do_decl_clause;
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_DEPEND:
+ case OMP_CLAUSE_DEVICE:
+ case OMP_CLAUSE_NUM_TEAMS:
+ case OMP_CLAUSE_THREAD_LIMIT:
+ case OMP_CLAUSE_SAFELEN:
wi->val_only = true;
wi->is_lhs = false;
convert_local_reference_op (&OMP_CLAUSE_OPERAND (clause, 0), &dummy,
wi);
break;
+ case OMP_CLAUSE_DIST_SCHEDULE:
+ if (OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (clause) != NULL)
+ {
+ wi->val_only = true;
+ wi->is_lhs = false;
+ convert_local_reference_op (&OMP_CLAUSE_OPERAND (clause, 0),
+ &dummy, wi);
+ }
+ break;
+
+ case OMP_CLAUSE_MAP:
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ if (OMP_CLAUSE_SIZE (clause))
+ {
+ wi->val_only = true;
+ wi->is_lhs = false;
+ convert_local_reference_op (&OMP_CLAUSE_SIZE (clause),
+ &dummy, wi);
+ }
+ if (DECL_P (OMP_CLAUSE_DECL (clause)))
+ goto do_decl_clause;
+ wi->val_only = true;
+ wi->is_lhs = false;
+ convert_local_reference_op (&OMP_CLAUSE_DECL (clause),
+ &dummy, wi);
+ break;
+
+
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_DEFAULT:
case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_MERGEABLE:
+ case OMP_CLAUSE_PROC_BIND:
break;
default:
pp_string (buffer, "alloc");
break;
case OMP_CLAUSE_MAP_TO:
+ case OMP_CLAUSE_MAP_TO_PSET:
pp_string (buffer, "to");
break;
case OMP_CLAUSE_MAP_FROM:
if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_POINTER)
pp_string (buffer, " [pointer assign, bias: ");
+ else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_TO_PSET)
+ pp_string (buffer, " [pointer set, len: ");
else
pp_string (buffer, " [len: ");
dump_generic_node (buffer, OMP_CLAUSE_SIZE (clause),
+2014-06-18 Jakub Jelinek <jakub@redhat.com>
+
+ * omp_lib.f90.in (openmp_version): Set to 201307.
+ * omp_lib.h.in (openmp_version): Likewise.
+ * testsuite/libgomp.c/target-8.c: New test.
+ * testsuite/libgomp.fortran/declare-simd-1.f90: Add notinbranch
+ and inbranch clauses.
+ * testsuite/libgomp.fortran/depend-3.f90: New test.
+ * testsuite/libgomp.fortran/openmp_version-1.f: Adjust for new
+ openmp_version.
+ * testsuite/libgomp.fortran/openmp_version-2.f90: Likewise.
+ * testsuite/libgomp.fortran/target1.f90: New test.
+ * testsuite/libgomp.fortran/target2.f90: New test.
+ * testsuite/libgomp.fortran/target3.f90: New test.
+ * testsuite/libgomp.fortran/target4.f90: New test.
+ * testsuite/libgomp.fortran/target5.f90: New test.
+ * testsuite/libgomp.fortran/target6.f90: New test.
+ * testsuite/libgomp.fortran/target7.f90: New test.
+
2014-06-10 Jakub Jelinek <jakub@redhat.com>
PR fortran/60928
module omp_lib
use omp_lib_kinds
implicit none
- integer, parameter :: openmp_version = 201107
+ integer, parameter :: openmp_version = 201307
interface
subroutine omp_init_lock (svar)
parameter (omp_proc_bind_master = 2)
parameter (omp_proc_bind_close = 3)
parameter (omp_proc_bind_spread = 4)
- parameter (openmp_version = 201107)
+ parameter (openmp_version = 201307)
external omp_init_lock, omp_init_nest_lock
external omp_destroy_lock, omp_destroy_nest_lock
--- /dev/null
+/* { dg-do run } */
+/* { dg-options "-fopenmp" } */
+
+void
+foo (int *p)
+{
+ int i;
+ #pragma omp parallel
+ #pragma omp single
+ #pragma omp target teams distribute parallel for map(p[0:24])
+ for (i = 0; i < 24; i++)
+ p[i] = p[i] + 1;
+}
+
+int
+main ()
+{
+ int p[24], i;
+ for (i = 0; i < 24; i++)
+ p[i] = i;
+ foo (p);
+ for (i = 0; i < 24; i++)
+ if (p[i] != i + 1)
+ __builtin_abort ();
+ return 0;
+}
module declare_simd_1_mod
contains
real function foo (a, b, c)
- !$omp declare simd (foo) simdlen (4) uniform (a) linear (b : 5)
+ !$omp declare simd (foo) simdlen (4) uniform (a) linear (b : 5) &
+ !$omp & notinbranch
double precision, value :: a
real, value :: c
!$omp declare simd (foo)
real, value :: c
real :: bar
!$omp declare simd (bar) simdlen (4) linear (b : 2)
+ !$omp declare simd (bar) simdlen (16) inbranch
double precision, value :: a
end function bar
end interface
--- /dev/null
+! { dg-do run }
+
+ integer :: x(2, 3)
+ integer, allocatable :: z(:, :)
+ allocate (z(-2:3, 2:4))
+ call foo (x, z)
+contains
+ subroutine foo (x, z)
+ integer :: x(:, :), y
+ integer, allocatable :: z(:, :)
+ y = 1
+ !$omp parallel shared (x, y, z)
+ !$omp single
+ !$omp taskgroup
+ !$omp task depend(in: x)
+ if (y.ne.1) call abort
+ !$omp end task
+ !$omp task depend(out: x(1:2, 1:3))
+ y = 2
+ !$omp end task
+ !$omp end taskgroup
+ !$omp taskgroup
+ !$omp task depend(in: z)
+ if (y.ne.2) call abort
+ !$omp end task
+ !$omp task depend(out: z(-2:3, 2:4))
+ y = 3
+ !$omp end task
+ !$omp end taskgroup
+ !$omp taskgroup
+ !$omp task depend(in: x)
+ if (y.ne.3) call abort
+ !$omp end task
+ !$omp task depend(out: x(1:, 1:))
+ y = 4
+ !$omp end task
+ !$omp end taskgroup
+ !$omp end single
+ !$omp end parallel
+ if (y.ne.4) call abort
+ end subroutine
+end
implicit none
include "omp_lib.h"
- if (openmp_version .ne. 201107) call abort;
+ if (openmp_version .ne. 201307) call abort;
end program main
use omp_lib
implicit none
- if (openmp_version .ne. 201107) call abort;
+ if (openmp_version .ne. 201307) call abort;
end program main
--- /dev/null
+! { dg-do run }
+
+module target1
+contains
+ subroutine foo (p, v, w, n)
+ double precision, pointer :: p(:), v(:), w(:)
+ double precision :: q(n)
+ integer :: i, n
+ !$omp target if (n > 256) map (to: v(1:n), w(:n)) map (from: p(1:n), q)
+ !$omp parallel do simd
+ do i = 1, n
+ p(i) = v(i) * w(i)
+ q(i) = p(i)
+ end do
+ !$omp end target
+ if (any (p /= q)) call abort
+ do i = 1, n
+ if (p(i) /= i * iand (i, 63)) call abort
+ end do
+ !$omp target data if (n > 256) map (to: v(1:n), w) map (from: p, q)
+ !$omp target if (n > 256)
+ do i = 1, n
+ p(i) = 1.0
+ q(i) = 2.0
+ end do
+ !$omp end target
+ !$omp target if (n > 256)
+ do i = 1, n
+ p(i) = p(i) + v(i) * w(i)
+ q(i) = q(i) + v(i) * w(i)
+ end do
+ !$omp end target
+ !$omp target if (n > 256)
+ !$omp teams distribute parallel do simd linear(i:1)
+ do i = 1, n
+ p(i) = p(i) + 2.0
+ q(i) = q(i) + 3.0
+ end do
+ !$omp end target
+ !$omp end target data
+ if (any (p + 2.0 /= q)) call abort
+ end subroutine
+end module target1
+ use target1, only : foo
+ integer :: n, i
+ double precision, pointer :: p(:), v(:), w(:)
+ n = 10000
+ allocate (p(n), v(n), w(n))
+ do i = 1, n
+ v(i) = i
+ w(i) = iand (i, 63)
+ end do
+ call foo (p, v, w, n)
+ do i = 1, n
+ if (p(i) /= i * iand (i, 63) + 3) call abort
+ end do
+ deallocate (p, v, w)
+end
--- /dev/null
+! { dg-do run }
+! { dg-options "-fopenmp -ffree-line-length-160" }
+
+module target2
+contains
+ subroutine foo (a, b, c, d, e, f, g, n, q)
+ integer :: n, q
+ integer :: a, b(3:n), c(5:), d(2:*), e(:,:)
+ integer, pointer :: f, g(:)
+ integer :: h, i(3:n)
+ integer, pointer :: j, k(:)
+ logical :: r
+ allocate (j, k(4:n))
+ h = 14
+ i = 15
+ j = 16
+ k = 17
+ !$omp target map (to: a, b, c, d(2:n+1), e, f, g, h, i, j, k, n) map (from: r)
+ r = a /= 7
+ r = r .or. (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+ r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+ r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+ r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+ r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+ r = r .or. (f /= 12)
+ r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+ r = r .or. (h /= 14)
+ r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+ r = r .or. (j /= 16)
+ r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+ !$omp end target
+ if (r) call abort
+ !$omp target map (to: b(3:n), c(5:n+4), d(2:n+1), e(1:,:2), g(3:n), i(3:n), k(4:n), n) map (from: r)
+ r = (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+ r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+ r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+ r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+ r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+ r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+ r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+ r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+ !$omp end target
+ if (r) call abort
+ !$omp target map (to: b(5:n-2), c(7:n), d(4:n-2), e(1:,2:), g(5:n-3), i(6:n-4), k(5:n-5), n) map (from: r)
+ r = (any (b(5:n-2) /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+ r = r .or. (any (c(7:n) /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+ r = r .or. (any (d(4:n-2) /= 10)) .or. (lbound (d, 1) /= 2)
+ r = r .or. (any (e(1:,2:) /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+ r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+ r = r .or. (any (g(5:n-3) /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+ r = r .or. (any (i(6:n-4) /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+ r = r .or. (any (k(5:n-5) /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+ !$omp end target
+ !$omp target map (to: b(q+5:n-2+q), c(q+7:q+n), d(q+4:q+n-2), e(1:q+2,2:q+2), g(5+q:n-3+q), &
+ !$omp & i(6+q:n-4+q), k(5+q:n-5+q), n) map (from: r)
+ r = (any (b(5:n-2) /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+ r = r .or. (any (c(7:n) /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+ r = r .or. (any (d(4:n-2) /= 10)) .or. (lbound (d, 1) /= 2)
+ r = r .or. (any (e(1:,2:) /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+ r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+ r = r .or. (any (g(5:n-3) /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+ r = r .or. (any (i(6:n-4) /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+ r = r .or. (any (k(5:n-5) /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+ !$omp end target
+ if (r) call abort
+ !$omp target map (to: d(2:n+1), n)
+ r = a /= 7
+ r = r .or. (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+ r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+ r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+ r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+ r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+ r = r .or. (f /= 12)
+ r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+ r = r .or. (h /= 14)
+ r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+ r = r .or. (j /= 16)
+ r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+ !$omp end target
+ if (r) call abort
+ end subroutine foo
+end module target2
+ use target2, only : foo
+ integer, parameter :: n = 15, q = 0
+ integer :: a, b(2:n-1), c(n), d(n), e(3:4, 3:4)
+ integer, pointer :: f, g(:)
+ allocate (f, g(3:n))
+ a = 7
+ b = 8
+ c = 9
+ d = 10
+ e = 11
+ f = 12
+ g = 13
+ call foo (a, b, c, d, e, f, g, n, q)
+end
--- /dev/null
+! { dg-do run }
+
+module target3
+contains
+ subroutine foo (f, g)
+ integer :: n
+ integer, pointer :: f, g(:)
+ integer, pointer :: j, k(:)
+ logical :: r
+ nullify (j)
+ k => null ()
+ !$omp target map (tofrom: f, g, j, k) map (from: r)
+ r = associated (f) .or. associated (g)
+ r = r .or. associated (j) .or. associated (k)
+ !$omp end target
+ if (r) call abort
+ !$omp target
+ r = associated (f) .or. associated (g)
+ r = r .or. associated (j) .or. associated (k)
+ !$omp end target
+ if (r) call abort
+ end subroutine foo
+end module target3
+ use target3, only : foo
+ integer, pointer :: f, g(:)
+ f => null ()
+ nullify (g)
+ call foo (f, g)
+end
--- /dev/null
+! { dg-do run }
+
+module target4
+contains
+ subroutine foo (a,m,n)
+ integer :: m,n,i,j
+ double precision :: a(m, n), t
+ !$omp target data map(a) map(to: m, n)
+ do i=1,n
+ t = 0.0d0
+ !$omp target
+ !$omp parallel do reduction(+:t)
+ do j=1,m
+ t = t + a(j,i) * a(j,i)
+ end do
+ !$omp end target
+ t = 2.0d0 * t
+ !$omp target
+ !$omp parallel do
+ do j=1,m
+ a(j,i) = a(j,i) * t
+ end do
+ !$omp end target
+ end do
+ !$omp end target data
+ end subroutine foo
+end module target4
+ use target4, only : foo
+ integer :: i, j
+ double precision :: a(8, 9), res(8, 9)
+ do i = 1, 8
+ do j = 1, 9
+ a(i, j) = i + j
+ end do
+ end do
+ call foo (a, 8, 9)
+ res = reshape ((/ 1136.0d0, 1704.0d0, 2272.0d0, 2840.0d0, 3408.0d0, 3976.0d0, &
+& 4544.0d0, 5112.0d0, 2280.0d0, 3040.0d0, 3800.0d0, 4560.0d0, 5320.0d0, 6080.0d0, &
+& 6840.0d0, 7600.0d0, 3936.0d0, 4920.0d0, 5904.0d0, 6888.0d0, 7872.0d0, 8856.0d0, &
+& 9840.0d0, 10824.0d0, 6200.0d0, 7440.0d0, 8680.0d0, 9920.0d0, 11160.0d0, 12400.0d0, &
+& 13640.0d0, 14880.0d0, 9168.0d0, 10696.0d0, 12224.0d0, 13752.0d0, 15280.0d0, 16808.0d0, &
+& 18336.0d0, 19864.0d0, 12936.0d0, 14784.0d0, 16632.0d0, 18480.0d0, 20328.0d0, 22176.0d0, &
+& 24024.0d0, 25872.0d0, 17600.0d0, 19800.0d0, 22000.0d0, 24200.0d0, 26400.0d0, 28600.0d0, &
+& 30800.0d0, 33000.0d0, 23256.0d0, 25840.0d0, 28424.0d0, 31008.0d0, 33592.0d0, 36176.0d0, &
+& 38760.0d0, 41344.0d0, 30000.0d0, 33000.0d0, 36000.0d0, 39000.0d0, 42000.0d0, 45000.0d0, &
+& 48000.0d0, 51000.0d0 /), (/ 8, 9 /))
+ if (any (a /= res)) call abort
+end
--- /dev/null
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+ integer :: r
+ r = 0
+ call foo (r)
+ if (r /= 11) call abort
+contains
+ subroutine foo (r)
+ integer :: i, r
+ !$omp parallel
+ !$omp single
+ !$omp target teams distribute parallel do reduction (+: r)
+ do i = 1, 10
+ r = r + 1
+ end do
+ r = r + 1
+ !$omp end single
+ !$omp end parallel
+ end subroutine
+end
--- /dev/null
+! { dg-do run }
+
+module target6
+contains
+ subroutine foo (p, v, w, n)
+ double precision, pointer :: p(:), v(:), w(:)
+ double precision :: q(n)
+ integer :: i, n
+ !$omp target data if (n > 256) map (to: v(1:n), w(:n)) map (from: p(1:n), q)
+ !$omp target if (n > 256)
+ !$omp parallel do simd
+ do i = 1, n
+ p(i) = v(i) * w(i)
+ q(i) = p(i)
+ end do
+ !$omp end target
+ !$omp target update if (n > 256) from (p)
+ do i = 1, n
+ if (p(i) /= i * iand (i, 63)) call abort
+ v(i) = v(i) + 1
+ end do
+ !$omp target update if (n > 256) to (v(1:n))
+ !$omp target if (n > 256)
+ !$omp parallel do simd
+ do i = 1, n
+ p(i) = v(i) * w(i)
+ end do
+ !$omp end target
+ !$omp end target data
+ do i = 1, n
+ if (q(i) /= (v(i) - 1) * w(i)) call abort
+ if (p(i) /= q(i) + w(i)) call abort
+ end do
+ end subroutine
+end module target6
+ use target6, only : foo
+ integer :: n, i
+ double precision, pointer :: p(:), v(:), w(:)
+ n = 10000
+ allocate (p(n), v(n), w(n))
+ do i = 1, n
+ v(i) = i
+ w(i) = iand (i, 63)
+ end do
+ call foo (p, v, w, n)
+ do i = 1, n
+ if (p(i) /= (i + 1) * iand (i, 63)) call abort
+ end do
+ deallocate (p, v, w)
+end
--- /dev/null
+! { dg-do run }
+
+ interface
+ real function foo (x)
+ !$omp declare target
+ real, intent(in) :: x
+ end function foo
+ end interface
+ integer, parameter :: n = 1000
+ integer, parameter :: c = 100
+ integer :: i, j
+ real :: a(n)
+ do i = 1, n
+ a(i) = i
+ end do
+ do i = 1, n, c
+ !$omp task shared(a)
+ !$omp target map(a(i:i+c-1))
+ !$omp parallel do
+ do j = i, i + c - 1
+ a(j) = foo (a(j))
+ end do
+ !$omp end target
+ !$omp end task
+ end do
+ do i = 1, n
+ if (a(i) /= i + 1) call abort
+ end do
+end
+real function foo (x)
+ !$omp declare target
+ real, intent(in) :: x
+ foo = x + 1
+end function foo
projects / platform / upstream / gcc.git / commitdiff