[PATCH] Fortran OpenMP 4.0 target support

classic Classic list List threaded Threaded
5 messages Options
Reply | Threaded
Open this post in threaded view
|

[PATCH] Fortran OpenMP 4.0 target support

Jakub Jelinek
Hi!

This patch adds the target directives.
Tested both normally plus with target.c/splay-tree.c from
gomp-4_0-branch@203409 plus the attached patch against
target.c to implement the new to_pset map kind (5) and
allow handling of NULL.  That patch will need to be forward
ported to whatever gomp-4_0-branch now has after this is merged
from trunk to that branch.

Does this look reasonable to Fortran maintainers?

2014-06-17  Jakub Jelinek  <[hidden email]>

        * 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.
gcc/cp/
        * cp-gimplify.c (cxx_omp_finish_clause): Add a gimple_seq *
        argument.
        * cp-tree.h (cxx_omp_finish_clause): Adjust prototype.
gcc/fortran/
        * 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.
gcc/testsuite/
        * 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.
libgomp/
        * 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.

--- gcc/gimplify.c.jj 2014-06-12 23:09:06.000000000 +0200
+++ gcc/gimplify.c 2014-06-17 15:36:01.513908231 +0200
@@ -5650,6 +5650,7 @@ omp_notice_variable (struct gimplify_omp
   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
     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,
   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_tre
       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_
     }
 
   /* 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, gim
   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_
   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_s
  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, gi
     }
   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
 
   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));
 
--- gcc/langhooks.c.jj 2014-05-22 10:18:10.000000000 +0200
+++ gcc/langhooks.c 2014-06-16 20:12:40.093228650 +0200
@@ -515,6 +515,13 @@ lhd_omp_assignment (tree clause ATTRIBUT
   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.  */
 
--- gcc/langhooks-def.h.jj 2014-05-22 10:18:11.000000000 +0200
+++ gcc/langhooks-def.h 2014-06-16 20:11:53.007462562 +0200
@@ -75,6 +75,7 @@ extern bool lhd_handle_option (size_t, c
 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_cod
 #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, \
--- gcc/langhooks.h.jj 2014-05-22 10:18:10.000000000 +0200
+++ gcc/langhooks.h 2014-06-16 20:20:15.134886638 +0200
@@ -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.  */
--- gcc/omp-low.c.jj 2014-06-12 23:09:06.000000000 +0200
+++ gcc/omp-low.c 2014-06-17 12:01:57.619782584 +0200
@@ -1678,6 +1678,11 @@ scan_sharing_clauses (tree clauses, omp_
  }
       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
       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 s
   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
  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);
--- gcc/tree-core.h.jj 2014-06-16 10:07:30.738832010 +0200
+++ gcc/tree-core.h 2014-06-16 10:34:19.403582595 +0200
@@ -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
 };
 
--- gcc/tree-nested.c.jj 2014-06-09 17:44:37.000000000 +0200
+++ gcc/tree-nested.c 2014-06-17 12:32:06.186450886 +0200
@@ -1085,6 +1085,10 @@ convert_nonlocal_omp_clauses (tree *pcla
  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 *pcla
  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 *pcla
  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 *pclause
  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 *pclause
  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 *pclause
  case OMP_CLAUSE_COLLAPSE:
  case OMP_CLAUSE_UNTIED:
  case OMP_CLAUSE_MERGEABLE:
+ case OMP_CLAUSE_PROC_BIND:
   break;
 
  default:
--- gcc/tree-pretty-print.c.jj 2014-06-16 10:06:40.359092881 +0200
+++ gcc/tree-pretty-print.c 2014-06-16 10:34:19.400582611 +0200
@@ -500,6 +500,7 @@ dump_omp_clause (pretty_printer *buffer,
   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,
   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),
--- gcc/cp/cp-gimplify.c.jj 2014-06-06 09:19:20.000000000 +0200
+++ gcc/cp/cp-gimplify.c 2014-06-16 20:13:41.766908514 +0200
@@ -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;
--- gcc/cp/cp-tree.h.jj 2014-06-02 10:34:11.000000000 +0200
+++ gcc/cp/cp-tree.h 2014-06-16 20:13:24.116999529 +0200
@@ -6228,7 +6228,7 @@ extern tree cxx_omp_clause_default_ctor
 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 */
--- gcc/fortran/cpp.c.jj 2014-03-22 08:17:22.000000000 +0100
+++ gcc/fortran/cpp.c 2014-06-16 18:45:43.736825272 +0200
@@ -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.
 
--- gcc/fortran/dump-parse-tree.c.jj 2014-06-16 10:06:39.374097978 +0200
+++ gcc/fortran/dump-parse-tree.c 2014-06-16 10:34:19.350582873 +0200
@@ -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)
--- gcc/fortran/f95-lang.c.jj 2014-06-16 10:06:39.513097254 +0200
+++ gcc/fortran/f95-lang.c 2014-06-16 19:30:52.256006765 +0200
@@ -87,6 +87,24 @@ static alias_set_type gfc_get_alias_set
 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;
 
--- gcc/fortran/frontend-passes.c.jj 2014-06-16 10:06:38.996099917 +0200
+++ gcc/fortran/frontend-passes.c 2014-06-16 10:34:19.349582875 +0200
@@ -2147,14 +2147,31 @@ gfc_code_walker (gfc_code **c, walk_code
       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
 
       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;
--- gcc/fortran/gfortran.h.jj 2014-06-16 10:06:39.247098599 +0200
+++ gcc/fortran/gfortran.h 2014-06-16 10:34:19.341582926 +0200
@@ -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
 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 *);
--- gcc/fortran/match.h.jj 2014-06-16 10:06:39.449097520 +0200
+++ gcc/fortran/match.h 2014-06-16 10:34:19.355582852 +0200
@@ -131,6 +131,11 @@ match gfc_match_omp_cancellation_point (
 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 (
 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);
--- gcc/fortran/module.c.jj 2014-06-16 10:06:39.116099328 +0200
+++ gcc/fortran/module.c 2014-06-16 10:34:19.374582762 +0200
@@ -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 *
  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 *
     case AB_VTAB:
       attr->vtab = 1;
       break;
+    case AB_OMP_DECLARE_TARGET:
+      attr->omp_declare_target = 1;
+      break;
     }
  }
     }
--- gcc/fortran/openmp.c.jj 2014-06-16 10:06:39.409097865 +0200
+++ gcc/fortran/openmp.c 2014-06-17 19:07:34.110844226 +0200
@@ -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 *
       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 *
       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 *
   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 *
    | 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 *
   (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, loc
   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, loc
     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, loc
  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, loc
   }
       }
     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, loc
       || 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, loc
  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, loc
   }
     }
  }
+    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, loc
  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, loc
  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, loc
  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, loc
        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, loc
  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_cod
   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 *cod
 
   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 *cod
     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_namesp
   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);
--- gcc/fortran/parse.c.jj 2014-06-16 10:06:39.491097364 +0200
+++ gcc/fortran/parse.c 2014-06-17 19:06:39.122116074 +0200
@@ -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_statemen
     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;
--- gcc/fortran/resolve.c.jj 2014-06-16 10:06:39.050099644 +0200
+++ gcc/fortran/resolve.c 2014-06-16 10:34:19.391582656 +0200
@@ -9027,6 +9027,10 @@ gfc_resolve_blocks (gfc_code *b, gfc_nam
 
  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:
@@ -9039,10 +9043,23 @@ gfc_resolve_blocks (gfc_code *b, gfc_nam
  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;
 
@@ -9808,11 +9825,23 @@ resolve_code (gfc_code *code, gfc_namesp
     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:
@@ -10139,6 +10168,10 @@ resolve_code (gfc_code *code, gfc_namesp
  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:
@@ -10146,9 +10179,23 @@ resolve_code (gfc_code *code, gfc_namesp
  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;
@@ -10158,7 +10205,6 @@ resolve_code (gfc_code *code, gfc_namesp
  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);
@@ -13520,6 +13566,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.  */
--- gcc/fortran/st.c.jj 2014-06-16 10:06:39.210098806 +0200
+++ gcc/fortran/st.c 2014-06-16 10:34:19.357582835 +0200
@@ -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);
--- gcc/fortran/symbol.c.jj 2014-06-16 10:06:39.480097325 +0200
+++ gcc/fortran/symbol.c 2014-06-17 19:06:39.139117070 +0200
@@ -367,6 +367,7 @@ check_conflict (symbol_attribute *attr,
     *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,
   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,
   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,
   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,
 
   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,
       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,
       conf2 (subroutine);
       conf2 (threadprivate);
       conf2 (result);
+      conf2 (omp_declare_target);
 
       if (attr->intent != INTENT_UNKNOWN)
  {
@@ -1207,6 +1215,22 @@ gfc_add_threadprivate (symbol_attribute
 
 
 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, s
   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))
--- gcc/fortran/trans.c.jj 2014-06-16 10:06:39.289098442 +0200
+++ gcc/fortran/trans.c 2014-06-16 10:34:19.375582759 +0200
@@ -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;
--- gcc/fortran/trans-common.c.jj 2014-06-16 10:07:26.193855540 +0200
+++ gcc/fortran/trans-common.c 2014-06-16 10:34:19.358582830 +0200
@@ -456,6 +456,11 @@ build_common_decl (gfc_common_head *com,
       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);
--- gcc/fortran/trans-decl.c.jj 2014-06-16 10:07:28.434843937 +0200
+++ gcc/fortran/trans-decl.c 2014-06-16 10:34:19.342582933 +0200
@@ -1219,6 +1219,10 @@ add_attributes_to_decl (symbol_attribute
  list = chainon (list, attr);
       }
 
+  if (sym_attr.omp_declare_target)
+    list = tree_cons (get_identifier ("omp declare target"),
+      NULL_TREE, list);
+
   return list;
 }
 
--- gcc/fortran/trans.h.jj 2014-06-16 10:07:28.000000000 +0200
+++ gcc/fortran/trans.h 2014-06-16 20:18:36.113414391 +0200
@@ -670,6 +670,7 @@ tree gfc_omp_clause_default_ctor (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);
--- gcc/fortran/trans-openmp.c.jj 2014-06-16 10:06:39.164099047 +0200
+++ gcc/fortran/trans-openmp.c 2014-06-17 19:32:58.939176877 +0200
@@ -873,6 +873,110 @@ gfc_omp_clause_dtor (tree clause, tree d
 }
 
 
+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_na
     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_na
  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 *bloc
       }
   }
   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 *bloc
       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 *bloc
     }
   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 *bloc
       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_ex
 
       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_ex
     }
 
   /* 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, g
   if (gfc_option.gfc_flag_openmp)
     omp_do_clauses
       = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_DO], code->loc);
-  pblock = &block;
-  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, g
 }
 
 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 = &block;
-  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 = &block;
+      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 *cod
 }
 
 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:
--- gcc/testsuite/gfortran.dg/gomp/declare-simd-1.f90.jj 2014-06-16 10:34:19.404582590 +0200
+++ gcc/testsuite/gfortran.dg/gomp/declare-simd-1.f90 2014-06-16 10:34:19.404582590 +0200
@@ -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
--- gcc/testsuite/gfortran.dg/gomp/depend-1.f90.jj 2014-06-16 10:34:19.404582590 +0200
+++ gcc/testsuite/gfortran.dg/gomp/depend-1.f90 2014-06-16 10:34:19.404582590 +0200
@@ -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
--- gcc/testsuite/gfortran.dg/gomp/target1.f90.jj 2014-06-16 10:34:19.404582590 +0200
+++ gcc/testsuite/gfortran.dg/gomp/target1.f90 2014-06-17 10:21:51.111817953 +0200
@@ -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
--- gcc/testsuite/gfortran.dg/gomp/target2.f90.jj 2014-06-17 09:40:12.646898771 +0200
+++ gcc/testsuite/gfortran.dg/gomp/target2.f90 2014-06-17 09:47:57.000000000 +0200
@@ -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
--- gcc/testsuite/gfortran.dg/gomp/target3.f90.jj 2014-06-17 12:04:22.852029201 +0200
+++ gcc/testsuite/gfortran.dg/gomp/target3.f90 2014-06-17 12:04:46.864901093 +0200
@@ -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
--- gcc/testsuite/gfortran.dg/gomp/udr4.f90.jj 2014-06-16 10:06:40.009094670 +0200
+++ gcc/testsuite/gfortran.dg/gomp/udr4.f90 2014-06-16 10:34:19.404582590 +0200
@@ -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)
--- gcc/testsuite/gfortran.dg/openmp-define-3.f90.jj 2013-10-17 22:30:48.854215211 +0200
+++ gcc/testsuite/gfortran.dg/openmp-define-3.f90 2014-06-17 21:25:25.385141631 +0200
@@ -6,6 +6,6 @@
 # error _OPENMP not defined
 #endif
 
-#if _OPENMP != 201107
+#if _OPENMP != 201307
 # error _OPENMP defined to wrong value
 #endif
--- libgomp/omp_lib.f90.in.jj 2014-01-03 11:41:28.000000000 +0100
+++ libgomp/omp_lib.f90.in 2014-06-16 18:45:08.582999946 +0200
@@ -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)
--- libgomp/omp_lib.h.in.jj 2014-01-03 11:41:29.000000000 +0100
+++ libgomp/omp_lib.h.in 2014-06-16 18:46:03.332717272 +0200
@@ -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
--- libgomp/testsuite/libgomp.c/target-8.c.jj 2014-06-17 11:43:34.942505638 +0200
+++ libgomp/testsuite/libgomp.c/target-8.c 2014-06-17 11:43:16.000000000 +0200
@@ -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;
+}
--- libgomp/testsuite/libgomp.fortran/declare-simd-1.f90.jj 2014-05-30 10:51:05.000000000 +0200
+++ libgomp/testsuite/libgomp.fortran/declare-simd-1.f90 2014-06-17 16:53:22.560503385 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/depend-3.f90.jj 2014-06-16 10:34:19.405582586 +0200
+++ libgomp/testsuite/libgomp.fortran/depend-3.f90 2014-06-16 10:34:19.405582586 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/openmp_version-1.f.jj 2013-10-11 17:19:46.000000000 +0200
+++ libgomp/testsuite/libgomp.fortran/openmp_version-1.f 2014-06-17 16:49:13.830772438 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/openmp_version-2.f90.jj 2013-10-11 17:19:46.000000000 +0200
+++ libgomp/testsuite/libgomp.fortran/openmp_version-2.f90 2014-06-17 16:49:23.315723912 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/target1.f90.jj 2014-06-16 10:34:19.405582586 +0200
+++ libgomp/testsuite/libgomp.fortran/target1.f90 2014-06-17 12:50:05.940874685 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/target2.f90.jj 2014-06-16 10:34:19.405582586 +0200
+++ libgomp/testsuite/libgomp.fortran/target2.f90 2014-06-17 12:50:37.560727850 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/target3.f90.jj 2014-06-16 18:16:33.319775473 +0200
+++ libgomp/testsuite/libgomp.fortran/target3.f90 2014-06-17 07:48:23.527848883 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/target4.f90.jj 2014-06-17 08:13:28.808148961 +0200
+++ libgomp/testsuite/libgomp.fortran/target4.f90 2014-06-17 12:53:00.195970409 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/target5.f90.jj 2014-06-17 12:03:58.592155260 +0200
+++ libgomp/testsuite/libgomp.fortran/target5.f90 2014-06-17 12:03:41.000000000 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/target6.f90.jj 2014-06-17 13:45:19.357750928 +0200
+++ libgomp/testsuite/libgomp.fortran/target6.f90 2014-06-17 16:31:45.213119705 +0200
@@ -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
--- libgomp/testsuite/libgomp.fortran/target7.f90.jj 2014-06-17 15:55:20.232999819 +0200
+++ libgomp/testsuite/libgomp.fortran/target7.f90 2014-06-17 17:07:08.606278813 +0200
@@ -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

        Jakub

target.patch (5K) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH] Fortran OpenMP 4.0 target support

Tobias Burnus
Jakub Jelinek wrote:
> This patch adds the target directives.
> Tested both normally plus with target.c/splay-tree.c from
> gomp-4_0-branch@203409 plus the attached patch against
> target.c to implement the new to_pset map kind (5) and
> allow handling of NULL.  That patch will need to be forward
> ported to whatever gomp-4_0-branch now has after this is merged
> from trunk to that branch.
>
> Does this look reasonable to Fortran maintainers?

Thanks for the patch! I browsed through the patch, and it looked good to
me. (However, given that the patch has "48 files changed, 3342
insertions(+), 330 deletions(-)", I didn't check every line.)

If I did the book keeping correctly, a patch for an alignment test case
is still missing. As are the changes for some corner cases for which the
OpenMP ARB has to provide some feedback. Any news from that side?
Otherwise and aside of 4.9.1 backporting, it now looks pretty complete.

Tobias
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH] Fortran OpenMP 4.0 target support

Jakub Jelinek
On Tue, Jun 17, 2014 at 11:59:22PM +0200, Tobias Burnus wrote:

> >This patch adds the target directives.
> >Tested both normally plus with target.c/splay-tree.c from
> >gomp-4_0-branch@203409 plus the attached patch against
> >target.c to implement the new to_pset map kind (5) and
> >allow handling of NULL.  That patch will need to be forward
> >ported to whatever gomp-4_0-branch now has after this is merged
> >from trunk to that branch.
> >
> >Does this look reasonable to Fortran maintainers?
>
> Thanks for the patch! I browsed through the patch, and it looked good to me.
> (However, given that the patch has "48 files changed, 3342 insertions(+),
> 330 deletions(-)", I didn't check every line.)
>
> If I did the book keeping correctly, a patch for an alignment test case is
> still missing. As are the changes for some corner cases for which the OpenMP
> ARB has to provide some feedback. Any news from that side? Otherwise and
> aside of 4.9.1 backporting, it now looks pretty complete.

I think some work is needed in tree-nested.c, ideally write a testcase
that tests all the new OpenMP 4.0 clauses in contained functions with and
without non-local decls (and with local decls used by contained functions).

One of the omp-lang answers shows some work is needed on the UDRs too, in
particular that the combiner/initializer should not be resolved as part of
the UDR directive, but only when used in a reduction clause where not only
the typespec, but also rank/shape, pointer/allocatable etc. are known.

Some further restriction checking is probably needed + backing that with
testcases.  And wait for further omp-lang/omp-f2003 feedback.

        Jakub
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH] Fortran OpenMP 4.0 target support

Thomas Schwinge-8
In reply to this post by Jakub Jelinek
Hi!

On Tue, 17 Jun 2014 23:03:47 +0200, Jakub Jelinek <[hidden email]> wrote:
> This patch adds the target directives.
> Tested both normally plus with target.c/splay-tree.c from
> gomp-4_0-branch@203409 plus the attached patch against
> target.c to implement the new to_pset map kind (5) and
> allow handling of NULL.  That patch will need to be forward
> ported to whatever gomp-4_0-branch now has after this is merged
> from trunk to that branch.

Thanks for the target.c patch.  In r212405, I have committed your patch
to gomp-4_0-branch.  All the new libgomp.fortran/target* tests pass with
the libgomp/plugin-host.c plugin.

commit fb4cfa7ef6e819fa0108ca75b52b86a5826a863f
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Wed Jul 9 19:59:13 2014 +0000

    libgomp: Support NULL mappings as well as mapping kind OMP_CLAUSE_MAP_TO_PSET.
   
    libgomp/
    * target.c (gomp_map_vars, gomp_unmap_vars, gomp_update): Support
    NULL mappings as well as mapping kind OMP_CLAUSE_MAP_TO_PSET.
    Also, some code reformatting.
   
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@212405 138bc75d-0d04-0410-961f-82ee72b054a4

That is your patch with the following folded in:

--- libgomp/target.c
+++ libgomp/target.c
@@ -357,10 +357,10 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
     /* FIXME: Perhaps add some smarts, like if copying
        several adjacent fields from host to target, use some
        host buffer to avoid sending each var individually.  */
-    devicep->device_host2dev_func((void *) (tgt->tgt_start
-    + k->tgt_offset),
-  (void *) k->host_start,
-  k->host_end - k->host_start);
+    devicep->device_host2dev_func
+      ((void *) (tgt->tgt_start + k->tgt_offset),
+       (void *) k->host_start,
+       k->host_end - k->host_start);
     break;
   case 4: /* POINTER */
     cur_node.host_start
@@ -368,11 +368,12 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
     if (cur_node.host_start == (uintptr_t) NULL)
       {
  cur_node.tgt_offset = (uintptr_t) NULL;
- /* FIXME: host to device copy, see above FIXME
-   comment.  */
- memcpy ((void *) (tgt->tgt_start + k->tgt_offset),
- (void *) &cur_node.tgt_offset,
- sizeof (void *));
+ /* Copy from host to device memory.  */
+ /* FIXME: see above FIXME comment.  */
+ devicep->device_host2dev_func
+  ((void *) (tgt->tgt_start + k->tgt_offset),
+   (void *) &cur_node.tgt_offset,
+   sizeof (void *));
  break;
       }
     /* Add bias to the pointer value.  */
@@ -406,19 +407,18 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
     cur_node.tgt_offset -= sizes[i];
     /* Copy from host to device memory.  */
     /* FIXME: see above FIXME comment.  */
-    devicep->device_host2dev_func ((void *) (tgt->tgt_start
-     + k->tgt_offset),
-   (void *) &cur_node.tgt_offset,
-   sizeof (void *));
+    devicep->device_host2dev_func
+      ((void *) (tgt->tgt_start + k->tgt_offset),
+       (void *) &cur_node.tgt_offset,
+       sizeof (void *));
     break;
   case 5: /* TO_PSET */
-    /* FIXME: This is supposed to be copy from host to device
-       memory.  Perhaps add some smarts, like if copying
-       several adjacent fields from host to target, use some
-       host buffer to avoid sending each var individually.  */
-    memcpy ((void *) (tgt->tgt_start + k->tgt_offset),
-    (void *) k->host_start,
-    k->host_end - k->host_start);
+    /* Copy from host to device memory.  */
+    /* FIXME: see above FIXME comment.  */
+    devicep->device_host2dev_func
+      ((void *) (tgt->tgt_start + k->tgt_offset),
+       (void *) k->host_start,
+       (k->host_end - k->host_start));
     for (j = i + 1; j < mapnum; j++)
       if ((kinds[j] & 7) != 4)
  break;
@@ -435,13 +435,14 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
   if (cur_node.host_start == (uintptr_t) NULL)
     {
       cur_node.tgt_offset = (uintptr_t) NULL;
-      /* FIXME: host to device copy, see above FIXME
- comment.  */
-      memcpy ((void *) (tgt->tgt_start + k->tgt_offset
- + ((uintptr_t) hostaddrs[j]
-   - k->host_start)),
-      (void *) &cur_node.tgt_offset,
-      sizeof (void *));
+      /* Copy from host to device memory.  */
+      /* FIXME: see above FIXME comment.  */
+      devicep->device_host2dev_func
+ ((void *) (tgt->tgt_start + k->tgt_offset
+   + ((uintptr_t) hostaddrs[j]
+      - k->host_start)),
+ (void *) &cur_node.tgt_offset,
+ sizeof (void *));
       i++;
       continue;
     }
@@ -475,13 +476,14 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
      array section.  Now subtract bias to get what we
      want to initialize the pointer with.  */
   cur_node.tgt_offset -= sizes[j];
-  /* FIXME: host to device copy, see above FIXME
-     comment.  */
-  memcpy ((void *) (tgt->tgt_start + k->tgt_offset
-    + ((uintptr_t) hostaddrs[j]
-       - k->host_start)),
-  (void *) &cur_node.tgt_offset,
-  sizeof (void *));
+  /* Copy from host to device memory.  */
+  /* FIXME: see above FIXME comment.  */
+  devicep->device_host2dev_func
+    ((void *) (tgt->tgt_start + k->tgt_offset
+       + ((uintptr_t) hostaddrs[j]
+  - k->host_start)),
+     (void *) &cur_node.tgt_offset,
+     sizeof (void *));
   i++;
  }
       break;
@@ -501,10 +503,10 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
   + tgt->list[i]->tgt_offset;
   /* Copy from host to device memory.  */
   /* FIXME: see above FIXME comment.  */
-  devicep->device_host2dev_func ((void *) (tgt->tgt_start
-   + i * sizeof (void *)),
- (void *) &cur_node.tgt_offset,
- sizeof (void *));
+  devicep->device_host2dev_func
+    ((void *) (tgt->tgt_start + i * sizeof (void *)),
+     (void *) &cur_node.tgt_offset,
+     sizeof (void *));
  }
     }
 
@@ -546,10 +548,10 @@ gomp_unmap_vars (struct target_mem_desc *tgt)
  splay_tree_key k = tgt->list[i];
  if (k->copy_from)
   /* Copy from device to host memory.  */
-  devicep->device_dev2host_func ((void *) k->host_start,
- (void *) (k->tgt->tgt_start
-   + k->tgt_offset),
- k->host_end - k->host_start);
+  devicep->device_dev2host_func
+    ((void *) k->host_start,
+     (void *) (k->tgt->tgt_start + k->tgt_offset),
+     k->host_end - k->host_start);
  splay_tree_remove (&devicep->dev_splay_tree, k);
  if (k->tgt->refcount > 1)
   k->tgt->refcount--;
@@ -597,22 +599,22 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum,
   (void *) n->host_end);
     if ((kinds[i] & 7) == 1)
       /* Copy from host to device memory.  */
-      devicep->device_host2dev_func ((void *) (n->tgt->tgt_start
-       + n->tgt_offset
-       + cur_node.host_start
-       - n->host_start),
-     (void *) cur_node.host_start,
-     cur_node.host_end
-     - cur_node.host_start);
+      devicep->device_host2dev_func
+ ((void *) (n->tgt->tgt_start
+   + n->tgt_offset
+   + cur_node.host_start
+   - n->host_start),
+ (void *) cur_node.host_start,
+ cur_node.host_end - cur_node.host_start);
     else if ((kinds[i] & 7) == 2)
       /* Copy from device to host memory.  */
-      devicep->device_dev2host_func ((void *) cur_node.host_start,
-     (void *) (n->tgt->tgt_start
-       + n->tgt_offset
-       + cur_node.host_start
-       - n->host_start),
-     cur_node.host_end
-     - cur_node.host_start);
+      devicep->device_dev2host_func
+ ((void *) cur_node.host_start,
+ (void *) (n->tgt->tgt_start
+   + n->tgt_offset
+   + cur_node.host_start
+   - n->host_start),
+ cur_node.host_end - cur_node.host_start);
   }
  else
   gomp_fatal ("Trying to update [%p..%p) object that is not mapped",


Grüße,
 Thomas

attachment0 (482 bytes) Download Attachment
Reply | Threaded
Open this post in threaded view
|

PR90030 "Fortran OpenACC subarray data alignment" (was: [PATCH] Fortran OpenMP 4.0 target support)

Thomas Schwinge-8
In reply to this post by Jakub Jelinek
Hi Jakub!

In context of PR90030 "Fortran OpenACC subarray data alignment" (which
actually is reproducible for OpenMP with nvptx offloading in the very
same way, see below), can you please explain the reason for the seven
"[var] = fold_convert (build_pointer_type (char_type_node), [var])"
instances that you've added as part of your 2014 trunk r211768 "Fortran
OpenMP 4.0 target support" commit?

Replacing all these with "gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)))"
(see the attached WIP patch, which also includes an OpenMP test case), I
don't see any ill effects for 'check-gcc-fortran', and
'check-target-libgomp' with nvptx offloading, and the errors 'libgomp:
cuStreamSynchronize error: misaligned address' are gone.  I added these
'gcc_assert's just for checking; Cesar in
<https://gcc.gnu.org/ml/gcc-patches/2015-09/msg01664.html>, and Julian in
<https://gcc.gnu.org/ml/gcc-patches/2018-08/msg01911.html> propose to
simply drop (a subset of) these casts.  Do we need (a) all, (b) some, (c)
none of these casts?  And do we want to replace them with 'gcc_assert's,
or not do that?

If approving such a patch (for all release branches), please respond with
"Reviewed-by: NAME <EMAIL>" so that your effort will be recorded in the
commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>.

For reference, see the seven 'char_type_node' instances:

On Tue, 17 Jun 2014 23:03:47 +0200, Jakub Jelinek <[hidden email]> wrote:

> --- gcc/fortran/trans-openmp.c.jj 2014-06-16 10:06:39.164099047 +0200
> +++ gcc/fortran/trans-openmp.c 2014-06-17 19:32:58.939176877 +0200
> @@ -873,6 +873,110 @@ gfc_omp_clause_dtor (tree clause, tree d
>  }
>  
>  
> +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_na
>      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_na
>   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 *bloc
>        }
>    }
>    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 *bloc
>        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 *bloc
>      }
>    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 *bloc
>        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);
>  }

Grüße
 Thomas



From b8b60e9c1b0bd100717bf52e312f9285b8ac8fbf Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <[hidden email]>
Date: Wed, 10 Apr 2019 13:13:28 +0200
Subject: [PATCH] [WIP] PR90030

---
 gcc/fortran/trans-openmp.c                    | 20 +++++--------
 libgomp/testsuite/libgomp.fortran/pr90030.f90 |  3 ++
 .../libgomp.oacc-fortran/pr90030.f90          | 29 +++++++++++++++++++
 3 files changed, 39 insertions(+), 13 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.fortran/pr90030.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr90030.f90

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 0eb5956cc531..07a63f4a45ce 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1125,7 +1125,7 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
       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);
+      gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
       ptr = build_fold_indirect_ref (ptr);
       OMP_CLAUSE_DECL (c) = ptr;
       c2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
@@ -2081,8 +2081,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
   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);
+      gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
       decl = build_fold_indirect_ref (decl);
     }
   else if (DECL_P (decl))
@@ -2105,8 +2104,7 @@ 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);
-  ptr = fold_convert (build_pointer_type (char_type_node),
-      ptr);
+  gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
   OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
  }
       switch (n->u.depend_op)
@@ -2172,8 +2170,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
     {
       tree type = TREE_TYPE (decl);
       tree ptr = gfc_conv_descriptor_data_get (decl);
-      ptr = fold_convert (build_pointer_type (char_type_node),
-  ptr);
+      gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
       ptr = build_fold_indirect_ref (ptr);
       OMP_CLAUSE_DECL (node) = ptr;
       node2 = build_omp_clause (input_location,
@@ -2266,8 +2263,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
        OMP_CLAUSE_SIZE (node), elemsz);
     }
   gfc_add_block_to_block (block, &se.post);
-  ptr = fold_convert (build_pointer_type (char_type_node),
-      ptr);
+  gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
   OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
 
   if (POINTER_TYPE_P (TREE_TYPE (decl))
@@ -2407,8 +2403,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
     {
       tree type = TREE_TYPE (decl);
       tree ptr = gfc_conv_descriptor_data_get (decl);
-      ptr = fold_convert (build_pointer_type (char_type_node),
-  ptr);
+      gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
       ptr = build_fold_indirect_ref (ptr);
       OMP_CLAUSE_DECL (node) = ptr;
       OMP_CLAUSE_SIZE (node)
@@ -2453,8 +2448,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
        OMP_CLAUSE_SIZE (node), elemsz);
     }
   gfc_add_block_to_block (block, &se.post);
-  ptr = fold_convert (build_pointer_type (char_type_node),
-      ptr);
+  gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
   OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
  }
       omp_clauses = gfc_trans_add_clause (node, omp_clauses);
diff --git a/libgomp/testsuite/libgomp.fortran/pr90030.f90 b/libgomp/testsuite/libgomp.fortran/pr90030.f90
new file mode 100644
index 000000000000..8c2432cb1783
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr90030.f90
@@ -0,0 +1,3 @@
+! { dg-do run }
+
+include '../libgomp.oacc-fortran/pr90030.f90'
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr90030.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr90030.f90
new file mode 100644
index 000000000000..bbfcff3a869f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr90030.f90
@@ -0,0 +1,29 @@
+! PR90030.
+! Test if the array data associated with c is properly aligned
+! on the accelerator.  If it is not, this program will crash.
+
+! This is also included from '../libgomp.fortran/pr90030.f90'.
+
+! { dg-do run }
+
+program routine_align_main
+  implicit none
+  integer :: i, n
+  real*8, dimension(:), allocatable :: c
+
+  n = 10
+
+  allocate (c(n))
+
+  !$omp target map(to: n) map(from: c(1:n))
+  !$acc parallel copyin(n) copyout(c(1:n))
+  do i = 1, n
+     c(i) = i
+  enddo
+  !$acc end parallel
+  !$omp end target
+
+  do i = 1, n
+     if (c(i) .ne. i) stop i
+  enddo
+end program routine_align_main
--
2.17.1