diff options
53 files changed, 3583 insertions, 330 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index ef2dbd0..0fc6716 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,34 @@ +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. diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog index 9705f64..9b4818e 100644 --- a/gcc/cp/ChangeLog +++ b/gcc/cp/ChangeLog @@ -1,3 +1,9 @@ +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 diff --git a/gcc/cp/cp-gimplify.c b/gcc/cp/cp-gimplify.c index 2798358..296bd5f 100644 --- a/gcc/cp/cp-gimplify.c +++ b/gcc/cp/cp-gimplify.c @@ -1592,7 +1592,7 @@ cxx_omp_predetermined_sharing (tree decl) /* 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; diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 3e4ec3d..71298ef 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -6228,7 +6228,7 @@ extern tree cxx_omp_clause_default_ctor (tree, tree, tree); 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 */ diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index cbaf2ab..3a5d6aa 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,177 @@ +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 diff --git a/gcc/fortran/cpp.c b/gcc/fortran/cpp.c index 1695990..7fb8d16 100644 --- a/gcc/fortran/cpp.c +++ b/gcc/fortran/cpp.c @@ -171,7 +171,7 @@ cpp_define_builtins (cpp_reader *pfile) 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. diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c index 5a96119..de942f8 100644 --- a/gcc/fortran/dump-parse-tree.c +++ b/gcc/fortran/dump-parse-tree.c @@ -1016,32 +1016,51 @@ show_code (int level, gfc_code *c) } 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) { @@ -1117,7 +1136,7 @@ show_omp_node (int level, gfc_code *c) 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; @@ -1226,18 +1245,12 @@ show_omp_node (int level, gfc_code *c) 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) @@ -1269,6 +1282,34 @@ show_omp_node (int level, gfc_code *c) } 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) @@ -1296,7 +1337,8 @@ show_omp_node (int level, gfc_code *c) 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) diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c index 1962144..e41f61a 100644 --- a/gcc/fortran/f95-lang.c +++ b/gcc/fortran/f95-lang.c @@ -87,6 +87,24 @@ static alias_set_type gfc_get_alias_set (tree); 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 @@ -109,6 +127,7 @@ static tree gfc_builtin_function (tree); #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 @@ -116,6 +135,7 @@ static tree gfc_builtin_function (tree); #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" @@ -139,13 +159,15 @@ static tree gfc_builtin_function (tree); #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; diff --git a/gcc/fortran/frontend-passes.c b/gcc/fortran/frontend-passes.c index c69bd0c..4646cc3 100644 --- a/gcc/fortran/frontend-passes.c +++ b/gcc/fortran/frontend-passes.c @@ -2147,14 +2147,31 @@ gfc_code_walker (gfc_code **c, walk_code_fn_t codefn, walk_expr_fn_t exprfn, 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. */ @@ -2163,28 +2180,28 @@ gfc_code_walker (gfc_code **c, walk_code_fn_t codefn, walk_expr_fn_t exprfn, 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; diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 1df79fd..a11ca3d 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -215,6 +215,24 @@ typedef enum 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 } @@ -821,6 +839,9 @@ typedef struct !$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; @@ -1060,6 +1081,23 @@ typedef enum } 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.). */ @@ -1067,7 +1105,12 @@ typedef struct gfc_omp_namelist { 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; } @@ -1086,8 +1129,10 @@ enum 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 }; @@ -1147,6 +1192,11 @@ typedef struct gfc_omp_clauses 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; @@ -1387,7 +1437,7 @@ struct gfc_undo_change_set 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; @@ -2217,7 +2267,17 @@ typedef enum 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; @@ -2682,6 +2742,7 @@ bool gfc_add_protected (symbol_attribute *, const char *, locus *); 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 *); diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h index 86c2d1b..d07db11 100644 --- a/gcc/fortran/match.h +++ b/gcc/fortran/match.h @@ -131,6 +131,11 @@ match gfc_match_omp_cancellation_point (void); 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); @@ -144,10 +149,23 @@ match gfc_match_omp_parallel_workshare (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); diff --git a/gcc/fortran/module.c b/gcc/fortran/module.c index 261c904..bdd9961 100644 --- a/gcc/fortran/module.c +++ b/gcc/fortran/module.c @@ -1877,7 +1877,7 @@ typedef enum 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; @@ -1932,6 +1932,7 @@ static const mstring attr_bits[] = 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) }; @@ -2110,6 +2111,8 @@ mio_symbol_attribute (symbol_attribute *attr) 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 (); @@ -2273,6 +2276,9 @@ mio_symbol_attribute (symbol_attribute *attr) case AB_VTAB: attr->vtab = 1; break; + case AB_OMP_DECLARE_TARGET: + attr->omp_declare_target = 1; + break; } } } diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index a6e5f6c..266ac3d 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -72,6 +72,10 @@ gfc_free_omp_clauses (gfc_omp_clauses *c) 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); @@ -283,38 +287,45 @@ cleanup: 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; @@ -474,7 +485,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true, else for (n = *head; n; n = n->next) { - n->rop = rop; + n->u.reduction_op = rop; n->udr = udr; } continue; @@ -570,13 +581,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true, 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; @@ -662,21 +673,94 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = 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; @@ -699,7 +783,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true, | 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 \ @@ -715,100 +799,97 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true, (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)); } @@ -830,18 +911,6 @@ gfc_match_omp_flush (void) 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; @@ -1235,6 +1304,13 @@ gfc_match_omp_declare_reduction (void) 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 (); @@ -1243,6 +1319,102 @@ gfc_match_omp_declare_reduction (void) 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; @@ -1299,6 +1471,12 @@ gfc_match_omp_threadprivate (void) goto syntax; } + if (gfc_match_omp_eos () != MATCH_YES) + { + gfc_error ("Unexpected junk after OMP THREADPRIVATE at %C"); + goto cleanup; + } + return MATCH_YES; syntax: @@ -1311,83 +1489,213 @@ cleanup: 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) @@ -1602,8 +1910,8 @@ resolve_omp_clauses (gfc_code *code, locus *where, 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; @@ -1692,8 +2000,10 @@ resolve_omp_clauses (gfc_code *code, locus *where, 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) @@ -1745,6 +2055,20 @@ resolve_omp_clauses (gfc_code *code, locus *where, 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) { @@ -1819,8 +2143,10 @@ resolve_omp_clauses (gfc_code *code, locus *where, } } 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) { @@ -1829,11 +2155,11 @@ resolve_omp_clauses (gfc_code *code, locus *where, || 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; @@ -1842,19 +2168,20 @@ resolve_omp_clauses (gfc_code *code, locus *where, 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 @@ -1868,6 +2195,17 @@ resolve_omp_clauses (gfc_code *code, locus *where, } } } + 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) @@ -1917,7 +2255,7 @@ resolve_omp_clauses (gfc_code *code, locus *where, switch (list) { case OMP_LIST_REDUCTION: - switch (n->rop) + switch (n->u.reduction_op) { case OMP_REDUCTION_PLUS: case OMP_REDUCTION_TIMES: @@ -1964,7 +2302,7 @@ resolve_omp_clauses (gfc_code *code, locus *where, if (n->udr == NULL) { if (udr_name == NULL) - switch (n->rop) + switch (n->u.reduction_op) { case OMP_REDUCTION_PLUS: case OMP_REDUCTION_TIMES: @@ -1974,7 +2312,7 @@ resolve_omp_clauses (gfc_code *code, locus *where, 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"; @@ -1999,7 +2337,7 @@ resolve_omp_clauses (gfc_code *code, locus *where, gfc_typename (&n->sym->ts), where); } else - n->rop = OMP_REDUCTION_USER; + n->u.reduction_op = OMP_REDUCTION_USER; } break; case OMP_LIST_LINEAR: @@ -2051,6 +2389,38 @@ resolve_omp_clauses (gfc_code *code, locus *where, 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); + } } @@ -2565,14 +2935,38 @@ gfc_resolve_omp_parallel_blocks (gfc_code *code, gfc_namespace *ns) 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); @@ -2660,13 +3054,52 @@ resolve_omp_do (gfc_code *code) 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 (); } @@ -2786,11 +3219,23 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED) 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: @@ -2799,11 +3244,24 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED) 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; @@ -2822,7 +3280,7 @@ gfc_resolve_omp_declare_simd (gfc_namespace *ns) 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); diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c index bdee831..e8dcb70 100644 --- a/gcc/fortran/parse.c +++ b/gcc/fortran/parse.c @@ -633,12 +633,29 @@ decode_omp_directive (void) 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); @@ -654,8 +671,29 @@ decode_omp_directive (void) 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; @@ -685,10 +723,37 @@ decode_omp_directive (void) 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; @@ -1094,8 +1159,8 @@ next_statement (void) 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. */ @@ -1108,14 +1173,27 @@ next_statement (void) 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(). */ @@ -1621,6 +1699,21 @@ gfc_ascii_statement (gfc_statement st) 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; @@ -1633,6 +1726,18 @@ gfc_ascii_statement (gfc_statement st) 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; @@ -1672,9 +1777,45 @@ gfc_ascii_statement (gfc_statement st) 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; @@ -1714,6 +1855,30 @@ gfc_ascii_statement (gfc_statement st) 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; @@ -1726,6 +1891,21 @@ gfc_ascii_statement (gfc_statement st) 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; @@ -3699,13 +3879,47 @@ parse_omp_do (gfc_statement omp_st) 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) @@ -3814,12 +4028,60 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) 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; @@ -4052,6 +4314,10 @@ parse_executable (gfc_statement st) 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); @@ -4062,11 +4328,23 @@ parse_executable (gfc_statement st) 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; diff --git a/gcc/fortran/resolve.c b/gcc/fortran/resolve.c index 7ea7c36..64f3489 100644 --- a/gcc/fortran/resolve.c +++ b/gcc/fortran/resolve.c @@ -9032,6 +9032,10 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns) 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: @@ -9044,10 +9048,23 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns) 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; @@ -9827,11 +9844,23 @@ resolve_code (gfc_code *code, gfc_namespace *ns) 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: @@ -10160,6 +10189,10 @@ resolve_code (gfc_code *code, gfc_namespace *ns) 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: @@ -10167,9 +10200,23 @@ resolve_code (gfc_code *code, gfc_namespace *ns) 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; @@ -10179,7 +10226,6 @@ resolve_code (gfc_code *code, gfc_namespace *ns) 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); @@ -13541,6 +13587,18 @@ resolve_symbol (gfc_symbol *sym) || 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. */ diff --git a/gcc/fortran/st.c b/gcc/fortran/st.c index a3df43e..0f18f78 100644 --- a/gcc/fortran/st.c +++ b/gcc/fortran/st.c @@ -187,6 +187,10 @@ gfc_free_statement (gfc_code *p) 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: @@ -197,7 +201,20 @@ gfc_free_statement (gfc_code *p) 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); diff --git a/gcc/fortran/symbol.c b/gcc/fortran/symbol.c index 922b421..aee7510 100644 --- a/gcc/fortran/symbol.c +++ b/gcc/fortran/symbol.c @@ -367,6 +367,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where) *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; @@ -453,6 +454,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where) conf (dummy, entry); conf (dummy, intrinsic); conf (dummy, threadprivate); + conf (dummy, omp_declare_target); conf (pointer, target); conf (pointer, intrinsic); conf (pointer, elemental); @@ -495,6 +497,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where) conf (in_equivalence, entry); conf (in_equivalence, allocatable); conf (in_equivalence, threadprivate); + conf (in_equivalence, omp_declare_target); conf (dummy, result); conf (entry, result); @@ -543,6 +546,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where) 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); @@ -596,6 +600,8 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where) conf (proc_pointer, abstract) + conf (entry, omp_declare_target) + a1 = gfc_code2string (flavors, attr->flavor); if (attr->in_namelist @@ -631,6 +637,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where) conf2 (function); conf2 (subroutine); conf2 (threadprivate); + conf2 (omp_declare_target); if (attr->access == ACCESS_PUBLIC || attr->access == ACCESS_PRIVATE) { @@ -712,6 +719,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where) conf2 (subroutine); conf2 (threadprivate); conf2 (result); + conf2 (omp_declare_target); if (attr->intent != INTENT_UNKNOWN) { @@ -1207,6 +1215,22 @@ gfc_add_threadprivate (symbol_attribute *attr, const char *name, locus *where) 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) { @@ -1761,6 +1785,9 @@ gfc_copy_attr (symbol_attribute *dest, symbol_attribute *src, 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)) diff --git a/gcc/fortran/trans-common.c b/gcc/fortran/trans-common.c index 5a52984..f28eda6 100644 --- a/gcc/fortran/trans-common.c +++ b/gcc/fortran/trans-common.c @@ -456,6 +456,11 @@ build_common_decl (gfc_common_head *com, tree union_type, bool is_init) 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); diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c index 2e129c9..f1a18c3 100644 --- a/gcc/fortran/trans-decl.c +++ b/gcc/fortran/trans-decl.c @@ -1222,6 +1222,10 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list) list = chainon (list, attr); } + if (sym_attr.omp_declare_target) + list = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, list); + return list; } diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 998d687..7667f25 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -873,6 +873,110 @@ gfc_omp_clause_dtor (tree clause, tree decl) } +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 @@ -1487,7 +1591,7 @@ gfc_trans_omp_reduction_list (gfc_omp_namelist *namelist, tree list, 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; @@ -1532,7 +1636,7 @@ gfc_trans_omp_reduction_list (gfc_omp_namelist *namelist, tree list, 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); @@ -1661,8 +1765,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } } 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) @@ -1671,9 +1774,19 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, 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 { @@ -1691,13 +1804,286 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } 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; @@ -1920,7 +2306,69 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, 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. */ @@ -2329,12 +2777,13 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock, 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) @@ -2554,7 +3003,13 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock, } /* 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); @@ -2610,6 +3065,9 @@ enum 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 }; @@ -2617,7 +3075,10 @@ enum { 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 @@ -2628,10 +3089,32 @@ gfc_split_omp_clauses (gfc_code *code, 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; @@ -2640,11 +3123,99 @@ gfc_split_omp_clauses (gfc_code *code, 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. */ @@ -2659,9 +3230,6 @@ gfc_split_omp_clauses (gfc_code *code, = 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) { @@ -2701,6 +3269,12 @@ gfc_split_omp_clauses (gfc_code *code, /* 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]; @@ -2722,6 +3296,9 @@ gfc_split_omp_clauses (gfc_code *code, /* 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]; @@ -2731,6 +3308,13 @@ gfc_split_omp_clauses (gfc_code *code, 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)) @@ -2738,14 +3322,17 @@ gfc_split_omp_clauses (gfc_code *code, } 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) { @@ -2755,13 +3342,17 @@ gfc_trans_omp_do_simd (gfc_code *code, gfc_omp_clauses *clausesa, 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); @@ -2776,29 +3367,45 @@ gfc_trans_omp_do_simd (gfc_code *code, gfc_omp_clauses *clausesa, } 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; @@ -2807,25 +3414,39 @@ gfc_trans_omp_parallel_do (gfc_code *code) } 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, @@ -2969,6 +3590,170 @@ gfc_trans_omp_taskyield (void) } 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; @@ -3141,12 +3926,17 @@ gfc_trans_omp_directive (gfc_code *code) 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: @@ -3156,9 +3946,9 @@ gfc_trans_omp_directive (gfc_code *code) 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: @@ -3167,6 +3957,17 @@ gfc_trans_omp_directive (gfc_code *code) 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: @@ -3175,6 +3976,12 @@ gfc_trans_omp_directive (gfc_code *code) 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: diff --git a/gcc/fortran/trans.c b/gcc/fortran/trans.c index cfb8038..1925506 100644 --- a/gcc/fortran/trans.c +++ b/gcc/fortran/trans.c @@ -1851,6 +1851,10 @@ trans_code (gfc_code * code, tree cond) 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: @@ -1864,10 +1868,23 @@ trans_code (gfc_code * code, tree cond) 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; diff --git a/gcc/fortran/trans.h b/gcc/fortran/trans.h index 7ab9dd4..c272c0d 100644 --- a/gcc/fortran/trans.h +++ b/gcc/fortran/trans.h @@ -671,6 +671,7 @@ tree gfc_omp_clause_default_ctor (tree, tree, tree); 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); diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 10f8ac6..2efc899 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -5650,6 +5650,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) 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))) @@ -5662,8 +5663,12 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) 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; } @@ -6201,13 +6206,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, 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; @@ -6308,15 +6321,21 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) 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; @@ -6521,7 +6540,10 @@ gimplify_adjust_omp_clauses (tree *list_p) } /* 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); @@ -6552,7 +6574,7 @@ gimplify_omp_parallel (tree *expr_p, gimple_seq *pre_p) 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), @@ -6588,7 +6610,7 @@ gimplify_omp_task (tree *expr_p, gimple_seq *pre_p) 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), @@ -6934,7 +6956,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) 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)) @@ -7034,7 +7056,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) } 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)) { @@ -7073,7 +7095,7 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) 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)); diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h index d9a1dfd..76bb907 100644 --- a/gcc/langhooks-def.h +++ b/gcc/langhooks-def.h @@ -75,6 +75,7 @@ extern bool lhd_handle_option (size_t, const char *, int, int, location_t, 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); @@ -215,7 +216,7 @@ extern tree lhd_make_node (enum tree_code); #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, \ diff --git a/gcc/langhooks.c b/gcc/langhooks.c index 8f65c68..add0856 100644 --- a/gcc/langhooks.c +++ b/gcc/langhooks.c @@ -515,6 +515,13 @@ lhd_omp_assignment (tree clause ATTRIBUTE_UNUSED, tree dst, tree src) 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. */ diff --git a/gcc/langhooks.h b/gcc/langhooks.h index 35b47bc..33aa558 100644 --- a/gcc/langhooks.h +++ b/gcc/langhooks.h @@ -230,7 +230,7 @@ struct lang_hooks_for_decls 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. */ diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 67254cc..a30ce5a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1678,6 +1678,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } 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 @@ -2011,6 +2016,7 @@ scan_omp_parallel (gimple_stmt_iterator *gsi, omp_context *outer_ctx) 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); @@ -2508,6 +2514,23 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) 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; } @@ -9041,7 +9064,10 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) 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); diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 80c4e2a..9b76f23 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,14 @@ +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 diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-simd-1.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-simd-1.f90 new file mode 100644 index 0000000..d6ae7c9 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-simd-1.f90 @@ -0,0 +1,9 @@ +! { 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 diff --git a/gcc/testsuite/gfortran.dg/gomp/depend-1.f90 b/gcc/testsuite/gfortran.dg/gomp/depend-1.f90 new file mode 100644 index 0000000..bd6d26a --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/depend-1.f90 @@ -0,0 +1,13 @@ +! { 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 diff --git a/gcc/testsuite/gfortran.dg/gomp/target1.f90 b/gcc/testsuite/gfortran.dg/gomp/target1.f90 new file mode 100644 index 0000000..14db497 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target1.f90 @@ -0,0 +1,520 @@ +! { 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 diff --git a/gcc/testsuite/gfortran.dg/gomp/target2.f90 b/gcc/testsuite/gfortran.dg/gomp/target2.f90 new file mode 100644 index 0000000..7521331 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target2.f90 @@ -0,0 +1,74 @@ +! { 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 diff --git a/gcc/testsuite/gfortran.dg/gomp/target3.f90 b/gcc/testsuite/gfortran.dg/gomp/target3.f90 new file mode 100644 index 0000000..53a9682 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target3.f90 @@ -0,0 +1,12 @@ +! { 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 diff --git a/gcc/testsuite/gfortran.dg/gomp/udr4.f90 b/gcc/testsuite/gfortran.dg/gomp/udr4.f90 index 223dfd0..7e86a75 100644 --- a/gcc/testsuite/gfortran.dg/gomp/udr4.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/udr4.f90 @@ -6,7 +6,7 @@ subroutine f3 !$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) diff --git a/gcc/testsuite/gfortran.dg/openmp-define-3.f90 b/gcc/testsuite/gfortran.dg/openmp-define-3.f90 index 3d55986..44d5c9d 100644 --- a/gcc/testsuite/gfortran.dg/openmp-define-3.f90 +++ b/gcc/testsuite/gfortran.dg/openmp-define-3.f90 @@ -6,6 +6,6 @@ # error _OPENMP not defined #endif -#if _OPENMP != 201107 +#if _OPENMP != 201307 # error _OPENMP defined to wrong value #endif diff --git a/gcc/tree-core.h b/gcc/tree-core.h index aa7498b..a176553 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -1152,6 +1152,11 @@ enum omp_clause_map_kind 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 }; diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c index 85c6a03..ea2fb72 100644 --- a/gcc/tree-nested.c +++ b/gcc/tree-nested.c @@ -1085,6 +1085,10 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) 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: @@ -1113,10 +1117,42 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) 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: @@ -1126,6 +1162,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_COLLAPSE: case OMP_CLAUSE_UNTIED: case OMP_CLAUSE_MERGEABLE: + case OMP_CLAUSE_PROC_BIND: break; default: @@ -1620,6 +1657,10 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) 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: @@ -1653,12 +1694,45 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) 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: @@ -1666,6 +1740,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_COLLAPSE: case OMP_CLAUSE_UNTIED: case OMP_CLAUSE_MERGEABLE: + case OMP_CLAUSE_PROC_BIND: break; default: diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 3f6152f..59a825c 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -500,6 +500,7 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) 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: @@ -520,6 +521,9 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) 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), diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 8e6d37a..e3fdb62 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,22 @@ +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 diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in index dda297a..757053c 100644 --- a/libgomp/omp_lib.f90.in +++ b/libgomp/omp_lib.f90.in @@ -42,7 +42,7 @@ 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) diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in index 7725396..691adb8 100644 --- a/libgomp/omp_lib.h.in +++ b/libgomp/omp_lib.h.in @@ -45,7 +45,7 @@ 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 diff --git a/libgomp/testsuite/libgomp.c/target-8.c b/libgomp/testsuite/libgomp.c/target-8.c new file mode 100644 index 0000000..3508457 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-8.c @@ -0,0 +1,26 @@ +/* { 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; +} diff --git a/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90 b/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90 index b141a3b..5cd592c 100644 --- a/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90 @@ -6,7 +6,8 @@ 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) @@ -22,6 +23,7 @@ end module declare_simd_1_mod 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 diff --git a/libgomp/testsuite/libgomp.fortran/depend-3.f90 b/libgomp/testsuite/libgomp.fortran/depend-3.f90 new file mode 100644 index 0000000..11be641 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/depend-3.f90 @@ -0,0 +1,42 @@ +! { 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 diff --git a/libgomp/testsuite/libgomp.fortran/openmp_version-1.f b/libgomp/testsuite/libgomp.fortran/openmp_version-1.f index aaa8881..be24adc 100644 --- a/libgomp/testsuite/libgomp.fortran/openmp_version-1.f +++ b/libgomp/testsuite/libgomp.fortran/openmp_version-1.f @@ -4,6 +4,6 @@ implicit none include "omp_lib.h" - if (openmp_version .ne. 201107) call abort; + if (openmp_version .ne. 201307) call abort; end program main diff --git a/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90 b/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90 index b2d1d26..62712c7 100644 --- a/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90 +++ b/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90 @@ -4,6 +4,6 @@ program main use omp_lib implicit none - if (openmp_version .ne. 201107) call abort; + if (openmp_version .ne. 201307) call abort; end program main diff --git a/libgomp/testsuite/libgomp.fortran/target1.f90 b/libgomp/testsuite/libgomp.fortran/target1.f90 new file mode 100644 index 0000000..c70daac --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target1.f90 @@ -0,0 +1,58 @@ +! { 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 diff --git a/libgomp/testsuite/libgomp.fortran/target2.f90 b/libgomp/testsuite/libgomp.fortran/target2.f90 new file mode 100644 index 0000000..42f704f --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target2.f90 @@ -0,0 +1,96 @@ +! { 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 diff --git a/libgomp/testsuite/libgomp.fortran/target3.f90 b/libgomp/testsuite/libgomp.fortran/target3.f90 new file mode 100644 index 0000000..1f197ac --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target3.f90 @@ -0,0 +1,29 @@ +! { 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 diff --git a/libgomp/testsuite/libgomp.fortran/target4.f90 b/libgomp/testsuite/libgomp.fortran/target4.f90 new file mode 100644 index 0000000..aa2f0a5 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target4.f90 @@ -0,0 +1,48 @@ +! { 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 diff --git a/libgomp/testsuite/libgomp.fortran/target5.f90 b/libgomp/testsuite/libgomp.fortran/target5.f90 new file mode 100644 index 0000000..c46faf2 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target5.f90 @@ -0,0 +1,21 @@ +! { 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 diff --git a/libgomp/testsuite/libgomp.fortran/target6.f90 b/libgomp/testsuite/libgomp.fortran/target6.f90 new file mode 100644 index 0000000..13f5a52 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target6.f90 @@ -0,0 +1,50 @@ +! { 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 diff --git a/libgomp/testsuite/libgomp.fortran/target7.f90 b/libgomp/testsuite/libgomp.fortran/target7.f90 new file mode 100644 index 0000000..4af0ee3 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target7.f90 @@ -0,0 +1,34 @@ +! { 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 |