* 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. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@211768 138bc75d-0d04-0410-961f-82ee72b054a4
author: jakub <jakub@138bc75d-0d04-0410-961f-82ee72b054a4> 2014-06-18 07:16:12 +0000
committer: jakub <jakub@138bc75d-0d04-0410-961f-82ee72b054a4> 2014-06-18 07:16:12 +0000
commit: 691447aba6b3f8f3ea8402a35e7c53e2dbb9a6bd (patch)
tree: 20e7887bc99ba095e639c296ace7f2c40b3b9f11
parent: 23096bb05120b4fdca8164e7fa512ae25a369f30 (diff)
download: gcc-691447aba6b3f8f3ea8402a35e7c53e2dbb9a6bd.tar.gz
53 files changed, 3583 insertions, 330 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index ef2dbd01773..0fc67169c8e 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,34 @@
+2014-06-18  Jakub Jelinek  <jakub@redhat.com>
+
+	* gimplify.c (omp_notice_variable): If n is non-NULL
+	and no flags change in ORT_TARGET region, don't jump to
+	do_outer.
+	(struct gimplify_adjust_omp_clauses_data): New type.
+	(gimplify_adjust_omp_clauses_1): Adjust for data being
+	a struct gimplify_adjust_omp_clauses_data pointer instead
+	of tree *.  Pass pre_p as a new argument to
+	lang_hooks.decls.omp_finish_clause hook.
+	(gimplify_adjust_omp_clauses): Add pre_p argument, adjust
+	splay_tree_foreach to pass both list_p and pre_p.
+	(gimplify_omp_parallel, gimplify_omp_task, gimplify_omp_for,
+	gimplify_omp_workshare, gimplify_omp_target_update): Adjust
+	gimplify_adjust_omp_clauses callers.
+	* langhooks.c (lhd_omp_finish_clause): New function.
+	* langhooks-def.h (lhd_omp_finish_clause): New prototype.
+	(LANG_HOOKS_OMP_FINISH_CLAUSE): Define to lhd_omp_finish_clause.
+	* langhooks.h (struct lang_hooks_for_decls): Add a new
+	gimple_seq * argument to omp_finish_clause hook.
+	* omp-low.c (scan_sharing_clauses): Call scan_omp_op on
+	non-DECL_P OMP_CLAUSE_DECL if ctx->outer.
+	(scan_omp_parallel, lower_omp_for): When adding
+	_LOOPTEMP_ clause var, add it to outer ctx's decl_map
+	as identity.
+	* tree-core.h (OMP_CLAUSE_MAP_TO_PSET): New map kind.
+	* tree-nested.c (convert_nonlocal_omp_clauses,
+	convert_local_omp_clauses): Handle various OpenMP 4.0 clauses.
+	* tree-pretty-print.c (dump_omp_clause): Handle
+	OMP_CLAUSE_MAP_TO_PSET.
+
 2014-06-17  Andrew MacLeod  <amacleod@redhat.com>
 
 	* tree-dfa.h (get_addr_base_and_unit_offset_1): Move from here.
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 9705f641893..9b4818ed971 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,9 @@
+2014-06-18  Jakub Jelinek  <jakub@redhat.com>
+
+	* cp-gimplify.c (cxx_omp_finish_clause): Add a gimple_seq *
+	argument.
+	* cp-tree.h (cxx_omp_finish_clause): Adjust prototype.
+
 2014-06-17  Jason Merrill  <jason@redhat.com>
 
 	PR c++/60605
diff --git a/gcc/cp/cp-gimplify.c b/gcc/cp/cp-gimplify.c
index 2798358be6b..296bd5f2443 100644
--- a/gcc/cp/cp-gimplify.c
+++ b/gcc/cp/cp-gimplify.c
@@ -1592,7 +1592,7 @@ cxx_omp_predetermined_sharing (tree decl)
 /* Finalize an implicitly determined clause.  */
 
 void
-cxx_omp_finish_clause (tree c)
+cxx_omp_finish_clause (tree c, gimple_seq *)
 {
   tree decl, inner_type;
   bool make_shared = false;
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 3e4ec3d72d9..71298efe504 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -6228,7 +6228,7 @@ extern tree cxx_omp_clause_default_ctor		(tree, tree, tree);
 extern tree cxx_omp_clause_copy_ctor		(tree, tree, tree);
 extern tree cxx_omp_clause_assign_op		(tree, tree, tree);
 extern tree cxx_omp_clause_dtor			(tree, tree);
-extern void cxx_omp_finish_clause		(tree);
+extern void cxx_omp_finish_clause		(tree, gimple_seq *);
 extern bool cxx_omp_privatize_by_reference	(const_tree);
 
 /* in name-lookup.c */
diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index cbaf2ab30ea..3a5d6aab441 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,177 @@
+2014-06-18  Jakub Jelinek  <jakub@redhat.com>
+
+	* cpp.c (cpp_define_builtins): Change _OPENMP macro to
+	201307.
+	* dump-parse-tree.c (show_omp_namelist): Add list_type
+	argument.  Adjust for rop being u.reduction_op now,
+	handle depend_op or map_op.
+	(show_omp_node): Adjust callers.  Print some new
+	OpenMP 4.0 clauses, adjust for OMP_LIST_DEPEND_{IN,OUT}
+	becoming a single OMP_LIST_DEPEND.
+	* f95-lang.c (gfc_handle_omp_declare_target_attribute): New
+	function.
+	(gfc_attribute_table): New variable.
+	(LANG_HOOKS_OMP_FINISH_CLAUSE, LANG_HOOKS_ATTRIBUTE_TABLE): Redefine.
+	* frontend-passes.c (gfc_code_walker): Handle new OpenMP target
+	EXEC_OMP_* codes and new clauses.
+	* gfortran.h (gfc_statement): Add ST_OMP_TARGET, ST_OMP_END_TARGET,
+	ST_OMP_TARGET_DATA, ST_OMP_END_TARGET_DATA, ST_OMP_TARGET_UPDATE,
+	ST_OMP_DECLARE_TARGET, ST_OMP_TEAMS, ST_OMP_END_TEAMS,
+	ST_OMP_DISTRIBUTE, ST_OMP_END_DISTRIBUTE, ST_OMP_DISTRIBUTE_SIMD,
+	ST_OMP_END_DISTRIBUTE_SIMD, ST_OMP_DISTRIBUTE_PARALLEL_DO,
+	ST_OMP_END_DISTRIBUTE_PARALLEL_DO, ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
+	ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD, ST_OMP_TARGET_TEAMS,
+	ST_OMP_END_TARGET_TEAMS, ST_OMP_TEAMS_DISTRIBUTE,
+	ST_OMP_END_TEAMS_DISTRIBUTE, ST_OMP_TEAMS_DISTRIBUTE_SIMD,
+	ST_OMP_END_TEAMS_DISTRIBUTE_SIMD, ST_OMP_TARGET_TEAMS_DISTRIBUTE,
+	ST_OMP_END_TARGET_TEAMS_DISTRIBUTE,
+	ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+	ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD,
+	ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+	ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO,
+	ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+	ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+	ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+	ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+	ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD and
+	ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD.
+	(symbol_attribute): Add omp_declare_target field.
+	(gfc_omp_depend_op, gfc_omp_map_op): New enums.
+	(gfc_omp_namelist): Replace rop field with union
+	containing reduction_op, depend_op and map_op.
+	(OMP_LIST_DEPEND_IN, OMP_LIST_DEPEND_OUT): Remove.
+	(OMP_LIST_DEPEND, OMP_LIST_MAP, OMP_LIST_TO, OMP_LIST_FROM): New.
+	(gfc_omp_clauses): Add num_teams, device, thread_limit,
+	dist_sched_kind, dist_chunk_size fields.
+	(gfc_common_head): Add omp_declare_target field.
+	(gfc_exec_op): Add EXEC_OMP_TARGET, EXEC_OMP_TARGET_DATA,
+	EXEC_OMP_TEAMS, EXEC_OMP_DISTRIBUTE, EXEC_OMP_DISTRIBUTE_SIMD,
+	EXEC_OMP_DISTRIBUTE_PARALLEL_DO, EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
+	EXEC_OMP_TARGET_TEAMS, EXEC_OMP_TEAMS_DISTRIBUTE,
+	EXEC_OMP_TEAMS_DISTRIBUTE_SIMD, EXEC_OMP_TARGET_TEAMS_DISTRIBUTE,
+	EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+	EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+	EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+	EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+	EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD and
+	EXEC_OMP_TARGET_UPDATE.
+	(gfc_add_omp_declare_target): New prototype.
+	* match.h (gfc_match_omp_declare_target, gfc_match_omp_distribute,
+	gfc_match_omp_distribute_parallel_do,
+	gfc_match_omp_distribute_parallel_do_simd,
+	gfc_match_omp_distribute_simd, gfc_match_omp_target,
+	gfc_match_omp_target_data, gfc_match_omp_target_teams,
+	gfc_match_omp_target_teams_distribute,
+	gfc_match_omp_target_teams_distribute_parallel_do,
+	gfc_match_omp_target_teams_distribute_parallel_do_simd,
+	gfc_match_omp_target_teams_distribute_simd,
+	gfc_match_omp_target_update, gfc_match_omp_teams,
+	gfc_match_omp_teams_distribute,
+	gfc_match_omp_teams_distribute_parallel_do,
+	gfc_match_omp_teams_distribute_parallel_do_simd,
+	gfc_match_omp_teams_distribute_simd): New prototypes.
+	* module.c (ab_attribute): Add AB_OMP_DECLARE_TARGET.
+	(attr_bits): Likewise.
+	(mio_symbol_attribute): Handle omp_declare_target attribute.
+	(gfc_free_omp_clauses): Free num_teams, device, thread_limit
+	and dist_chunk_size expressions.
+	(OMP_CLAUSE_PRIVATE, OMP_CLAUSE_FIRSTPRIVATE, OMP_CLAUSE_LASTPRIVATE,
+	OMP_CLAUSE_COPYPRIVATE, OMP_CLAUSE_SHARED, OMP_CLAUSE_COPYIN,
+	OMP_CLAUSE_REDUCTION, OMP_CLAUSE_IF, OMP_CLAUSE_NUM_THREADS,
+	OMP_CLAUSE_SCHEDULE, OMP_CLAUSE_DEFAULT, OMP_CLAUSE_ORDERED,
+	OMP_CLAUSE_COLLAPSE, OMP_CLAUSE_UNTIED, OMP_CLAUSE_FINAL,
+	OMP_CLAUSE_MERGEABLE, OMP_CLAUSE_ALIGNED, OMP_CLAUSE_DEPEND,
+	OMP_CLAUSE_INBRANCH, OMP_CLAUSE_LINEAR, OMP_CLAUSE_NOTINBRANCH,
+	OMP_CLAUSE_PROC_BIND, OMP_CLAUSE_SAFELEN, OMP_CLAUSE_SIMDLEN,
+	OMP_CLAUSE_UNIFORM): Use 1U instead of 1.
+	(OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP, OMP_CLAUSE_TO, OMP_CLAUSE_FROM,
+	OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_THREAD_LIMIT,
+	OMP_CLAUSE_DIST_SCHEDULE): Define.
+	(gfc_match_omp_clauses): Change mask parameter to unsigned int.
+	Adjust for rop becoming u.reduction_op.  Disallow inbranch with
+	notinbranch.  For depend clause, always create OMP_LIST_DEPEND
+	and fill in u.depend_op.  Handle num_teams, device, map,
+	to, from, thread_limit and dist_schedule clauses.
+	(OMP_DECLARE_SIMD_CLAUSES): Or in OMP_CLAUSE_INBRANCH and
+	OMP_CLAUSE_NOTINBRANCH.
+	(OMP_TARGET_CLAUSES, OMP_TARGET_DATA_CLAUSES,
+	OMP_TARGET_UPDATE_CLAUSES, OMP_TEAMS_CLAUSES,
+	OMP_DISTRIBUTE_CLAUSES): Define.
+	(match_omp): New function.
+	(gfc_match_omp_do, gfc_match_omp_do_simd, gfc_match_omp_parallel,
+	gfc_match_omp_parallel_do, gfc_match_omp_parallel_do_simd,
+	gfc_match_omp_parallel_sections, gfc_match_omp_parallel_workshare,
+	gfc_match_omp_sections, gfc_match_omp_simd, gfc_match_omp_single,
+	gfc_match_omp_task): Rewritten using match_omp.
+	(gfc_match_omp_threadprivate, gfc_match_omp_declare_reduction):
+	Diagnose if the directives are followed by unexpected junk.
+	(gfc_match_omp_distribute, gfc_match_omp_distribute_parallel_do,
+	gfc_match_omp_distribute_parallel_do_simd,
+	gfc_match_omp_distrbute_simd, gfc_match_omp_declare_target,
+	gfc_match_omp_target, gfc_match_omp_target_data,
+	gfc_match_omp_target_teams, gfc_match_omp_target_teams_distribute,
+	gfc_match_omp_target_teams_distribute_parallel_do,
+	gfc_match_omp_target_teams_distribute_parallel_do_simd,
+	gfc_match_omp_target_teams_distrbute_simd, gfc_match_omp_target_update,
+	gfc_match_omp_teams, gfc_match_omp_teams_distribute,
+	gfc_match_omp_teams_distribute_parallel_do,
+	gfc_match_omp_teams_distribute_parallel_do_simd,
+	gfc_match_omp_teams_distrbute_simd): New functions.
+	* openmp.c (resolve_omp_clauses): Adjust for
+	OMP_LIST_DEPEND_{IN,OUT} being changed to OMP_LIST_DEPEND.  Handle
+	OMP_LIST_MAP, OMP_LIST_FROM, OMP_LIST_TO, num_teams, device,
+	dist_chunk_size and thread_limit.
+	(gfc_resolve_omp_parallel_blocks): Only put sharing clauses into
+	ctx.sharing_clauses.  Call gfc_resolve_omp_do_blocks for various
+	new EXEC_OMP_* codes.
+	(resolve_omp_do): Handle various new EXEC_OMP_* codes.
+	(gfc_resolve_omp_directive): Likewise.
+	(gfc_resolve_omp_declare_simd): Add missing space to diagnostics.
+	* parse.c (decode_omp_directive): Handle parsing of OpenMP 4.0
+	offloading related directives.
+	(case_executable): Add ST_OMP_TARGET_UPDATE.
+	(case_exec_markers): Add ST_OMP_TARGET*, ST_OMP_TEAMS*,
+	ST_OMP_DISTRIBUTE*.
+	(case_decl): Add ST_OMP_DECLARE_TARGET.
+	(gfc_ascii_statement): Handle new ST_OMP_* codes.
+	(parse_omp_do): Handle various new ST_OMP_* codes.
+	(parse_executable): Likewise.
+	* resolve.c (gfc_resolve_blocks): Handle various new EXEC_OMP_*
+	codes.
+	(resolve_code): Likewise.
+	(resolve_symbol): Change that !$OMP DECLARE TARGET variables
+	are saved.
+	* st.c (gfc_free_statement): Handle various new EXEC_OMP_* codes.
+	* symbol.c (check_conflict): Check omp_declare_target conflicts.
+	(gfc_add_omp_declare_target): New function.
+	(gfc_copy_attr): Copy omp_declare_target.
+	* trans.c (trans_code): Handle various new EXEC_OMP_* codes.
+	* trans-common.c (build_common_decl): Add "omp declare target"
+	attribute if needed.
+	* trans-decl.c (add_attributes_to_decl): Likewise.
+	* trans.h (gfc_omp_finish_clause): New prototype.
+	* trans-openmp.c (gfc_omp_finish_clause): New function.
+	(gfc_trans_omp_reduction_list): Adjust for rop being renamed
+	to u.reduction_op.
+	(gfc_trans_omp_clauses): Adjust for OMP_LIST_DEPEND_{IN,OUT}
+	change to OMP_LIST_DEPEND and fix up depend handling.
+	Handle OMP_LIST_MAP, OMP_LIST_TO, OMP_LIST_FROM, num_teams,
+	thread_limit, device, dist_chunk_size and dist_sched_kind.
+	(gfc_trans_omp_do): Handle EXEC_OMP_DISTRIBUTE.
+	(GFC_OMP_SPLIT_DISTRIBUTE, GFC_OMP_SPLIT_TEAMS,
+	GFC_OMP_SPLIT_TARGET, GFC_OMP_SPLIT_NUM, GFC_OMP_MASK_DISTRIBUTE,
+	GFC_OMP_MASK_TEAMS, GFC_OMP_MASK_TARGET, GFC_OMP_MASK_NUM): New.
+	(gfc_split_omp_clauses): Handle splitting of clauses for new
+	EXEC_OMP_* codes.
+	(gfc_trans_omp_do_simd): Add pblock argument, adjust for being
+	callable for combined constructs.
+	(gfc_trans_omp_parallel_do, gfc_trans_omp_parallel_do_simd): Likewise.
+	(gfc_trans_omp_distribute, gfc_trans_omp_teams,
+	gfc_trans_omp_target, gfc_trans_omp_target_data,
+	gfc_trans_omp_target_update): New functions.
+	(gfc_trans_omp_directive): Adjust gfc_trans_omp_* callers, handle
+	new EXEC_OMP_* codes.
+
 2014-06-18  Tobias Burnus  <burnus@net-b.de>
 
 	PR fortran/61126
diff --git a/gcc/fortran/cpp.c b/gcc/fortran/cpp.c
index 169599003db..7fb8d160267 100644
--- a/gcc/fortran/cpp.c
+++ b/gcc/fortran/cpp.c
@@ -171,7 +171,7 @@ cpp_define_builtins (cpp_reader *pfile)
   cpp_define (pfile, "_LANGUAGE_FORTRAN=1");
 
   if (gfc_option.gfc_flag_openmp)
-    cpp_define (pfile, "_OPENMP=201107");
+    cpp_define (pfile, "_OPENMP=201307");
 
   /* The defines below are necessary for the TARGET_* macros.
 
diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c
index 5a9611923d8..de942f83819 100644
--- a/gcc/fortran/dump-parse-tree.c
+++ b/gcc/fortran/dump-parse-tree.c
@@ -1016,32 +1016,51 @@ show_code (int level, gfc_code *c)
 }
 
 static void
-show_omp_namelist (gfc_omp_namelist *n)
+show_omp_namelist (int list_type, gfc_omp_namelist *n)
 {
   for (; n; n = n->next)
     {
-      switch (n->rop)
-	{
-	case OMP_REDUCTION_PLUS:
-	case OMP_REDUCTION_TIMES:
-	case OMP_REDUCTION_MINUS:
-	case OMP_REDUCTION_AND:
-	case OMP_REDUCTION_OR:
-	case OMP_REDUCTION_EQV:
-	case OMP_REDUCTION_NEQV:
-	  fprintf (dumpfile, "%s:", gfc_op2string ((gfc_intrinsic_op) n->rop));
-	  break;
-	case OMP_REDUCTION_MAX: fputs ("max:", dumpfile); break;
-	case OMP_REDUCTION_MIN: fputs ("min:", dumpfile); break;
-	case OMP_REDUCTION_IAND: fputs ("iand:", dumpfile); break;
-	case OMP_REDUCTION_IOR: fputs ("ior:", dumpfile); break;
-	case OMP_REDUCTION_IEOR: fputs ("ieor:", dumpfile); break;
-	case OMP_REDUCTION_USER:
-	  if (n->udr)
-	    fprintf (dumpfile, "%s:", n->udr->name);
-	  break;
-	default: break;
-	}
+      if (list_type == OMP_LIST_REDUCTION)
+	switch (n->u.reduction_op)
+	  {
+	  case OMP_REDUCTION_PLUS:
+	  case OMP_REDUCTION_TIMES:
+	  case OMP_REDUCTION_MINUS:
+	  case OMP_REDUCTION_AND:
+	  case OMP_REDUCTION_OR:
+	  case OMP_REDUCTION_EQV:
+	  case OMP_REDUCTION_NEQV:
+	    fprintf (dumpfile, "%s:",
+		     gfc_op2string ((gfc_intrinsic_op) n->u.reduction_op));
+	    break;
+	  case OMP_REDUCTION_MAX: fputs ("max:", dumpfile); break;
+	  case OMP_REDUCTION_MIN: fputs ("min:", dumpfile); break;
+	  case OMP_REDUCTION_IAND: fputs ("iand:", dumpfile); break;
+	  case OMP_REDUCTION_IOR: fputs ("ior:", dumpfile); break;
+	  case OMP_REDUCTION_IEOR: fputs ("ieor:", dumpfile); break;
+	  case OMP_REDUCTION_USER:
+	    if (n->udr)
+	      fprintf (dumpfile, "%s:", n->udr->name);
+	    break;
+	  default: break;
+	  }
+      else if (list_type == OMP_LIST_DEPEND)
+	switch (n->u.depend_op)
+	  {
+	  case OMP_DEPEND_IN: fputs ("in:", dumpfile); break;
+	  case OMP_DEPEND_OUT: fputs ("out:", dumpfile); break;
+	  case OMP_DEPEND_INOUT: fputs ("inout:", dumpfile); break;
+	  default: break;
+	  }
+      else if (list_type == OMP_LIST_MAP)
+	switch (n->u.map_op)
+	  {
+	  case OMP_MAP_ALLOC: fputs ("alloc:", dumpfile); break;
+	  case OMP_MAP_TO: fputs ("to:", dumpfile); break;
+	  case OMP_MAP_FROM: fputs ("from:", dumpfile); break;
+	  case OMP_MAP_TOFROM: fputs ("tofrom:", dumpfile); break;
+	  default: break;
+	  }
       fprintf (dumpfile, "%s", n->sym->name);
       if (n->expr)
 	{
@@ -1117,7 +1136,7 @@ show_omp_node (int level, gfc_code *c)
       if (c->ext.omp_namelist)
 	{
 	  fputs (" (", dumpfile);
-	  show_omp_namelist (c->ext.omp_namelist);
+	  show_omp_namelist (OMP_LIST_NUM, c->ext.omp_namelist);
 	  fputc (')', dumpfile);
 	}
       return;
@@ -1226,18 +1245,12 @@ show_omp_node (int level, gfc_code *c)
 	      case OMP_LIST_ALIGNED: type = "ALIGNED"; break;
 	      case OMP_LIST_LINEAR: type = "LINEAR"; break;
 	      case OMP_LIST_REDUCTION: type = "REDUCTION"; break;
-	      case OMP_LIST_DEPEND_IN:
-		fprintf (dumpfile, " DEPEND(IN:");
-		break;
-	      case OMP_LIST_DEPEND_OUT:
-		fprintf (dumpfile, " DEPEND(OUT:");
-		break;
+	      case OMP_LIST_DEPEND: type = "DEPEND"; break;
 	      default:
 		gcc_unreachable ();
 	      }
-	    if (type)
-	      fprintf (dumpfile, " %s(", type);
-	    show_omp_namelist (omp_clauses->lists[list_type]);
+	    fprintf (dumpfile, " %s(", type);
+	    show_omp_namelist (list_type, omp_clauses->lists[list_type]);
 	    fputc (')', dumpfile);
 	  }
       if (omp_clauses->safelen_expr)
@@ -1269,6 +1282,34 @@ show_omp_node (int level, gfc_code *c)
 	    }
 	  fprintf (dumpfile, " PROC_BIND(%s)", type);
 	}
+      if (omp_clauses->num_teams)
+	{
+	  fputs (" NUM_TEAMS(", dumpfile);
+	  show_expr (omp_clauses->num_teams);
+	  fputc (')', dumpfile);
+	}
+      if (omp_clauses->device)
+	{
+	  fputs (" DEVICE(", dumpfile);
+	  show_expr (omp_clauses->device);
+	  fputc (')', dumpfile);
+	}
+      if (omp_clauses->thread_limit)
+	{
+	  fputs (" THREAD_LIMIT(", dumpfile);
+	  show_expr (omp_clauses->thread_limit);
+	  fputc (')', dumpfile);
+	}
+      if (omp_clauses->dist_sched_kind != OMP_SCHED_NONE)
+	{
+	  fprintf (dumpfile, " DIST_SCHEDULE (static");
+	  if (omp_clauses->dist_chunk_size)
+	    {
+	      fputc (',', dumpfile);
+	      show_expr (omp_clauses->dist_chunk_size);
+	    }
+	  fputc (')', dumpfile);
+	}
     }
   fputc ('\n', dumpfile);
   if (c->op == EXEC_OMP_SECTIONS || c->op == EXEC_OMP_PARALLEL_SECTIONS)
@@ -1296,7 +1337,8 @@ show_omp_node (int level, gfc_code *c)
       if (omp_clauses->lists[OMP_LIST_COPYPRIVATE])
 	{
 	  fputs (" COPYPRIVATE(", dumpfile);
-	  show_omp_namelist (omp_clauses->lists[OMP_LIST_COPYPRIVATE]);
+	  show_omp_namelist (OMP_LIST_COPYPRIVATE,
+			     omp_clauses->lists[OMP_LIST_COPYPRIVATE]);
 	  fputc (')', dumpfile);
 	}
       else if (omp_clauses->nowait)
diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c
index 19621449767..e41f61a9f5d 100644
--- a/gcc/fortran/f95-lang.c
+++ b/gcc/fortran/f95-lang.c
@@ -87,6 +87,24 @@ static alias_set_type gfc_get_alias_set (tree);
 static void gfc_init_ts (void);
 static tree gfc_builtin_function (tree);
 
+/* Handle an "omp declare target" attribute; arguments as in
+   struct attribute_spec.handler.  */
+static tree
+gfc_handle_omp_declare_target_attribute (tree *, tree, tree, int, bool *)
+{
+  return NULL_TREE;
+}
+
+/* Table of valid Fortran attributes.  */
+static const struct attribute_spec gfc_attribute_table[] =
+{
+  /* { name, min_len, max_len, decl_req, type_req, fn_type_req, handler,
+       affects_type_identity } */
+  { "omp declare target", 0, 0, true,  false, false,
+    gfc_handle_omp_declare_target_attribute, false },
+  { NULL,		  0, 0, false, false, false, NULL, false }
+};
+
 #undef LANG_HOOKS_NAME
 #undef LANG_HOOKS_INIT
 #undef LANG_HOOKS_FINISH
@@ -109,6 +127,7 @@ static tree gfc_builtin_function (tree);
 #undef LANG_HOOKS_OMP_CLAUSE_COPY_CTOR
 #undef LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP
 #undef LANG_HOOKS_OMP_CLAUSE_DTOR
+#undef LANG_HOOKS_OMP_FINISH_CLAUSE
 #undef LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR
 #undef LANG_HOOKS_OMP_PRIVATE_DEBUG_CLAUSE
 #undef LANG_HOOKS_OMP_PRIVATE_OUTER_REF
@@ -116,6 +135,7 @@ static tree gfc_builtin_function (tree);
 #undef LANG_HOOKS_BUILTIN_FUNCTION
 #undef LANG_HOOKS_BUILTIN_FUNCTION
 #undef LANG_HOOKS_GET_ARRAY_DESCR_INFO
+#undef LANG_HOOKS_ATTRIBUTE_TABLE
 
 /* Define lang hooks.  */
 #define LANG_HOOKS_NAME                 "GNU Fortran"
@@ -139,13 +159,15 @@ static tree gfc_builtin_function (tree);
 #define LANG_HOOKS_OMP_CLAUSE_COPY_CTOR		gfc_omp_clause_copy_ctor
 #define LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP		gfc_omp_clause_assign_op
 #define LANG_HOOKS_OMP_CLAUSE_DTOR		gfc_omp_clause_dtor
+#define LANG_HOOKS_OMP_FINISH_CLAUSE		gfc_omp_finish_clause
 #define LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR	gfc_omp_disregard_value_expr
 #define LANG_HOOKS_OMP_PRIVATE_DEBUG_CLAUSE	gfc_omp_private_debug_clause
 #define LANG_HOOKS_OMP_PRIVATE_OUTER_REF	gfc_omp_private_outer_ref
 #define LANG_HOOKS_OMP_FIRSTPRIVATIZE_TYPE_SIZES \
   gfc_omp_firstprivatize_type_sizes
-#define LANG_HOOKS_BUILTIN_FUNCTION          gfc_builtin_function
-#define LANG_HOOKS_GET_ARRAY_DESCR_INFO	     gfc_get_array_descr_info
+#define LANG_HOOKS_BUILTIN_FUNCTION	gfc_builtin_function
+#define LANG_HOOKS_GET_ARRAY_DESCR_INFO	gfc_get_array_descr_info
+#define LANG_HOOKS_ATTRIBUTE_TABLE	gfc_attribute_table
 
 struct lang_hooks lang_hooks = LANG_HOOKS_INITIALIZER;
 
diff --git a/gcc/fortran/frontend-passes.c b/gcc/fortran/frontend-passes.c
index c69bd0cf179..4646cc33fd3 100644
--- a/gcc/fortran/frontend-passes.c
+++ b/gcc/fortran/frontend-passes.c
@@ -2147,14 +2147,31 @@ gfc_code_walker (gfc_code **c, walk_code_fn_t codefn, walk_expr_fn_t exprfn,
 	      in_omp_workshare = true;
 
 	      /* Fall through  */
-	      
+
+	    case EXEC_OMP_DISTRIBUTE:
+	    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+	    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+	    case EXEC_OMP_DISTRIBUTE_SIMD:
 	    case EXEC_OMP_DO:
 	    case EXEC_OMP_DO_SIMD:
 	    case EXEC_OMP_SECTIONS:
 	    case EXEC_OMP_SINGLE:
 	    case EXEC_OMP_END_SINGLE:
 	    case EXEC_OMP_SIMD:
+	    case EXEC_OMP_TARGET:
+	    case EXEC_OMP_TARGET_DATA:
+	    case EXEC_OMP_TARGET_TEAMS:
+	    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+	    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+	    case EXEC_OMP_TARGET_UPDATE:
 	    case EXEC_OMP_TASK:
+	    case EXEC_OMP_TEAMS:
+	    case EXEC_OMP_TEAMS_DISTRIBUTE:
+	    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
 
 	      /* Come to this label only from the
 		 EXEC_OMP_PARALLEL_* cases above.  */
@@ -2163,28 +2180,28 @@ gfc_code_walker (gfc_code **c, walk_code_fn_t codefn, walk_expr_fn_t exprfn,
 
 	      if (co->ext.omp_clauses)
 		{
+		  gfc_omp_namelist *n;
+		  static int list_types[]
+		    = { OMP_LIST_ALIGNED, OMP_LIST_LINEAR, OMP_LIST_DEPEND,
+			OMP_LIST_MAP, OMP_LIST_TO, OMP_LIST_FROM };
+		  size_t idx;
 		  WALK_SUBEXPR (co->ext.omp_clauses->if_expr);
 		  WALK_SUBEXPR (co->ext.omp_clauses->final_expr);
 		  WALK_SUBEXPR (co->ext.omp_clauses->num_threads);
 		  WALK_SUBEXPR (co->ext.omp_clauses->chunk_size);
 		  WALK_SUBEXPR (co->ext.omp_clauses->safelen_expr);
 		  WALK_SUBEXPR (co->ext.omp_clauses->simdlen_expr);
+		  WALK_SUBEXPR (co->ext.omp_clauses->num_teams);
+		  WALK_SUBEXPR (co->ext.omp_clauses->device);
+		  WALK_SUBEXPR (co->ext.omp_clauses->thread_limit);
+		  WALK_SUBEXPR (co->ext.omp_clauses->dist_chunk_size);
+		  for (idx = 0;
+		       idx < sizeof (list_types) / sizeof (list_types[0]);
+		       idx++)
+		    for (n = co->ext.omp_clauses->lists[list_types[idx]];
+			 n; n = n->next)
+		      WALK_SUBEXPR (n->expr);
 		}
-	      {
-		gfc_omp_namelist *n;
-		for (n = co->ext.omp_clauses->lists[OMP_LIST_ALIGNED];
-		     n; n = n->next)
-		  WALK_SUBEXPR (n->expr);
-		for (n = co->ext.omp_clauses->lists[OMP_LIST_LINEAR];
-		     n; n = n->next)
-		  WALK_SUBEXPR (n->expr);
-		for (n = co->ext.omp_clauses->lists[OMP_LIST_DEPEND_IN];
-		     n; n = n->next)
-		  WALK_SUBEXPR (n->expr);
-		for (n = co->ext.omp_clauses->lists[OMP_LIST_DEPEND_OUT];
-		     n; n = n->next)
-		  WALK_SUBEXPR (n->expr);
-	      }
 	      break;
 	    default:
 	      break;
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 1df79fdbe05..a11ca3d704f 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -215,6 +215,24 @@ typedef enum
   ST_OMP_TASKGROUP, ST_OMP_END_TASKGROUP, ST_OMP_SIMD, ST_OMP_END_SIMD,
   ST_OMP_DO_SIMD, ST_OMP_END_DO_SIMD, ST_OMP_PARALLEL_DO_SIMD,
   ST_OMP_END_PARALLEL_DO_SIMD, ST_OMP_DECLARE_SIMD, ST_OMP_DECLARE_REDUCTION,
+  ST_OMP_TARGET, ST_OMP_END_TARGET, ST_OMP_TARGET_DATA, ST_OMP_END_TARGET_DATA,
+  ST_OMP_TARGET_UPDATE, ST_OMP_DECLARE_TARGET,
+  ST_OMP_TEAMS, ST_OMP_END_TEAMS, ST_OMP_DISTRIBUTE, ST_OMP_END_DISTRIBUTE,
+  ST_OMP_DISTRIBUTE_SIMD, ST_OMP_END_DISTRIBUTE_SIMD,
+  ST_OMP_DISTRIBUTE_PARALLEL_DO, ST_OMP_END_DISTRIBUTE_PARALLEL_DO,
+  ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD, ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD,
+  ST_OMP_TARGET_TEAMS, ST_OMP_END_TARGET_TEAMS, ST_OMP_TEAMS_DISTRIBUTE,
+  ST_OMP_END_TEAMS_DISTRIBUTE, ST_OMP_TEAMS_DISTRIBUTE_SIMD,
+  ST_OMP_END_TEAMS_DISTRIBUTE_SIMD, ST_OMP_TARGET_TEAMS_DISTRIBUTE,
+  ST_OMP_END_TARGET_TEAMS_DISTRIBUTE, ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+  ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD, ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+  ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO,
+  ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+  ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+  ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+  ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+  ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+  ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
   ST_PROCEDURE, ST_GENERIC, ST_CRITICAL, ST_END_CRITICAL,
   ST_GET_FCN_CHARACTERISTICS, ST_LOCK, ST_UNLOCK, ST_NONE
 }
@@ -821,6 +839,9 @@ typedef struct
      !$OMP DECLARE REDUCTION.  */
   unsigned omp_udr_artificial_var:1;
 
+  /* Mentioned in OMP DECLARE TARGET.  */
+  unsigned omp_declare_target:1;
+
   /* Attributes set by compiler extensions (!GCC$ ATTRIBUTES).  */
   unsigned ext_attr:EXT_ATTR_NUM;
 
@@ -1060,6 +1081,23 @@ typedef enum
 }
 gfc_omp_reduction_op;
 
+typedef enum
+{
+  OMP_DEPEND_IN,
+  OMP_DEPEND_OUT,
+  OMP_DEPEND_INOUT
+}
+gfc_omp_depend_op;
+
+typedef enum
+{
+  OMP_MAP_ALLOC,
+  OMP_MAP_TO,
+  OMP_MAP_FROM,
+  OMP_MAP_TOFROM
+}
+gfc_omp_map_op;
+
 /* For use in OpenMP clauses in case we need extra information
    (aligned clause alignment, linear clause step, etc.).  */
 
@@ -1067,7 +1105,12 @@ typedef struct gfc_omp_namelist
 {
   struct gfc_symbol *sym;
   struct gfc_expr *expr;
-  gfc_omp_reduction_op rop;
+  union
+    {
+      gfc_omp_reduction_op reduction_op;
+      gfc_omp_depend_op depend_op;
+      gfc_omp_map_op map_op;
+    } u;
   struct gfc_omp_udr *udr;
   struct gfc_omp_namelist *next;
 }
@@ -1086,8 +1129,10 @@ enum
   OMP_LIST_UNIFORM,
   OMP_LIST_ALIGNED,
   OMP_LIST_LINEAR,
-  OMP_LIST_DEPEND_IN,
-  OMP_LIST_DEPEND_OUT,
+  OMP_LIST_DEPEND,
+  OMP_LIST_MAP,
+  OMP_LIST_TO,
+  OMP_LIST_FROM,
   OMP_LIST_REDUCTION,
   OMP_LIST_NUM
 };
@@ -1147,6 +1192,11 @@ typedef struct gfc_omp_clauses
   enum gfc_omp_proc_bind_kind proc_bind;
   struct gfc_expr *safelen_expr;
   struct gfc_expr *simdlen_expr;
+  struct gfc_expr *num_teams;
+  struct gfc_expr *device;
+  struct gfc_expr *thread_limit;
+  enum gfc_omp_sched_kind dist_sched_kind;
+  struct gfc_expr *dist_chunk_size;
 }
 gfc_omp_clauses;
 
@@ -1387,7 +1437,7 @@ struct gfc_undo_change_set
 typedef struct gfc_common_head
 {
   locus where;
-  char use_assoc, saved, threadprivate;
+  char use_assoc, saved, threadprivate, omp_declare_target;
   char name[GFC_MAX_SYMBOL_LEN + 1];
   struct gfc_symbol *head;
   const char* binding_label;
@@ -2217,7 +2267,17 @@ typedef enum
   EXEC_OMP_END_SINGLE, EXEC_OMP_TASK, EXEC_OMP_TASKWAIT,
   EXEC_OMP_TASKYIELD, EXEC_OMP_CANCEL, EXEC_OMP_CANCELLATION_POINT,
   EXEC_OMP_TASKGROUP, EXEC_OMP_SIMD, EXEC_OMP_DO_SIMD,
-  EXEC_OMP_PARALLEL_DO_SIMD
+  EXEC_OMP_PARALLEL_DO_SIMD, EXEC_OMP_TARGET, EXEC_OMP_TARGET_DATA,
+  EXEC_OMP_TEAMS, EXEC_OMP_DISTRIBUTE, EXEC_OMP_DISTRIBUTE_SIMD,
+  EXEC_OMP_DISTRIBUTE_PARALLEL_DO, EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
+  EXEC_OMP_TARGET_TEAMS, EXEC_OMP_TEAMS_DISTRIBUTE,
+  EXEC_OMP_TEAMS_DISTRIBUTE_SIMD, EXEC_OMP_TARGET_TEAMS_DISTRIBUTE,
+  EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+  EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+  EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+  EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+  EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+  EXEC_OMP_TARGET_UPDATE
 }
 gfc_exec_op;
 
@@ -2682,6 +2742,7 @@ bool gfc_add_protected (symbol_attribute *, const char *, locus *);
 bool gfc_add_result (symbol_attribute *, const char *, locus *);
 bool gfc_add_save (symbol_attribute *, save_state, const char *, locus *);
 bool gfc_add_threadprivate (symbol_attribute *, const char *, locus *);
+bool gfc_add_omp_declare_target (symbol_attribute *, const char *, locus *);
 bool gfc_add_saved_common (symbol_attribute *, locus *);
 bool gfc_add_target (symbol_attribute *, locus *);
 bool gfc_add_dummy (symbol_attribute *, const char *, locus *);
diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h
index 86c2d1bd46d..d07db11ef30 100644
--- a/gcc/fortran/match.h
+++ b/gcc/fortran/match.h
@@ -131,6 +131,11 @@ match gfc_match_omp_cancellation_point (void);
 match gfc_match_omp_critical (void);
 match gfc_match_omp_declare_reduction (void);
 match gfc_match_omp_declare_simd (void);
+match gfc_match_omp_declare_target (void);
+match gfc_match_omp_distribute (void);
+match gfc_match_omp_distribute_parallel_do (void);
+match gfc_match_omp_distribute_parallel_do_simd (void);
+match gfc_match_omp_distribute_simd (void);
 match gfc_match_omp_do (void);
 match gfc_match_omp_do_simd (void);
 match gfc_match_omp_flush (void);
@@ -144,10 +149,23 @@ match gfc_match_omp_parallel_workshare (void);
 match gfc_match_omp_sections (void);
 match gfc_match_omp_simd (void);
 match gfc_match_omp_single (void);
+match gfc_match_omp_target (void);
+match gfc_match_omp_target_data (void);
+match gfc_match_omp_target_teams (void);
+match gfc_match_omp_target_teams_distribute (void);
+match gfc_match_omp_target_teams_distribute_parallel_do (void);
+match gfc_match_omp_target_teams_distribute_parallel_do_simd (void);
+match gfc_match_omp_target_teams_distribute_simd (void);
+match gfc_match_omp_target_update (void);
 match gfc_match_omp_task (void);
 match gfc_match_omp_taskgroup (void);
 match gfc_match_omp_taskwait (void);
 match gfc_match_omp_taskyield (void);
+match gfc_match_omp_teams (void);
+match gfc_match_omp_teams_distribute (void);
+match gfc_match_omp_teams_distribute_parallel_do (void);
+match gfc_match_omp_teams_distribute_parallel_do_simd (void);
+match gfc_match_omp_teams_distribute_simd (void);
 match gfc_match_omp_threadprivate (void);
 match gfc_match_omp_workshare (void);
 match gfc_match_omp_end_nowait (void);
diff --git a/gcc/fortran/module.c b/gcc/fortran/module.c
index 261c9044332..bdd9961652d 100644
--- a/gcc/fortran/module.c
+++ b/gcc/fortran/module.c
@@ -1877,7 +1877,7 @@ typedef enum
   AB_IS_BIND_C, AB_IS_C_INTEROP, AB_IS_ISO_C, AB_ABSTRACT, AB_ZERO_COMP,
   AB_IS_CLASS, AB_PROCEDURE, AB_PROC_POINTER, AB_ASYNCHRONOUS, AB_CODIMENSION,
   AB_COARRAY_COMP, AB_VTYPE, AB_VTAB, AB_CONTIGUOUS, AB_CLASS_POINTER,
-  AB_IMPLICIT_PURE, AB_ARTIFICIAL, AB_UNLIMITED_POLY
+  AB_IMPLICIT_PURE, AB_ARTIFICIAL, AB_UNLIMITED_POLY, AB_OMP_DECLARE_TARGET
 }
 ab_attribute;
 
@@ -1932,6 +1932,7 @@ static const mstring attr_bits[] =
     minit ("CLASS_POINTER", AB_CLASS_POINTER),
     minit ("IMPLICIT_PURE", AB_IMPLICIT_PURE),
     minit ("UNLIMITED_POLY", AB_UNLIMITED_POLY),
+    minit ("OMP_DECLARE_TARGET", AB_OMP_DECLARE_TARGET),
     minit (NULL, -1)
 };
 
@@ -2110,6 +2111,8 @@ mio_symbol_attribute (symbol_attribute *attr)
 	MIO_NAME (ab_attribute) (AB_VTYPE, attr_bits);
       if (attr->vtab)
 	MIO_NAME (ab_attribute) (AB_VTAB, attr_bits);
+      if (attr->omp_declare_target)
+	MIO_NAME (ab_attribute) (AB_OMP_DECLARE_TARGET, attr_bits);
 
       mio_rparen ();
 
@@ -2273,6 +2276,9 @@ mio_symbol_attribute (symbol_attribute *attr)
 	    case AB_VTAB:
 	      attr->vtab = 1;
 	      break;
+	    case AB_OMP_DECLARE_TARGET:
+	      attr->omp_declare_target = 1;
+	      break;
 	    }
 	}
     }
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index a6e5f6c2cf8..266ac3d9a9d 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -72,6 +72,10 @@ gfc_free_omp_clauses (gfc_omp_clauses *c)
   gfc_free_expr (c->chunk_size);
   gfc_free_expr (c->safelen_expr);
   gfc_free_expr (c->simdlen_expr);
+  gfc_free_expr (c->num_teams);
+  gfc_free_expr (c->device);
+  gfc_free_expr (c->thread_limit);
+  gfc_free_expr (c->dist_chunk_size);
   for (i = 0; i < OMP_LIST_NUM; i++)
     gfc_free_omp_namelist (c->lists[i]);
   free (c);
@@ -283,38 +287,45 @@ cleanup:
   return MATCH_ERROR;
 }
 
-#define OMP_CLAUSE_PRIVATE	(1 << 0)
-#define OMP_CLAUSE_FIRSTPRIVATE	(1 << 1)
-#define OMP_CLAUSE_LASTPRIVATE	(1 << 2)
-#define OMP_CLAUSE_COPYPRIVATE	(1 << 3)
-#define OMP_CLAUSE_SHARED	(1 << 4)
-#define OMP_CLAUSE_COPYIN	(1 << 5)
-#define OMP_CLAUSE_REDUCTION	(1 << 6)
-#define OMP_CLAUSE_IF		(1 << 7)
-#define OMP_CLAUSE_NUM_THREADS	(1 << 8)
-#define OMP_CLAUSE_SCHEDULE	(1 << 9)
-#define OMP_CLAUSE_DEFAULT	(1 << 10)
-#define OMP_CLAUSE_ORDERED	(1 << 11)
-#define OMP_CLAUSE_COLLAPSE	(1 << 12)
-#define OMP_CLAUSE_UNTIED	(1 << 13)
-#define OMP_CLAUSE_FINAL	(1 << 14)
-#define OMP_CLAUSE_MERGEABLE	(1 << 15)
-#define OMP_CLAUSE_ALIGNED	(1 << 16)
-#define OMP_CLAUSE_DEPEND	(1 << 17)
-#define OMP_CLAUSE_INBRANCH	(1 << 18)
-#define OMP_CLAUSE_LINEAR	(1 << 19)
-#define OMP_CLAUSE_NOTINBRANCH	(1 << 20)
-#define OMP_CLAUSE_PROC_BIND	(1 << 21)
-#define OMP_CLAUSE_SAFELEN	(1 << 22)
-#define OMP_CLAUSE_SIMDLEN	(1 << 23)
-#define OMP_CLAUSE_UNIFORM	(1 << 24)
+#define OMP_CLAUSE_PRIVATE	(1U << 0)
+#define OMP_CLAUSE_FIRSTPRIVATE	(1U << 1)
+#define OMP_CLAUSE_LASTPRIVATE	(1U << 2)
+#define OMP_CLAUSE_COPYPRIVATE	(1U << 3)
+#define OMP_CLAUSE_SHARED	(1U << 4)
+#define OMP_CLAUSE_COPYIN	(1U << 5)
+#define OMP_CLAUSE_REDUCTION	(1U << 6)
+#define OMP_CLAUSE_IF		(1U << 7)
+#define OMP_CLAUSE_NUM_THREADS	(1U << 8)
+#define OMP_CLAUSE_SCHEDULE	(1U << 9)
+#define OMP_CLAUSE_DEFAULT	(1U << 10)
+#define OMP_CLAUSE_ORDERED	(1U << 11)
+#define OMP_CLAUSE_COLLAPSE	(1U << 12)
+#define OMP_CLAUSE_UNTIED	(1U << 13)
+#define OMP_CLAUSE_FINAL	(1U << 14)
+#define OMP_CLAUSE_MERGEABLE	(1U << 15)
+#define OMP_CLAUSE_ALIGNED	(1U << 16)
+#define OMP_CLAUSE_DEPEND	(1U << 17)
+#define OMP_CLAUSE_INBRANCH	(1U << 18)
+#define OMP_CLAUSE_LINEAR	(1U << 19)
+#define OMP_CLAUSE_NOTINBRANCH	(1U << 20)
+#define OMP_CLAUSE_PROC_BIND	(1U << 21)
+#define OMP_CLAUSE_SAFELEN	(1U << 22)
+#define OMP_CLAUSE_SIMDLEN	(1U << 23)
+#define OMP_CLAUSE_UNIFORM	(1U << 24)
+#define OMP_CLAUSE_DEVICE	(1U << 25)
+#define OMP_CLAUSE_MAP		(1U << 26)
+#define OMP_CLAUSE_TO		(1U << 27)
+#define OMP_CLAUSE_FROM		(1U << 28)
+#define OMP_CLAUSE_NUM_TEAMS	(1U << 29)
+#define OMP_CLAUSE_THREAD_LIMIT	(1U << 30)
+#define OMP_CLAUSE_DIST_SCHEDULE	(1U << 31)
 
 /* Match OpenMP directive clauses. MASK is a bitmask of
    clauses that are allowed for a particular directive.  */
 
 static match
-gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true,
-		       bool needs_space = true)
+gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned int mask,
+		       bool first = true, bool needs_space = true)
 {
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
   locus old_loc;
@@ -474,7 +485,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true,
 	      else
 		for (n = *head; n; n = n->next)
 		  {
-		    n->rop = rop;
+		    n->u.reduction_op = rop;
 		    n->udr = udr;
 		  }
 	      continue;
@@ -570,13 +581,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true,
 	      continue;
 	    }
 	}
-      if ((mask & OMP_CLAUSE_INBRANCH) && !c->inbranch
+      if ((mask & OMP_CLAUSE_INBRANCH) && !c->inbranch && !c->notinbranch
 	  && gfc_match ("inbranch") == MATCH_YES)
 	{
 	  c->inbranch = needs_space = true;
 	  continue;
 	}
-      if ((mask & OMP_CLAUSE_NOTINBRANCH) && !c->notinbranch
+      if ((mask & OMP_CLAUSE_NOTINBRANCH) && !c->notinbranch && !c->inbranch
 	  && gfc_match ("notinbranch") == MATCH_YES)
 	{
 	  c->notinbranch = needs_space = true;
@@ -662,21 +673,94 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true,
 	  continue;
 	}
       if ((mask & OMP_CLAUSE_DEPEND)
-	  && gfc_match_omp_variable_list ("depend ( in : ",
-					  &c->lists[OMP_LIST_DEPEND_IN], false,
-					  NULL, NULL, true)
-	     == MATCH_YES)
+	  && gfc_match ("depend ( ") == MATCH_YES)
+	{
+	  match m = MATCH_YES;
+	  gfc_omp_depend_op depend_op = OMP_DEPEND_OUT;
+	  if (gfc_match ("inout") == MATCH_YES)
+	    depend_op = OMP_DEPEND_INOUT;
+	  else if (gfc_match ("in") == MATCH_YES)
+	    depend_op = OMP_DEPEND_IN;
+	  else if (gfc_match ("out") == MATCH_YES)
+	    depend_op = OMP_DEPEND_OUT;
+	  else
+	    m = MATCH_NO;
+	  head = NULL;
+	  if (m == MATCH_YES
+	      && gfc_match_omp_variable_list (" : ",
+					      &c->lists[OMP_LIST_DEPEND],
+					      false, NULL, &head, true)
+		 == MATCH_YES)
+	    {
+	      gfc_omp_namelist *n;
+	      for (n = *head; n; n = n->next)
+		n->u.depend_op = depend_op;
+	      continue;
+	    }
+	  else
+	    gfc_current_locus = old_loc;
+	}
+      if ((mask & OMP_CLAUSE_DIST_SCHEDULE)
+	  && c->dist_sched_kind == OMP_SCHED_NONE
+	  && gfc_match ("dist_schedule ( static") == MATCH_YES)
+	{
+	  match m = MATCH_NO;
+	  c->dist_sched_kind = OMP_SCHED_STATIC;
+	  m = gfc_match (" , %e )", &c->dist_chunk_size);
+	  if (m != MATCH_YES)
+	    m = gfc_match_char (')');
+	  if (m != MATCH_YES)
+	    {
+	      c->dist_sched_kind = OMP_SCHED_NONE;
+	      gfc_current_locus = old_loc;
+	    }
+	  else
+	    continue;
+	}
+      if ((mask & OMP_CLAUSE_NUM_TEAMS) && c->num_teams == NULL
+	  && gfc_match ("num_teams ( %e )", &c->num_teams) == MATCH_YES)
 	continue;
-      if ((mask & OMP_CLAUSE_DEPEND)
-	  && gfc_match_omp_variable_list ("depend ( out : ",
-					  &c->lists[OMP_LIST_DEPEND_OUT], false,
-					  NULL, NULL, true)
+      if ((mask & OMP_CLAUSE_DEVICE) && c->device == NULL
+	  && gfc_match ("device ( %e )", &c->device) == MATCH_YES)
+	continue;
+      if ((mask & OMP_CLAUSE_THREAD_LIMIT) && c->thread_limit == NULL
+	  && gfc_match ("thread_limit ( %e )", &c->thread_limit) == MATCH_YES)
+	continue;
+      if ((mask & OMP_CLAUSE_MAP)
+	  && gfc_match ("map ( ") == MATCH_YES)
+	{
+	  gfc_omp_map_op map_op = OMP_MAP_TOFROM;
+	  if (gfc_match ("alloc : ") == MATCH_YES)
+	    map_op = OMP_MAP_ALLOC;
+	  else if (gfc_match ("tofrom : ") == MATCH_YES)
+	    map_op = OMP_MAP_TOFROM;
+	  else if (gfc_match ("to : ") == MATCH_YES)
+	    map_op = OMP_MAP_TO;
+	  else if (gfc_match ("from : ") == MATCH_YES)
+	    map_op = OMP_MAP_FROM;
+	  head = NULL;
+	  if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP],
+					   false, NULL, &head, true)
+	      == MATCH_YES)
+	    {
+	      gfc_omp_namelist *n;
+	      for (n = *head; n; n = n->next)
+		n->u.map_op = map_op;
+	      continue;
+	    }
+	  else
+	    gfc_current_locus = old_loc;
+	}
+      if ((mask & OMP_CLAUSE_TO)
+	  && gfc_match_omp_variable_list ("to (",
+					  &c->lists[OMP_LIST_TO], false,
+					  NULL, &head, true)
 	     == MATCH_YES)
 	continue;
-      if ((mask & OMP_CLAUSE_DEPEND)
-	  && gfc_match_omp_variable_list ("depend ( inout : ",
-					  &c->lists[OMP_LIST_DEPEND_OUT], false,
-					  NULL, NULL, true)
+      if ((mask & OMP_CLAUSE_FROM)
+	  && gfc_match_omp_variable_list ("from (",
+					  &c->lists[OMP_LIST_FROM], false,
+					  NULL, &head, true)
 	     == MATCH_YES)
 	continue;
 
@@ -699,7 +783,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true,
    | OMP_CLAUSE_NUM_THREADS | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PROC_BIND)
 #define OMP_DECLARE_SIMD_CLAUSES \
   (OMP_CLAUSE_SIMDLEN | OMP_CLAUSE_LINEAR | OMP_CLAUSE_UNIFORM		\
-   | OMP_CLAUSE_ALIGNED)
+   | OMP_CLAUSE_ALIGNED | OMP_CLAUSE_INBRANCH | OMP_CLAUSE_NOTINBRANCH)
 #define OMP_DO_CLAUSES \
   (OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE				\
    | OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_REDUCTION			\
@@ -715,100 +799,97 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, int mask, bool first = true,
   (OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_SHARED	\
    | OMP_CLAUSE_IF | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED		\
    | OMP_CLAUSE_FINAL | OMP_CLAUSE_MERGEABLE | OMP_CLAUSE_DEPEND)
-
-match
-gfc_match_omp_parallel (void)
-{
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_PARALLEL_CLAUSES) != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_PARALLEL;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
-}
+#define OMP_TARGET_CLAUSES \
+  (OMP_CLAUSE_DEVICE | OMP_CLAUSE_MAP | OMP_CLAUSE_IF)
+#define OMP_TARGET_DATA_CLAUSES \
+  (OMP_CLAUSE_DEVICE | OMP_CLAUSE_MAP | OMP_CLAUSE_IF)
+#define OMP_TARGET_UPDATE_CLAUSES \
+  (OMP_CLAUSE_DEVICE | OMP_CLAUSE_IF | OMP_CLAUSE_TO | OMP_CLAUSE_FROM)
+#define OMP_TEAMS_CLAUSES \
+  (OMP_CLAUSE_NUM_TEAMS | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_DEFAULT	\
+   | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_SHARED	\
+   | OMP_CLAUSE_REDUCTION)
+#define OMP_DISTRIBUTE_CLAUSES \
+  (OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_COLLAPSE	\
+   | OMP_CLAUSE_DIST_SCHEDULE)
 
 
-match
-gfc_match_omp_task (void)
+static match
+match_omp (gfc_exec_op op, unsigned int mask)
 {
   gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_TASK_CLAUSES) != MATCH_YES)
+  if (gfc_match_omp_clauses (&c, mask) != MATCH_YES)
     return MATCH_ERROR;
-  new_st.op = EXEC_OMP_TASK;
+  new_st.op = op;
   new_st.ext.omp_clauses = c;
   return MATCH_YES;
 }
 
 
 match
-gfc_match_omp_taskwait (void)
+gfc_match_omp_critical (void)
 {
+  char n[GFC_MAX_SYMBOL_LEN+1];
+
+  if (gfc_match (" ( %n )", n) != MATCH_YES)
+    n[0] = '\0';
   if (gfc_match_omp_eos () != MATCH_YES)
     {
-      gfc_error ("Unexpected junk after TASKWAIT clause at %C");
+      gfc_error ("Unexpected junk after $OMP CRITICAL statement at %C");
       return MATCH_ERROR;
     }
-  new_st.op = EXEC_OMP_TASKWAIT;
-  new_st.ext.omp_clauses = NULL;
+  new_st.op = EXEC_OMP_CRITICAL;
+  new_st.ext.omp_name = n[0] ? xstrdup (n) : NULL;
   return MATCH_YES;
 }
 
 
 match
-gfc_match_omp_taskyield (void)
+gfc_match_omp_distribute (void)
 {
-  if (gfc_match_omp_eos () != MATCH_YES)
-    {
-      gfc_error ("Unexpected junk after TASKYIELD clause at %C");
-      return MATCH_ERROR;
-    }
-  new_st.op = EXEC_OMP_TASKYIELD;
-  new_st.ext.omp_clauses = NULL;
-  return MATCH_YES;
+  return match_omp (EXEC_OMP_DISTRIBUTE, OMP_DISTRIBUTE_CLAUSES);
 }
 
 
 match
-gfc_match_omp_critical (void)
+gfc_match_omp_distribute_parallel_do (void)
 {
-  char n[GFC_MAX_SYMBOL_LEN+1];
+  return match_omp (EXEC_OMP_DISTRIBUTE_PARALLEL_DO,
+		    OMP_DISTRIBUTE_CLAUSES | OMP_PARALLEL_CLAUSES
+		    | OMP_DO_CLAUSES);
+}
 
-  if (gfc_match (" ( %n )", n) != MATCH_YES)
-    n[0] = '\0';
-  if (gfc_match_omp_eos () != MATCH_YES)
-    {
-      gfc_error ("Unexpected junk after $OMP CRITICAL statement at %C");
-      return MATCH_ERROR;
-    }
-  new_st.op = EXEC_OMP_CRITICAL;
-  new_st.ext.omp_name = n[0] ? xstrdup (n) : NULL;
-  return MATCH_YES;
+
+match
+gfc_match_omp_distribute_parallel_do_simd (void)
+{
+  return match_omp (EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD,
+		    (OMP_DISTRIBUTE_CLAUSES | OMP_PARALLEL_CLAUSES
+		     | OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
+		    & ~OMP_CLAUSE_ORDERED);
+}
+
+
+match
+gfc_match_omp_distribute_simd (void)
+{
+  return match_omp (EXEC_OMP_DISTRIBUTE_SIMD,
+		    OMP_DISTRIBUTE_CLAUSES | OMP_SIMD_CLAUSES);
 }
 
 
 match
 gfc_match_omp_do (void)
 {
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_DO_CLAUSES) != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_DO;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
+  return match_omp (EXEC_OMP_DO, OMP_DO_CLAUSES);
 }
 
 
 match
 gfc_match_omp_do_simd (void)
 {
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, ((OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
-				  & ~OMP_CLAUSE_ORDERED))
-      != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_DO_SIMD;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
+  return match_omp (EXEC_OMP_DO_SIMD, ((OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
+				       & ~OMP_CLAUSE_ORDERED));
 }
 
 
@@ -830,18 +911,6 @@ gfc_match_omp_flush (void)
 
 
 match
-gfc_match_omp_simd (void)
-{
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_SIMD_CLAUSES) != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_SIMD;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
-}
-
-
-match
 gfc_match_omp_declare_simd (void)
 {
   locus where = gfc_current_locus;
@@ -1235,6 +1304,13 @@ gfc_match_omp_declare_reduction (void)
   if (end_loc_set)
     {
       gfc_current_locus = end_loc;
+      if (gfc_match_omp_eos () != MATCH_YES)
+	{
+	  gfc_error ("Unexpected junk after !$OMP DECLARE REDUCTION at %C");
+	  gfc_current_locus = where;
+	  return MATCH_ERROR;
+	}
+
       return MATCH_YES;
     }
   gfc_clear_error ();
@@ -1243,6 +1319,102 @@ gfc_match_omp_declare_reduction (void)
 
 
 match
+gfc_match_omp_declare_target (void)
+{
+  locus old_loc;
+  char n[GFC_MAX_SYMBOL_LEN+1];
+  gfc_symbol *sym;
+  match m;
+  gfc_symtree *st;
+
+  old_loc = gfc_current_locus;
+
+  m = gfc_match (" (");
+
+  if (gfc_current_ns->proc_name
+      && gfc_current_ns->proc_name->attr.if_source == IFSRC_IFBODY
+      && m == MATCH_YES)
+    {
+      gfc_error ("Only the !$OMP DECLARE TARGET form without "
+		 "list is allowed in interface block at %C");
+      goto cleanup;
+    }
+
+  if (m == MATCH_NO
+      && gfc_current_ns->proc_name
+      && gfc_match_omp_eos () == MATCH_YES)
+    {
+      if (!gfc_add_omp_declare_target (&gfc_current_ns->proc_name->attr,
+				       gfc_current_ns->proc_name->name,
+				       &old_loc))
+	goto cleanup;
+      return MATCH_YES;
+    }
+
+  if (m != MATCH_YES)
+    return m;
+
+  for (;;)
+    {
+      m = gfc_match_symbol (&sym, 0);
+      switch (m)
+	{
+	case MATCH_YES:
+	  if (sym->attr.in_common)
+	    gfc_error_now ("OMP DECLARE TARGET on a variable at %C is an "
+			   "element of a COMMON block");
+	  else if (!gfc_add_omp_declare_target (&sym->attr, sym->name,
+						&sym->declared_at))
+	    goto cleanup;
+	  goto next_item;
+	case MATCH_NO:
+	  break;
+	case MATCH_ERROR:
+	  goto cleanup;
+	}
+
+      m = gfc_match (" / %n /", n);
+      if (m == MATCH_ERROR)
+	goto cleanup;
+      if (m == MATCH_NO || n[0] == '\0')
+	goto syntax;
+
+      st = gfc_find_symtree (gfc_current_ns->common_root, n);
+      if (st == NULL)
+	{
+	  gfc_error ("COMMON block /%s/ not found at %C", n);
+	  goto cleanup;
+	}
+      st->n.common->omp_declare_target = 1;
+      for (sym = st->n.common->head; sym; sym = sym->common_next)
+	if (!gfc_add_omp_declare_target (&sym->attr, sym->name,
+					 &sym->declared_at))
+	  goto cleanup;
+
+    next_item:
+      if (gfc_match_char (')') == MATCH_YES)
+	break;
+      if (gfc_match_char (',') != MATCH_YES)
+	goto syntax;
+    }
+
+  if (gfc_match_omp_eos () != MATCH_YES)
+    {
+      gfc_error ("Unexpected junk after !$OMP DECLARE TARGET at %C");
+      goto cleanup;
+    }
+  return MATCH_YES;
+
+syntax:
+  gfc_error ("Syntax error in !$OMP DECLARE TARGET list at %C");
+
+cleanup:
+  gfc_current_locus = old_loc;
+  return MATCH_ERROR;
+}
+
+
+match
 gfc_match_omp_threadprivate (void)
 {
   locus old_loc;
@@ -1299,6 +1471,12 @@ gfc_match_omp_threadprivate (void)
 	goto syntax;
     }
 
+  if (gfc_match_omp_eos () != MATCH_YES)
+    {
+      gfc_error ("Unexpected junk after OMP THREADPRIVATE at %C");
+      goto cleanup;
+    }
+
   return MATCH_YES;
 
 syntax:
@@ -1311,83 +1489,213 @@ cleanup:
 
 
 match
+gfc_match_omp_parallel (void)
+{
+  return match_omp (EXEC_OMP_PARALLEL, OMP_PARALLEL_CLAUSES);
+}
+
+
+match
 gfc_match_omp_parallel_do (void)
 {
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES)
-      != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_PARALLEL_DO;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
+  return match_omp (EXEC_OMP_PARALLEL_DO,
+		    OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES);
 }
 
 
 match
 gfc_match_omp_parallel_do_simd (void)
 {
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, (OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES
-				  | OMP_SIMD_CLAUSES) & ~OMP_CLAUSE_ORDERED)
-      != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_PARALLEL_DO_SIMD;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
+  return match_omp (EXEC_OMP_PARALLEL_DO_SIMD,
+		    (OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
+		    & ~OMP_CLAUSE_ORDERED);
 }
 
 
 match
 gfc_match_omp_parallel_sections (void)
 {
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_PARALLEL_CLAUSES | OMP_SECTIONS_CLAUSES)
-      != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_PARALLEL_SECTIONS;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
+  return match_omp (EXEC_OMP_PARALLEL_SECTIONS,
+		    OMP_PARALLEL_CLAUSES | OMP_SECTIONS_CLAUSES);
 }
 
 
 match
 gfc_match_omp_parallel_workshare (void)
 {
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_PARALLEL_CLAUSES) != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_PARALLEL_WORKSHARE;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
+  return match_omp (EXEC_OMP_PARALLEL_WORKSHARE, OMP_PARALLEL_CLAUSES);
 }
 
 
 match
 gfc_match_omp_sections (void)
 {
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_SECTIONS_CLAUSES) != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_SECTIONS;
-  new_st.ext.omp_clauses = c;
-  return MATCH_YES;
+  return match_omp (EXEC_OMP_SECTIONS, OMP_SECTIONS_CLAUSES);
+}
+
+
+match
+gfc_match_omp_simd (void)
+{
+  return match_omp (EXEC_OMP_SIMD, OMP_SIMD_CLAUSES);
 }
 
 
 match
 gfc_match_omp_single (void)
 {
-  gfc_omp_clauses *c;
-  if (gfc_match_omp_clauses (&c, OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE)
-      != MATCH_YES)
-    return MATCH_ERROR;
-  new_st.op = EXEC_OMP_SINGLE;
-  new_st.ext.omp_clauses = c;
+  return match_omp (EXEC_OMP_SINGLE,
+		    OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE);
+}
+
+
+match
+gfc_match_omp_task (void)
+{
+  return match_omp (EXEC_OMP_TASK, OMP_TASK_CLAUSES);
+}
+
+
+match
+gfc_match_omp_taskwait (void)
+{
+  if (gfc_match_omp_eos () != MATCH_YES)
+    {
+      gfc_error ("Unexpected junk after TASKWAIT clause at %C");
+      return MATCH_ERROR;
+    }
+  new_st.op = EXEC_OMP_TASKWAIT;
+  new_st.ext.omp_clauses = NULL;
   return MATCH_YES;
 }
 
 
 match
+gfc_match_omp_taskyield (void)
+{
+  if (gfc_match_omp_eos () != MATCH_YES)
+    {
+      gfc_error ("Unexpected junk after TASKYIELD clause at %C");
+      return MATCH_ERROR;
+    }
+  new_st.op = EXEC_OMP_TASKYIELD;
+  new_st.ext.omp_clauses = NULL;
+  return MATCH_YES;
+}
+
+
+match
+gfc_match_omp_target (void)
+{
+  return match_omp (EXEC_OMP_TARGET, OMP_TARGET_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_data (void)
+{
+  return match_omp (EXEC_OMP_TARGET_DATA, OMP_TARGET_DATA_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_teams (void)
+{
+  return match_omp (EXEC_OMP_TARGET_TEAMS,
+		    OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_teams_distribute (void)
+{
+  return match_omp (EXEC_OMP_TARGET_TEAMS_DISTRIBUTE,
+		    OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES
+		    | OMP_DISTRIBUTE_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_teams_distribute_parallel_do (void)
+{
+  return match_omp (EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO,
+		    OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES
+		    | OMP_DISTRIBUTE_CLAUSES | OMP_PARALLEL_CLAUSES
+		    | OMP_DO_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_teams_distribute_parallel_do_simd (void)
+{
+  return match_omp (EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+		    (OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES
+		     | OMP_DISTRIBUTE_CLAUSES | OMP_PARALLEL_CLAUSES
+		     | OMP_DO_CLAUSES | OMP_SIMD_CLAUSES)
+		    & ~OMP_CLAUSE_ORDERED);
+}
+
+
+match
+gfc_match_omp_target_teams_distribute_simd (void)
+{
+  return match_omp (EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD,
+		    OMP_TARGET_CLAUSES | OMP_TEAMS_CLAUSES
+		    | OMP_DISTRIBUTE_CLAUSES | OMP_SIMD_CLAUSES);
+}
+
+
+match
+gfc_match_omp_target_update (void)
+{
+  return match_omp (EXEC_OMP_TARGET_UPDATE, OMP_TARGET_UPDATE_CLAUSES);
+}
+
+
+match
+gfc_match_omp_teams (void)
+{
+  return match_omp (EXEC_OMP_TEAMS, OMP_TEAMS_CLAUSES);
+}
+
+
+match
+gfc_match_omp_teams_distribute (void)
+{
+  return match_omp (EXEC_OMP_TEAMS_DISTRIBUTE,
+		    OMP_TEAMS_CLAUSES | OMP_DISTRIBUTE_CLAUSES);
+}
+
+
+match
+gfc_match_omp_teams_distribute_parallel_do (void)
+{
+  return match_omp (EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO,
+		    OMP_TEAMS_CLAUSES | OMP_DISTRIBUTE_CLAUSES
+		    | OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES);
+}
+
+
+match
+gfc_match_omp_teams_distribute_parallel_do_simd (void)
+{
+  return match_omp (EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD,
+		    (OMP_TEAMS_CLAUSES | OMP_DISTRIBUTE_CLAUSES
+		     | OMP_PARALLEL_CLAUSES | OMP_DO_CLAUSES
+		     | OMP_SIMD_CLAUSES) & ~OMP_CLAUSE_ORDERED);
+}
+
+
+match
+gfc_match_omp_teams_distribute_simd (void)
+{
+  return match_omp (EXEC_OMP_TEAMS_DISTRIBUTE_SIMD,
+		    OMP_TEAMS_CLAUSES | OMP_DISTRIBUTE_CLAUSES
+		    | OMP_SIMD_CLAUSES);
+}
+
+
+match
 gfc_match_omp_workshare (void)
 {
   if (gfc_match_omp_eos () != MATCH_YES)
@@ -1602,8 +1910,8 @@ resolve_omp_clauses (gfc_code *code, locus *where,
   int list;
   static const char *clause_names[]
     = { "PRIVATE", "FIRSTPRIVATE", "LASTPRIVATE", "COPYPRIVATE", "SHARED",
-	"COPYIN", "UNIFORM", "ALIGNED", "LINEAR", "DEPEND", "DEPEND",
-	"REDUCTION" };
+	"COPYIN", "UNIFORM", "ALIGNED", "LINEAR", "DEPEND", "MAP",
+	"TO", "FROM", "REDUCTION" };
 
   if (omp_clauses == NULL)
     return;
@@ -1692,8 +2000,10 @@ resolve_omp_clauses (gfc_code *code, locus *where,
     if (list != OMP_LIST_FIRSTPRIVATE
 	&& list != OMP_LIST_LASTPRIVATE
 	&& list != OMP_LIST_ALIGNED
-	&& list != OMP_LIST_DEPEND_IN
-	&& list != OMP_LIST_DEPEND_OUT)
+	&& list != OMP_LIST_DEPEND
+	&& list != OMP_LIST_MAP
+	&& list != OMP_LIST_FROM
+	&& list != OMP_LIST_TO)
       for (n = omp_clauses->lists[list]; n; n = n->next)
 	{
 	  if (n->sym->mark)
@@ -1745,6 +2055,20 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 	n->sym->mark = 1;
     }
 
+  for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
+    n->sym->mark = 0;
+  for (n = omp_clauses->lists[OMP_LIST_FROM]; n; n = n->next)
+    if (n->expr == NULL)
+      n->sym->mark = 1;
+  for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
+    {
+      if (n->expr == NULL && n->sym->mark)
+	gfc_error ("Symbol '%s' present on both FROM and TO clauses at %L",
+		   n->sym->name, where);
+      else
+	n->sym->mark = 1;
+    }
+
   for (list = 0; list < OMP_LIST_NUM; list++)
     if ((n = omp_clauses->lists[list]) != NULL)
       {
@@ -1819,8 +2143,10 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 		  }
 	      }
 	    break;
-	  case OMP_LIST_DEPEND_IN:
-	  case OMP_LIST_DEPEND_OUT:
+	  case OMP_LIST_DEPEND:
+	  case OMP_LIST_MAP:
+	  case OMP_LIST_TO:
+	  case OMP_LIST_FROM:
 	    for (; n != NULL; n = n->next)
 	      if (n->expr)
 		{
@@ -1829,11 +2155,11 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 		      || n->expr->ref == NULL
 		      || n->expr->ref->next
 		      || n->expr->ref->type != REF_ARRAY)
-		    gfc_error ("'%s' in DEPEND clause at %L is not a proper "
-			       "array section", n->sym->name, where);
+		    gfc_error ("'%s' in %s clause at %L is not a proper "
+			       "array section", n->sym->name, name, where);
 		  else if (n->expr->ref->u.ar.codimen)
-		    gfc_error ("Coarrays not supported in DEPEND clause at %L",
-			       where);
+		    gfc_error ("Coarrays not supported in %s clause at %L",
+			       name, where);
 		  else
 		    {
 		      int i;
@@ -1842,19 +2168,20 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 			if (ar->stride[i])
 			  {
 			    gfc_error ("Stride should not be specified for "
-				       "array section in DEPEND clause at %L",
-				       where);
+				       "array section in %s clause at %L",
+				       name, where);
 			    break;
 			  }
 			else if (ar->dimen_type[i] != DIMEN_ELEMENT
 				 && ar->dimen_type[i] != DIMEN_RANGE)
 			  {
-			    gfc_error ("'%s' in DEPEND clause at %L is not a "
+			    gfc_error ("'%s' in %s clause at %L is not a "
 				       "proper array section",
-				       n->sym->name, where);
+				       n->sym->name, name, where);
 			    break;
 			  }
-			else if (ar->start[i]
+			else if (list == OMP_LIST_DEPEND
+				 && ar->start[i]
 				 && ar->start[i]->expr_type == EXPR_CONSTANT
 				 && ar->end[i]
 				 && ar->end[i]->expr_type == EXPR_CONSTANT
@@ -1868,6 +2195,17 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 			  }
 		    }
 		}
+	    if (list != OMP_LIST_DEPEND)
+	      for (n = omp_clauses->lists[list]; n != NULL; n = n->next)
+		{
+		  n->sym->attr.referenced = 1;
+		  if (n->sym->attr.threadprivate)
+		    gfc_error ("THREADPRIVATE object '%s' in %s clause at %L",
+			       n->sym->name, name, where);
+		  if (n->sym->attr.cray_pointee)
+		    gfc_error ("Cray pointee '%s' in %s clause at %L",
+			       n->sym->name, name, where);
+		}
 	    break;
 	  default:
 	    for (; n != NULL; n = n->next)
@@ -1917,7 +2255,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 		switch (list)
 		  {
 		  case OMP_LIST_REDUCTION:
-		    switch (n->rop)
+		    switch (n->u.reduction_op)
 		      {
 		      case OMP_REDUCTION_PLUS:
 		      case OMP_REDUCTION_TIMES:
@@ -1964,7 +2302,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 			if (n->udr == NULL)
 			  {
 			    if (udr_name == NULL)
-			      switch (n->rop)
+			      switch (n->u.reduction_op)
 				{
 				case OMP_REDUCTION_PLUS:
 				case OMP_REDUCTION_TIMES:
@@ -1974,7 +2312,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 				case OMP_REDUCTION_EQV:
 				case OMP_REDUCTION_NEQV:
 				  udr_name = gfc_op2string ((gfc_intrinsic_op)
-							    n->rop);
+							    n->u.reduction_op);
 				  break;
 				case OMP_REDUCTION_MAX:
 				  udr_name = "max";
@@ -1999,7 +2337,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 				       gfc_typename (&n->sym->ts), where);
 			  }
 			else
-			  n->rop = OMP_REDUCTION_USER;
+			  n->u.reduction_op = OMP_REDUCTION_USER;
 		      }
 		    break;
 		  case OMP_LIST_LINEAR:
@@ -2051,6 +2389,38 @@ resolve_omp_clauses (gfc_code *code, locus *where,
 	gfc_error ("SIMDLEN clause at %L requires a scalar "
 		   "INTEGER expression", &expr->where);
     }
+  if (omp_clauses->num_teams)
+    {
+      gfc_expr *expr = omp_clauses->num_teams;
+      if (!gfc_resolve_expr (expr)
+	  || expr->ts.type != BT_INTEGER || expr->rank != 0)
+	gfc_error ("NUM_TEAMS clause at %L requires a scalar "
+		   "INTEGER expression", &expr->where);
+    }
+  if (omp_clauses->device)
+    {
+      gfc_expr *expr = omp_clauses->device;
+      if (!gfc_resolve_expr (expr)
+	  || expr->ts.type != BT_INTEGER || expr->rank != 0)
+	gfc_error ("DEVICE clause at %L requires a scalar "
+		   "INTEGER expression", &expr->where);
+    }
+  if (omp_clauses->dist_chunk_size)
+    {
+      gfc_expr *expr = omp_clauses->dist_chunk_size;
+      if (!gfc_resolve_expr (expr)
+	  || expr->ts.type != BT_INTEGER || expr->rank != 0)
+	gfc_error ("DIST_SCHEDULE clause's chunk_size at %L requires "
+		   "a scalar INTEGER expression", &expr->where);
+    }
+  if (omp_clauses->thread_limit)
+    {
+      gfc_expr *expr = omp_clauses->thread_limit;
+      if (!gfc_resolve_expr (expr)
+	  || expr->ts.type != BT_INTEGER || expr->rank != 0)
+	gfc_error ("THREAD_LIMIT clause at %L requires a scalar "
+		   "INTEGER expression", &expr->where);
+    }
 }
 
 
@@ -2565,14 +2935,38 @@ gfc_resolve_omp_parallel_blocks (gfc_code *code, gfc_namespace *ns)
   omp_current_ctx = &ctx;
 
   for (list = 0; list < OMP_LIST_NUM; list++)
-    for (n = omp_clauses->lists[list]; n; n = n->next)
-      pointer_set_insert (ctx.sharing_clauses, n->sym);
+    switch (list)
+      {
+      case OMP_LIST_SHARED:
+      case OMP_LIST_PRIVATE:
+      case OMP_LIST_FIRSTPRIVATE:
+      case OMP_LIST_LASTPRIVATE:
+      case OMP_LIST_REDUCTION:
+      case OMP_LIST_LINEAR:
+	for (n = omp_clauses->lists[list]; n; n = n->next)
+	  pointer_set_insert (ctx.sharing_clauses, n->sym);
+	break;
+      default:
+	break;
+      }
 
-  if (code->op == EXEC_OMP_PARALLEL_DO
-      || code->op == EXEC_OMP_PARALLEL_DO_SIMD)
-    gfc_resolve_omp_do_blocks (code, ns);
-  else
-    gfc_resolve_blocks (code->block, ns);
+  switch (code->op)
+    {
+    case EXEC_OMP_PARALLEL_DO:
+    case EXEC_OMP_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+    case EXEC_OMP_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+      gfc_resolve_omp_do_blocks (code, ns);
+      break;
+    default:
+      gfc_resolve_blocks (code->block, ns);
+    }
 
   omp_current_ctx = ctx.previous;
   pointer_set_destroy (ctx.sharing_clauses);
@@ -2660,13 +3054,52 @@ resolve_omp_do (gfc_code *code)
 
   switch (code->op)
     {
+    case EXEC_OMP_DISTRIBUTE: name = "!$OMP DISTRIBUTE"; break;
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+      name = "!$OMP DISTRIBUTE PARALLEL DO";
+      break;
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+      name = "!$OMP DISTRIBUTE PARALLEL DO SIMD";
+      is_simd = true;
+      break;
+    case EXEC_OMP_DISTRIBUTE_SIMD:
+      name = "!$OMP DISTRIBUTE SIMD";
+      is_simd = true;
+      break;
     case EXEC_OMP_DO: name = "!$OMP DO"; break;
     case EXEC_OMP_DO_SIMD: name = "!$OMP DO SIMD"; is_simd = true; break;
     case EXEC_OMP_PARALLEL_DO: name = "!$OMP PARALLEL DO"; break;
     case EXEC_OMP_PARALLEL_DO_SIMD:
       name = "!$OMP PARALLEL DO SIMD";
-      is_simd = true; break;
+      is_simd = true;
+      break;
     case EXEC_OMP_SIMD: name = "!$OMP SIMD"; is_simd = true; break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+      name = "!$OMP TARGET TEAMS_DISTRIBUTE";
+      break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      name = "!$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO";
+      break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      name = "!$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD";
+      is_simd = true;
+      break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+      name = "!$OMP TARGET TEAMS DISTRIBUTE SIMD";
+      is_simd = true;
+      break;
+    case EXEC_OMP_TEAMS_DISTRIBUTE: name = "!$OMP TEAMS_DISTRIBUTE"; break;
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      name = "!$OMP TEAMS DISTRIBUTE PARALLEL DO";
+      break;
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      name = "!$OMP TEAMS DISTRIBUTE PARALLEL DO SIMD";
+      is_simd = true;
+      break;
+    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+      name = "!$OMP TEAMS DISTRIBUTE SIMD";
+      is_simd = true;
+      break;
     default: gcc_unreachable ();
     }
 
@@ -2786,11 +3219,23 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED)
 
   switch (code->op)
     {
+    case EXEC_OMP_DISTRIBUTE:
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_DISTRIBUTE_SIMD:
     case EXEC_OMP_DO:
     case EXEC_OMP_DO_SIMD:
     case EXEC_OMP_PARALLEL_DO:
     case EXEC_OMP_PARALLEL_DO_SIMD:
     case EXEC_OMP_SIMD:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+    case EXEC_OMP_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
       resolve_omp_do (code);
       break;
     case EXEC_OMP_CANCEL:
@@ -2799,11 +3244,24 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED)
     case EXEC_OMP_PARALLEL_SECTIONS:
     case EXEC_OMP_SECTIONS:
     case EXEC_OMP_SINGLE:
+    case EXEC_OMP_TARGET:
+    case EXEC_OMP_TARGET_DATA:
+    case EXEC_OMP_TARGET_TEAMS:
     case EXEC_OMP_TASK:
+    case EXEC_OMP_TEAMS:
     case EXEC_OMP_WORKSHARE:
       if (code->ext.omp_clauses)
 	resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL);
       break;
+    case EXEC_OMP_TARGET_UPDATE:
+      if (code->ext.omp_clauses)
+	resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL);
+      if (code->ext.omp_clauses == NULL
+	  || (code->ext.omp_clauses->lists[OMP_LIST_TO] == NULL
+	      && code->ext.omp_clauses->lists[OMP_LIST_FROM] == NULL))
+	gfc_error ("OMP TARGET UPDATE at %L requires at least one TO or "
+		   "FROM clause", &code->loc);
+      break;
     case EXEC_OMP_ATOMIC:
       resolve_omp_atomic (code);
       break;
@@ -2822,7 +3280,7 @@ gfc_resolve_omp_declare_simd (gfc_namespace *ns)
   for (ods = ns->omp_declare_simd; ods; ods = ods->next)
     {
       if (ods->proc_name != ns->proc_name)
-	gfc_error ("!$OMP DECLARE SIMD should refer to containing procedure"
+	gfc_error ("!$OMP DECLARE SIMD should refer to containing procedure "
 		   "'%s' at %L", ns->proc_name->name, &ods->where);
       if (ods->clauses)
 	resolve_omp_clauses (NULL, &ods->where, ods->clauses, ns);
diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c
index bdee831ae4d..e8dcb70edb2 100644
--- a/gcc/fortran/parse.c
+++ b/gcc/fortran/parse.c
@@ -633,12 +633,29 @@ decode_omp_directive (void)
 	      ST_OMP_DECLARE_REDUCTION);
       matchs ("declare simd", gfc_match_omp_declare_simd,
 	      ST_OMP_DECLARE_SIMD);
+      matcho ("declare target", gfc_match_omp_declare_target,
+	      ST_OMP_DECLARE_TARGET);
+      matchs ("distribute parallel do simd",
+	      gfc_match_omp_distribute_parallel_do_simd,
+	      ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD);
+      matcho ("distribute parallel do", gfc_match_omp_distribute_parallel_do,
+	      ST_OMP_DISTRIBUTE_PARALLEL_DO);
+      matchs ("distribute simd", gfc_match_omp_distribute_simd,
+	      ST_OMP_DISTRIBUTE_SIMD);
+      matcho ("distribute", gfc_match_omp_distribute, ST_OMP_DISTRIBUTE);
       matchs ("do simd", gfc_match_omp_do_simd, ST_OMP_DO_SIMD);
       matcho ("do", gfc_match_omp_do, ST_OMP_DO);
       break;
     case 'e':
       matcho ("end atomic", gfc_match_omp_eos, ST_OMP_END_ATOMIC);
       matcho ("end critical", gfc_match_omp_critical, ST_OMP_END_CRITICAL);
+      matchs ("end distribute parallel do simd", gfc_match_omp_eos,
+	      ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD);
+      matcho ("end distribute parallel do", gfc_match_omp_eos,
+	      ST_OMP_END_DISTRIBUTE_PARALLEL_DO);
+      matchs ("end distribute simd", gfc_match_omp_eos,
+	      ST_OMP_END_DISTRIBUTE_SIMD);
+      matcho ("end distribute", gfc_match_omp_eos, ST_OMP_END_DISTRIBUTE);
       matchs ("end do simd", gfc_match_omp_end_nowait, ST_OMP_END_DO_SIMD);
       matcho ("end do", gfc_match_omp_end_nowait, ST_OMP_END_DO);
       matchs ("end simd", gfc_match_omp_eos, ST_OMP_END_SIMD);
@@ -654,8 +671,29 @@ decode_omp_directive (void)
       matcho ("end parallel", gfc_match_omp_eos, ST_OMP_END_PARALLEL);
       matcho ("end sections", gfc_match_omp_end_nowait, ST_OMP_END_SECTIONS);
       matcho ("end single", gfc_match_omp_end_single, ST_OMP_END_SINGLE);
+      matcho ("end target data", gfc_match_omp_eos, ST_OMP_END_TARGET_DATA);
+      matchs ("end target teams distribute parallel do simd",
+	      gfc_match_omp_eos,
+	      ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
+      matcho ("end target teams distribute parallel do", gfc_match_omp_eos,
+	      ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO);
+      matchs ("end target teams distribute simd", gfc_match_omp_eos,
+	      ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD);
+      matcho ("end target teams distribute", gfc_match_omp_eos,
+	      ST_OMP_END_TARGET_TEAMS_DISTRIBUTE);
+      matcho ("end target teams", gfc_match_omp_eos, ST_OMP_END_TARGET_TEAMS);
+      matcho ("end target", gfc_match_omp_eos, ST_OMP_END_TARGET);
       matcho ("end taskgroup", gfc_match_omp_eos, ST_OMP_END_TASKGROUP);
       matcho ("end task", gfc_match_omp_eos, ST_OMP_END_TASK);
+      matchs ("end teams distribute parallel do simd", gfc_match_omp_eos,
+	      ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
+      matcho ("end teams distribute parallel do", gfc_match_omp_eos,
+	      ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO);
+      matchs ("end teams distribute simd", gfc_match_omp_eos,
+	      ST_OMP_END_TEAMS_DISTRIBUTE_SIMD);
+      matcho ("end teams distribute", gfc_match_omp_eos,
+	      ST_OMP_END_TEAMS_DISTRIBUTE);
+      matcho ("end teams", gfc_match_omp_eos, ST_OMP_END_TEAMS);
       matcho ("end workshare", gfc_match_omp_end_nowait,
 	      ST_OMP_END_WORKSHARE);
       break;
@@ -685,10 +723,37 @@ decode_omp_directive (void)
       matcho ("single", gfc_match_omp_single, ST_OMP_SINGLE);
       break;
     case 't':
+      matcho ("target data", gfc_match_omp_target_data, ST_OMP_TARGET_DATA);
+      matchs ("target teams distribute parallel do simd",
+	      gfc_match_omp_target_teams_distribute_parallel_do_simd,
+	      ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
+      matcho ("target teams distribute parallel do",
+	      gfc_match_omp_target_teams_distribute_parallel_do,
+	      ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO);
+      matchs ("target teams distribute simd",
+	      gfc_match_omp_target_teams_distribute_simd,
+	      ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD);
+      matcho ("target teams distribute", gfc_match_omp_target_teams_distribute,
+	      ST_OMP_TARGET_TEAMS_DISTRIBUTE);
+      matcho ("target teams", gfc_match_omp_target_teams, ST_OMP_TARGET_TEAMS);
+      matcho ("target update", gfc_match_omp_target_update,
+	      ST_OMP_TARGET_UPDATE);
+      matcho ("target", gfc_match_omp_target, ST_OMP_TARGET);
       matcho ("taskgroup", gfc_match_omp_taskgroup, ST_OMP_TASKGROUP);
       matcho ("taskwait", gfc_match_omp_taskwait, ST_OMP_TASKWAIT);
       matcho ("taskyield", gfc_match_omp_taskyield, ST_OMP_TASKYIELD);
       matcho ("task", gfc_match_omp_task, ST_OMP_TASK);
+      matchs ("teams distribute parallel do simd",
+	      gfc_match_omp_teams_distribute_parallel_do_simd,
+	      ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD);
+      matcho ("teams distribute parallel do",
+	      gfc_match_omp_teams_distribute_parallel_do,
+	      ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO);
+      matchs ("teams distribute simd", gfc_match_omp_teams_distribute_simd,
+	      ST_OMP_TEAMS_DISTRIBUTE_SIMD);
+      matcho ("teams distribute", gfc_match_omp_teams_distribute,
+	      ST_OMP_TEAMS_DISTRIBUTE);
+      matcho ("teams", gfc_match_omp_teams, ST_OMP_TEAMS);
       matcho ("threadprivate", gfc_match_omp_threadprivate,
 	      ST_OMP_THREADPRIVATE);
       break;
@@ -1094,8 +1159,8 @@ next_statement (void)
   case ST_LABEL_ASSIGNMENT: case ST_FLUSH: case ST_OMP_FLUSH: \
   case ST_OMP_BARRIER: case ST_OMP_TASKWAIT: case ST_OMP_TASKYIELD: \
   case ST_OMP_CANCEL: case ST_OMP_CANCELLATION_POINT: \
-  case ST_ERROR_STOP: case ST_SYNC_ALL: case ST_SYNC_IMAGES: \
-  case ST_SYNC_MEMORY: case ST_LOCK: case ST_UNLOCK
+  case ST_OMP_TARGET_UPDATE: case ST_ERROR_STOP: case ST_SYNC_ALL: \
+  case ST_SYNC_IMAGES: case ST_SYNC_MEMORY: case ST_LOCK: case ST_UNLOCK
 
 /* Statements that mark other executable statements.  */
 
@@ -1108,14 +1173,27 @@ next_statement (void)
   case ST_OMP_DO: case ST_OMP_PARALLEL_DO: case ST_OMP_ATOMIC: \
   case ST_OMP_WORKSHARE: case ST_OMP_PARALLEL_WORKSHARE: \
   case ST_OMP_TASK: case ST_OMP_TASKGROUP: case ST_OMP_SIMD: \
-  case ST_OMP_DO_SIMD: case ST_OMP_PARALLEL_DO_SIMD: case ST_CRITICAL
+  case ST_OMP_DO_SIMD: case ST_OMP_PARALLEL_DO_SIMD: case ST_OMP_TARGET: \
+  case ST_OMP_TARGET_DATA: case ST_OMP_TARGET_TEAMS: \
+  case ST_OMP_TARGET_TEAMS_DISTRIBUTE: \
+  case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: \
+  case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: \
+  case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: \
+  case ST_OMP_TEAMS: case ST_OMP_TEAMS_DISTRIBUTE: \
+  case ST_OMP_TEAMS_DISTRIBUTE_SIMD: \
+  case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: \
+  case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: case ST_OMP_DISTRIBUTE: \
+  case ST_OMP_DISTRIBUTE_SIMD: case ST_OMP_DISTRIBUTE_PARALLEL_DO: \
+  case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD: \
+  case ST_CRITICAL
 
 /* Declaration statements */
 
 #define case_decl case ST_ATTR_DECL: case ST_COMMON: case ST_DATA_DECL: \
   case ST_EQUIVALENCE: case ST_NAMELIST: case ST_STATEMENT_FUNCTION: \
   case ST_TYPE: case ST_INTERFACE: case ST_OMP_THREADPRIVATE: \
-  case ST_PROCEDURE: case ST_OMP_DECLARE_SIMD: case ST_OMP_DECLARE_REDUCTION
+  case ST_PROCEDURE: case ST_OMP_DECLARE_SIMD: case ST_OMP_DECLARE_REDUCTION: \
+  case ST_OMP_DECLARE_TARGET
 
 /* Block end statements.  Errors associated with interchanging these
    are detected in gfc_match_end().  */
@@ -1621,6 +1699,21 @@ gfc_ascii_statement (gfc_statement st)
     case ST_OMP_DECLARE_SIMD:
       p = "!$OMP DECLARE SIMD";
       break;
+    case ST_OMP_DECLARE_TARGET:
+      p = "!$OMP DECLARE TARGET";
+      break;
+    case ST_OMP_DISTRIBUTE:
+      p = "!$OMP DISTRIBUTE";
+      break;
+    case ST_OMP_DISTRIBUTE_PARALLEL_DO:
+      p = "!$OMP DISTRIBUTE PARALLEL DO";
+      break;
+    case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+      p = "!$OMP DISTRIBUTE PARALLEL DO SIMD";
+      break;
+    case ST_OMP_DISTRIBUTE_SIMD:
+      p = "!$OMP DISTRIBUTE SIMD";
+      break;
     case ST_OMP_DO:
       p = "!$OMP DO";
       break;
@@ -1633,6 +1726,18 @@ gfc_ascii_statement (gfc_statement st)
     case ST_OMP_END_CRITICAL:
       p = "!$OMP END CRITICAL";
       break;
+    case ST_OMP_END_DISTRIBUTE:
+      p = "!$OMP END DISTRIBUTE";
+      break;
+    case ST_OMP_END_DISTRIBUTE_PARALLEL_DO:
+      p = "!$OMP END DISTRIBUTE PARALLEL DO";
+      break;
+    case ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD:
+      p = "!$OMP END DISTRIBUTE PARALLEL DO SIMD";
+      break;
+    case ST_OMP_END_DISTRIBUTE_SIMD:
+      p = "!$OMP END DISTRIBUTE SIMD";
+      break;
     case ST_OMP_END_DO:
       p = "!$OMP END DO";
       break;
@@ -1672,9 +1777,45 @@ gfc_ascii_statement (gfc_statement st)
     case ST_OMP_END_TASK:
       p = "!$OMP END TASK";
       break;
+    case ST_OMP_END_TARGET:
+      p = "!$OMP END TARGET";
+      break;
+    case ST_OMP_END_TARGET_DATA:
+      p = "!$OMP END TARGET DATA";
+      break;
+    case ST_OMP_END_TARGET_TEAMS:
+      p = "!$OMP END TARGET TEAMS";
+      break;
+    case ST_OMP_END_TARGET_TEAMS_DISTRIBUTE:
+      p = "!$OMP END TARGET TEAMS DISTRIBUTE";
+      break;
+    case ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      p = "!$OMP END TARGET TEAMS DISTRIBUTE PARALLEL DO";
+      break;
+    case ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      p = "!$OMP END TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD";
+      break;
+    case ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD:
+      p = "!$OMP END TARGET TEAMS DISTRIBUTE SIMD";
+      break;
     case ST_OMP_END_TASKGROUP:
       p = "!$OMP END TASKGROUP";
       break;
+    case ST_OMP_END_TEAMS:
+      p = "!$OMP END TEAMS";
+      break;
+    case ST_OMP_END_TEAMS_DISTRIBUTE:
+      p = "!$OMP END TEAMS DISTRIBUTE";
+      break;
+    case ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      p = "!$OMP END TEAMS DISTRIBUTE PARALLEL DO";
+      break;
+    case ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      p = "!$OMP END TEAMS DISTRIBUTE PARALLEL DO SIMD";
+      break;
+    case ST_OMP_END_TEAMS_DISTRIBUTE_SIMD:
+      p = "!$OMP END TEAMS DISTRIBUTE SIMD";
+      break;
     case ST_OMP_END_WORKSHARE:
       p = "!$OMP END WORKSHARE";
       break;
@@ -1714,6 +1855,30 @@ gfc_ascii_statement (gfc_statement st)
     case ST_OMP_SINGLE:
       p = "!$OMP SINGLE";
       break;
+    case ST_OMP_TARGET:
+      p = "!$OMP TARGET";
+      break;
+    case ST_OMP_TARGET_DATA:
+      p = "!$OMP TARGET DATA";
+      break;
+    case ST_OMP_TARGET_TEAMS:
+      p = "!$OMP TARGET TEAMS";
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
+      p = "!$OMP TARGET TEAMS DISTRIBUTE";
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      p = "!$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO";
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      p = "!$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO SIMD";
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+      p = "!$OMP TARGET TEAMS DISTRIBUTE SIMD";
+      break;
+    case ST_OMP_TARGET_UPDATE:
+      p = "!$OMP TARGET UPDATE";
+      break;
     case ST_OMP_TASK:
       p = "!$OMP TASK";
       break;
@@ -1726,6 +1891,21 @@ gfc_ascii_statement (gfc_statement st)
     case ST_OMP_TASKYIELD:
       p = "!$OMP TASKYIELD";
       break;
+    case ST_OMP_TEAMS:
+      p = "!$OMP TEAMS";
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE:
+      p = "!$OMP TEAMS DISTRIBUTE";
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      p = "!$OMP TEAMS DISTRIBUTE PARALLEL DO";
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      p = "!$OMP TEAMS DISTRIBUTE PARALLEL DO SIMD";
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
+      p = "!$OMP TEAMS DISTRIBUTE SIMD";
+      break;
     case ST_OMP_THREADPRIVATE:
       p = "!$OMP THREADPRIVATE";
       break;
@@ -3699,13 +3879,47 @@ parse_omp_do (gfc_statement omp_st)
   gfc_statement omp_end_st = ST_OMP_END_DO;
   switch (omp_st)
     {
-    case ST_OMP_SIMD: omp_end_st = ST_OMP_END_SIMD; break;
+    case ST_OMP_DISTRIBUTE: omp_end_st = ST_OMP_END_DISTRIBUTE; break;
+    case ST_OMP_DISTRIBUTE_PARALLEL_DO:
+      omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO;
+      break;
+    case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+      omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD;
+      break;
+    case ST_OMP_DISTRIBUTE_SIMD:
+      omp_end_st = ST_OMP_END_DISTRIBUTE_SIMD;
+      break;
     case ST_OMP_DO: omp_end_st = ST_OMP_END_DO; break;
     case ST_OMP_DO_SIMD: omp_end_st = ST_OMP_END_DO_SIMD; break;
     case ST_OMP_PARALLEL_DO: omp_end_st = ST_OMP_END_PARALLEL_DO; break;
     case ST_OMP_PARALLEL_DO_SIMD:
       omp_end_st = ST_OMP_END_PARALLEL_DO_SIMD;
       break;
+    case ST_OMP_SIMD: omp_end_st = ST_OMP_END_SIMD; break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE;
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO;
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD;
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE:
+      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE;
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO;
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
+      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_SIMD;
+      break;
     default: gcc_unreachable ();
     }
   if (st == omp_end_st)
@@ -3814,12 +4028,60 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only)
     case ST_OMP_SINGLE:
       omp_end_st = ST_OMP_END_SINGLE;
       break;
+    case ST_OMP_TARGET:
+      omp_end_st = ST_OMP_END_TARGET;
+      break;
+    case ST_OMP_TARGET_DATA:
+      omp_end_st = ST_OMP_END_TARGET_DATA;
+      break;
+    case ST_OMP_TARGET_TEAMS:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS;
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE;
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO;
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
+      break;
+    case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+      omp_end_st = ST_OMP_END_TARGET_TEAMS_DISTRIBUTE_SIMD;
+      break;
     case ST_OMP_TASK:
       omp_end_st = ST_OMP_END_TASK;
       break;
     case ST_OMP_TASKGROUP:
       omp_end_st = ST_OMP_END_TASKGROUP;
       break;
+    case ST_OMP_TEAMS:
+      omp_end_st = ST_OMP_END_TEAMS;
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE:
+      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE;
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO;
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD;
+      break;
+    case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
+      omp_end_st = ST_OMP_END_TEAMS_DISTRIBUTE_SIMD;
+      break;
+    case ST_OMP_DISTRIBUTE:
+      omp_end_st = ST_OMP_END_DISTRIBUTE;
+      break;
+    case ST_OMP_DISTRIBUTE_PARALLEL_DO:
+      omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO;
+      break;
+    case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+      omp_end_st = ST_OMP_END_DISTRIBUTE_PARALLEL_DO_SIMD;
+      break;
+    case ST_OMP_DISTRIBUTE_SIMD:
+      omp_end_st = ST_OMP_END_DISTRIBUTE_SIMD;
+      break;
     case ST_OMP_WORKSHARE:
       omp_end_st = ST_OMP_END_WORKSHARE;
       break;
@@ -4052,6 +4314,10 @@ parse_executable (gfc_statement st)
 	case ST_OMP_CRITICAL:
 	case ST_OMP_MASTER:
 	case ST_OMP_SINGLE:
+	case ST_OMP_TARGET:
+	case ST_OMP_TARGET_DATA:
+	case ST_OMP_TARGET_TEAMS:
+	case ST_OMP_TEAMS:
 	case ST_OMP_TASK:
 	case ST_OMP_TASKGROUP:
 	  parse_omp_structured_block (st, false);
@@ -4062,11 +4328,23 @@ parse_executable (gfc_statement st)
 	  parse_omp_structured_block (st, true);
 	  break;
 
+	case ST_OMP_DISTRIBUTE:
+	case ST_OMP_DISTRIBUTE_PARALLEL_DO:
+	case ST_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case ST_OMP_DISTRIBUTE_SIMD:
 	case ST_OMP_DO:
 	case ST_OMP_DO_SIMD:
 	case ST_OMP_PARALLEL_DO:
 	case ST_OMP_PARALLEL_DO_SIMD:
 	case ST_OMP_SIMD:
+	case ST_OMP_TARGET_TEAMS_DISTRIBUTE:
+	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case ST_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+	case ST_OMP_TEAMS_DISTRIBUTE:
+	case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case ST_OMP_TEAMS_DISTRIBUTE_SIMD:
 	  st = parse_omp_do (st);
 	  if (st == ST_IMPLIED_ENDDO)
 	    return st;
diff --git a/gcc/fortran/resolve.c b/gcc/fortran/resolve.c
index 7ea7c36e8f9..64f34898770 100644
--- a/gcc/fortran/resolve.c
+++ b/gcc/fortran/resolve.c
@@ -9032,6 +9032,10 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns)
 
 	case EXEC_OMP_ATOMIC:
 	case EXEC_OMP_CRITICAL:
+	case EXEC_OMP_DISTRIBUTE:
+	case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_DISTRIBUTE_SIMD:
 	case EXEC_OMP_DO:
 	case EXEC_OMP_DO_SIMD:
 	case EXEC_OMP_MASTER:
@@ -9044,10 +9048,23 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns)
 	case EXEC_OMP_SECTIONS:
 	case EXEC_OMP_SIMD:
 	case EXEC_OMP_SINGLE:
+	case EXEC_OMP_TARGET:
+	case EXEC_OMP_TARGET_DATA:
+	case EXEC_OMP_TARGET_TEAMS:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+	case EXEC_OMP_TARGET_UPDATE:
 	case EXEC_OMP_TASK:
 	case EXEC_OMP_TASKGROUP:
 	case EXEC_OMP_TASKWAIT:
 	case EXEC_OMP_TASKYIELD:
+	case EXEC_OMP_TEAMS:
+	case EXEC_OMP_TEAMS_DISTRIBUTE:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
 	case EXEC_OMP_WORKSHARE:
 	  break;
 
@@ -9827,11 +9844,23 @@ resolve_code (gfc_code *code, gfc_namespace *ns)
 	    case EXEC_OMP_PARALLEL_DO:
 	    case EXEC_OMP_PARALLEL_DO_SIMD:
 	    case EXEC_OMP_PARALLEL_SECTIONS:
+	    case EXEC_OMP_TARGET_TEAMS:
+	    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+	    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
 	    case EXEC_OMP_TASK:
+	    case EXEC_OMP_TEAMS:
+	    case EXEC_OMP_TEAMS_DISTRIBUTE:
+	    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
 	      omp_workshare_save = omp_workshare_flag;
 	      omp_workshare_flag = 0;
 	      gfc_resolve_omp_parallel_blocks (code, ns);
 	      break;
+	    case EXEC_OMP_DISTRIBUTE:
+	    case EXEC_OMP_DISTRIBUTE_SIMD:
 	    case EXEC_OMP_DO:
 	    case EXEC_OMP_DO_SIMD:
 	    case EXEC_OMP_SIMD:
@@ -10160,6 +10189,10 @@ resolve_code (gfc_code *code, gfc_namespace *ns)
 	case EXEC_OMP_CANCELLATION_POINT:
 	case EXEC_OMP_CRITICAL:
 	case EXEC_OMP_FLUSH:
+	case EXEC_OMP_DISTRIBUTE:
+	case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_DISTRIBUTE_SIMD:
 	case EXEC_OMP_DO:
 	case EXEC_OMP_DO_SIMD:
 	case EXEC_OMP_MASTER:
@@ -10167,9 +10200,23 @@ resolve_code (gfc_code *code, gfc_namespace *ns)
 	case EXEC_OMP_SECTIONS:
 	case EXEC_OMP_SIMD:
 	case EXEC_OMP_SINGLE:
+	case EXEC_OMP_TARGET:
+	case EXEC_OMP_TARGET_DATA:
+	case EXEC_OMP_TARGET_TEAMS:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+	case EXEC_OMP_TARGET_UPDATE:
+	case EXEC_OMP_TASK:
 	case EXEC_OMP_TASKGROUP:
 	case EXEC_OMP_TASKWAIT:
 	case EXEC_OMP_TASKYIELD:
+	case EXEC_OMP_TEAMS:
+	case EXEC_OMP_TEAMS_DISTRIBUTE:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
 	case EXEC_OMP_WORKSHARE:
 	  gfc_resolve_omp_directive (code, ns);
 	  break;
@@ -10179,7 +10226,6 @@ resolve_code (gfc_code *code, gfc_namespace *ns)
 	case EXEC_OMP_PARALLEL_DO_SIMD:
 	case EXEC_OMP_PARALLEL_SECTIONS:
 	case EXEC_OMP_PARALLEL_WORKSHARE:
-	case EXEC_OMP_TASK:
 	  omp_workshare_save = omp_workshare_flag;
 	  omp_workshare_flag = 0;
 	  gfc_resolve_omp_directive (code, ns);
@@ -13541,6 +13587,18 @@ resolve_symbol (gfc_symbol *sym)
 	      || sym->ns->proc_name->attr.flavor != FL_MODULE)))
     gfc_error ("Threadprivate at %L isn't SAVEd", &sym->declared_at);
 
+  /* Check omp declare target restrictions.  */
+  if (sym->attr.omp_declare_target
+      && sym->attr.flavor == FL_VARIABLE
+      && !sym->attr.save
+      && !sym->ns->save_all
+      && (!sym->attr.in_common
+	  && sym->module == NULL
+	  && (sym->ns->proc_name == NULL
+	      || sym->ns->proc_name->attr.flavor != FL_MODULE)))
+    gfc_error ("!$OMP DECLARE TARGET variable '%s' at %L isn't SAVEd",
+	       sym->name, &sym->declared_at);
+
   /* If we have come this far we can apply default-initializers, as
      described in 14.7.5, to those variables that have not already
      been assigned one.  */
diff --git a/gcc/fortran/st.c b/gcc/fortran/st.c
index a3df43ed386..0f18f787231 100644
--- a/gcc/fortran/st.c
+++ b/gcc/fortran/st.c
@@ -187,6 +187,10 @@ gfc_free_statement (gfc_code *p)
 
     case EXEC_OMP_CANCEL:
     case EXEC_OMP_CANCELLATION_POINT:
+    case EXEC_OMP_DISTRIBUTE:
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_DISTRIBUTE_SIMD:
     case EXEC_OMP_DO:
     case EXEC_OMP_DO_SIMD:
     case EXEC_OMP_END_SINGLE:
@@ -197,7 +201,20 @@ gfc_free_statement (gfc_code *p)
     case EXEC_OMP_SECTIONS:
     case EXEC_OMP_SIMD:
     case EXEC_OMP_SINGLE:
+    case EXEC_OMP_TARGET:
+    case EXEC_OMP_TARGET_DATA:
+    case EXEC_OMP_TARGET_TEAMS:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+    case EXEC_OMP_TARGET_UPDATE:
     case EXEC_OMP_TASK:
+    case EXEC_OMP_TEAMS:
+    case EXEC_OMP_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
     case EXEC_OMP_WORKSHARE:
     case EXEC_OMP_PARALLEL_WORKSHARE:
       gfc_free_omp_clauses (p->ext.omp_clauses);
diff --git a/gcc/fortran/symbol.c b/gcc/fortran/symbol.c
index 922b421b5e1..aee7510a463 100644
--- a/gcc/fortran/symbol.c
+++ b/gcc/fortran/symbol.c
@@ -367,6 +367,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
     *asynchronous = "ASYNCHRONOUS", *codimension = "CODIMENSION",
     *contiguous = "CONTIGUOUS", *generic = "GENERIC";
   static const char *threadprivate = "THREADPRIVATE";
+  static const char *omp_declare_target = "OMP DECLARE TARGET";
 
   const char *a1, *a2;
   int standard;
@@ -453,6 +454,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
   conf (dummy, entry);
   conf (dummy, intrinsic);
   conf (dummy, threadprivate);
+  conf (dummy, omp_declare_target);
   conf (pointer, target);
   conf (pointer, intrinsic);
   conf (pointer, elemental);
@@ -495,6 +497,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
   conf (in_equivalence, entry);
   conf (in_equivalence, allocatable);
   conf (in_equivalence, threadprivate);
+  conf (in_equivalence, omp_declare_target);
 
   conf (dummy, result);
   conf (entry, result);
@@ -543,6 +546,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
   conf (cray_pointee, in_common);
   conf (cray_pointee, in_equivalence);
   conf (cray_pointee, threadprivate);
+  conf (cray_pointee, omp_declare_target);
 
   conf (data, dummy);
   conf (data, function);
@@ -596,6 +600,8 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
 
   conf (proc_pointer, abstract)
 
+  conf (entry, omp_declare_target)
+
   a1 = gfc_code2string (flavors, attr->flavor);
 
   if (attr->in_namelist
@@ -631,6 +637,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
       conf2 (function);
       conf2 (subroutine);
       conf2 (threadprivate);
+      conf2 (omp_declare_target);
 
       if (attr->access == ACCESS_PUBLIC || attr->access == ACCESS_PRIVATE)
 	{
@@ -712,6 +719,7 @@ check_conflict (symbol_attribute *attr, const char *name, locus *where)
       conf2 (subroutine);
       conf2 (threadprivate);
       conf2 (result);
+      conf2 (omp_declare_target);
 
       if (attr->intent != INTENT_UNKNOWN)
 	{
@@ -1207,6 +1215,22 @@ gfc_add_threadprivate (symbol_attribute *attr, const char *name, locus *where)
 
 
 bool
+gfc_add_omp_declare_target (symbol_attribute *attr, const char *name,
+			    locus *where)
+{
+
+  if (check_used (attr, name, where))
+    return false;
+
+  if (attr->omp_declare_target)
+    return true;
+
+  attr->omp_declare_target = 1;
+  return check_conflict (attr, name, where);
+}
+
+
+bool
 gfc_add_target (symbol_attribute *attr, locus *where)
 {
 
@@ -1761,6 +1785,9 @@ gfc_copy_attr (symbol_attribute *dest, symbol_attribute *src, locus *where)
   if (src->threadprivate
       && !gfc_add_threadprivate (dest, NULL, where))
     goto fail;
+  if (src->omp_declare_target
+      && !gfc_add_omp_declare_target (dest, NULL, where))
+    goto fail;
   if (src->target && !gfc_add_target (dest, where))
     goto fail;
   if (src->dummy && !gfc_add_dummy (dest, NULL, where))
diff --git a/gcc/fortran/trans-common.c b/gcc/fortran/trans-common.c
index 5a52984602f..f28eda6b634 100644
--- a/gcc/fortran/trans-common.c
+++ b/gcc/fortran/trans-common.c
@@ -456,6 +456,11 @@ build_common_decl (gfc_common_head *com, tree union_type, bool is_init)
       if (com->threadprivate)
 	set_decl_tls_model (decl, decl_default_tls_model (decl));
 
+      if (com->omp_declare_target)
+	DECL_ATTRIBUTES (decl)
+	  = tree_cons (get_identifier ("omp declare target"),
+		       NULL_TREE, DECL_ATTRIBUTES (decl));
+
       /* Place the back end declaration for this common block in
          GLOBAL_BINDING_LEVEL.  */
       gfc_map_of_all_commons[identifier] = pushdecl_top_level (decl);
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index 2e129c96118..f1a18c3cfd2 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -1222,6 +1222,10 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list)
 	list = chainon (list, attr);
       }
 
+  if (sym_attr.omp_declare_target)
+    list = tree_cons (get_identifier ("omp declare target"),
+		      NULL_TREE, list);
+
   return list;
 }
 
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 998d687761b..7667f2534f7 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -873,6 +873,110 @@ gfc_omp_clause_dtor (tree clause, tree decl)
 }
 
 
+void
+gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
+{
+  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+    return;
+
+  tree decl = OMP_CLAUSE_DECL (c);
+  tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
+  if (POINTER_TYPE_P (TREE_TYPE (decl)))
+    {
+      if (!gfc_omp_privatize_by_reference (decl)
+	  && !GFC_DECL_GET_SCALAR_POINTER (decl)
+	  && !GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
+	  && !GFC_DECL_CRAY_POINTEE (decl)
+	  && !GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
+	return;
+      c4 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+      OMP_CLAUSE_MAP_KIND (c4) = OMP_CLAUSE_MAP_POINTER;
+      OMP_CLAUSE_DECL (c4) = decl;
+      OMP_CLAUSE_SIZE (c4) = size_int (0);
+      decl = build_fold_indirect_ref (decl);
+      OMP_CLAUSE_DECL (c) = decl;
+    }
+  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+    {
+      stmtblock_t block;
+      gfc_start_block (&block);
+      tree type = TREE_TYPE (decl);
+      tree ptr = gfc_conv_descriptor_data_get (decl);
+      ptr = fold_convert (build_pointer_type (char_type_node), ptr);
+      ptr = build_fold_indirect_ref (ptr);
+      OMP_CLAUSE_DECL (c) = ptr;
+      c2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+      OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_TO_PSET;
+      OMP_CLAUSE_DECL (c2) = decl;
+      OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (type);
+      c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+      OMP_CLAUSE_MAP_KIND (c3) = OMP_CLAUSE_MAP_POINTER;
+      OMP_CLAUSE_DECL (c3) = gfc_conv_descriptor_data_get (decl);
+      OMP_CLAUSE_SIZE (c3) = size_int (0);
+      tree size = create_tmp_var (gfc_array_index_type, NULL);
+      tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+      elemsz = fold_convert (gfc_array_index_type, elemsz);
+      if (GFC_TYPE_ARRAY_AKIND (type) == GFC_ARRAY_POINTER
+	  || GFC_TYPE_ARRAY_AKIND (type) == GFC_ARRAY_POINTER_CONT)
+	{
+	  stmtblock_t cond_block;
+	  tree tem, then_b, else_b, zero, cond;
+
+	  gfc_init_block (&cond_block);
+	  tem = gfc_full_array_size (&cond_block, decl,
+				     GFC_TYPE_ARRAY_RANK (type));
+	  gfc_add_modify (&cond_block, size, tem);
+	  gfc_add_modify (&cond_block, size,
+			  fold_build2 (MULT_EXPR, gfc_array_index_type,
+				       size, elemsz));
+	  then_b = gfc_finish_block (&cond_block);
+	  gfc_init_block (&cond_block);
+	  zero = build_int_cst (gfc_array_index_type, 0);
+	  gfc_add_modify (&cond_block, size, zero);
+	  else_b = gfc_finish_block (&cond_block);
+	  tem = gfc_conv_descriptor_data_get (decl);
+	  tem = fold_convert (pvoid_type_node, tem);
+	  cond = fold_build2_loc (input_location, NE_EXPR,
+				  boolean_type_node, tem, null_pointer_node);
+	  gfc_add_expr_to_block (&block, build3_loc (input_location, COND_EXPR,
+						     void_type_node, cond,
+						     then_b, else_b));
+	}
+      else
+	{
+	  gfc_add_modify (&block, size,
+			  gfc_full_array_size (&block, decl,
+					       GFC_TYPE_ARRAY_RANK (type)));
+	  gfc_add_modify (&block, size,
+			  fold_build2 (MULT_EXPR, gfc_array_index_type,
+				       size, elemsz));
+	}
+      OMP_CLAUSE_SIZE (c) = size;
+      tree stmt = gfc_finish_block (&block);
+      gimplify_and_add (stmt, pre_p);
+    }
+  tree last = c;
+  if (c2)
+    {
+      OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (last);
+      OMP_CLAUSE_CHAIN (last) = c2;
+      last = c2;
+    }
+  if (c3)
+    {
+      OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (last);
+      OMP_CLAUSE_CHAIN (last) = c3;
+      last = c3;
+    }
+  if (c4)
+    {
+      OMP_CLAUSE_CHAIN (c4) = OMP_CLAUSE_CHAIN (last);
+      OMP_CLAUSE_CHAIN (last) = c4;
+      last = c4;
+    }
+}
+
+
 /* Return true if DECL's DECL_VALUE_EXPR (if any) should be
    disregarded in OpenMP construct, because it is going to be
    remapped during OpenMP lowering.  SHARED is true if DECL
@@ -1487,7 +1591,7 @@ gfc_trans_omp_reduction_list (gfc_omp_namelist *namelist, tree list,
 	    tree node = build_omp_clause (where.lb->location,
 					  OMP_CLAUSE_REDUCTION);
 	    OMP_CLAUSE_DECL (node) = t;
-	    switch (namelist->rop)
+	    switch (namelist->u.reduction_op)
 	      {
 	      case OMP_REDUCTION_PLUS:
 		OMP_CLAUSE_REDUCTION_CODE (node) = PLUS_EXPR;
@@ -1532,7 +1636,7 @@ gfc_trans_omp_reduction_list (gfc_omp_namelist *namelist, tree list,
 		gcc_unreachable ();
 	      }
 	    if (namelist->sym->attr.dimension
-		|| namelist->rop == OMP_REDUCTION_USER
+		|| namelist->u.reduction_op == OMP_REDUCTION_USER
 		|| namelist->sym->attr.allocatable)
 	      gfc_trans_omp_array_reduction_or_udr (node, namelist, where);
 	    list = gfc_trans_add_clause (node, list);
@@ -1661,8 +1765,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	      }
 	  }
 	  break;
-	case OMP_LIST_DEPEND_IN:
-	case OMP_LIST_DEPEND_OUT:
+	case OMP_LIST_DEPEND:
 	  for (; n != NULL; n = n->next)
 	    {
 	      if (!n->sym->attr.referenced)
@@ -1671,9 +1774,19 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	      tree node = build_omp_clause (input_location, OMP_CLAUSE_DEPEND);
 	      if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
 		{
-		  OMP_CLAUSE_DECL (node) = gfc_get_symbol_decl (n->sym);
-		  if (DECL_P (OMP_CLAUSE_DECL (node)))
-		    TREE_ADDRESSABLE (OMP_CLAUSE_DECL (node)) = 1;
+		  tree decl = gfc_get_symbol_decl (n->sym);
+		  if (gfc_omp_privatize_by_reference (decl))
+		    decl = build_fold_indirect_ref (decl);
+		  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+		    {
+		      decl = gfc_conv_descriptor_data_get (decl);
+		      decl = fold_convert (build_pointer_type (char_type_node),
+					   decl);
+		      decl = build_fold_indirect_ref (decl);
+		    }
+		  else if (DECL_P (decl))
+		    TREE_ADDRESSABLE (decl) = 1;
+		  OMP_CLAUSE_DECL (node) = decl;
 		}
 	      else
 		{
@@ -1691,13 +1804,286 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		    }
 		  gfc_add_block_to_block (block, &se.pre);
 		  gfc_add_block_to_block (block, &se.post);
-		  OMP_CLAUSE_DECL (node)
-		    = fold_build1_loc (input_location, INDIRECT_REF,
-				       TREE_TYPE (TREE_TYPE (ptr)), ptr);
+		  ptr = fold_convert (build_pointer_type (char_type_node),
+				      ptr);
+		  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+		}
+	      switch (n->u.depend_op)
+		{
+		case OMP_DEPEND_IN:
+		  OMP_CLAUSE_DEPEND_KIND (node) = OMP_CLAUSE_DEPEND_IN;
+		  break;
+		case OMP_DEPEND_OUT:
+		  OMP_CLAUSE_DEPEND_KIND (node) = OMP_CLAUSE_DEPEND_OUT;
+		  break;
+		case OMP_DEPEND_INOUT:
+		  OMP_CLAUSE_DEPEND_KIND (node) = OMP_CLAUSE_DEPEND_INOUT;
+		  break;
+		default:
+		  gcc_unreachable ();
+		}
+	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+	    }
+	  break;
+	case OMP_LIST_MAP:
+	  for (; n != NULL; n = n->next)
+	    {
+	      if (!n->sym->attr.referenced)
+		continue;
+
+	      tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+	      tree node2 = NULL_TREE;
+	      tree node3 = NULL_TREE;
+	      tree node4 = NULL_TREE;
+	      tree decl = gfc_get_symbol_decl (n->sym);
+	      if (DECL_P (decl))
+		TREE_ADDRESSABLE (decl) = 1;
+	      if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
+		{
+		  if (POINTER_TYPE_P (TREE_TYPE (decl)))
+		    {
+		      node4 = build_omp_clause (input_location,
+						OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_MAP_KIND (node4) = OMP_CLAUSE_MAP_POINTER;
+		      OMP_CLAUSE_DECL (node4) = decl;
+		      OMP_CLAUSE_SIZE (node4) = size_int (0);
+		      decl = build_fold_indirect_ref (decl);
+		    }
+		  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+		    {
+		      tree type = TREE_TYPE (decl);
+		      tree ptr = gfc_conv_descriptor_data_get (decl);
+		      ptr = fold_convert (build_pointer_type (char_type_node),
+					  ptr);
+		      ptr = build_fold_indirect_ref (ptr);
+		      OMP_CLAUSE_DECL (node) = ptr;
+		      node2 = build_omp_clause (input_location,
+						OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_MAP_KIND (node2) = OMP_CLAUSE_MAP_TO_PSET;
+		      OMP_CLAUSE_DECL (node2) = decl;
+		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+		      node3 = build_omp_clause (input_location,
+						OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_MAP_KIND (node3) = OMP_CLAUSE_MAP_POINTER;
+		      OMP_CLAUSE_DECL (node3)
+			= gfc_conv_descriptor_data_get (decl);
+		      OMP_CLAUSE_SIZE (node3) = size_int (0);
+		      if (n->sym->attr.pointer)
+			{
+			  stmtblock_t cond_block;
+			  tree size
+			    = gfc_create_var (gfc_array_index_type, NULL);
+			  tree tem, then_b, else_b, zero, cond;
+
+			  gfc_init_block (&cond_block);
+			  tem
+			    = gfc_full_array_size (&cond_block, decl,
+						   GFC_TYPE_ARRAY_RANK (type));
+			  gfc_add_modify (&cond_block, size, tem);
+			  then_b = gfc_finish_block (&cond_block);
+			  gfc_init_block (&cond_block);
+			  zero = build_int_cst (gfc_array_index_type, 0);
+			  gfc_add_modify (&cond_block, size, zero);
+			  else_b = gfc_finish_block (&cond_block);
+			  tem = gfc_conv_descriptor_data_get (decl);
+			  tem = fold_convert (pvoid_type_node, tem);
+			  cond = fold_build2_loc (input_location, NE_EXPR,
+						  boolean_type_node,
+						  tem, null_pointer_node);
+			  gfc_add_expr_to_block (block,
+						 build3_loc (input_location,
+							     COND_EXPR,
+							     void_type_node,
+							     cond, then_b,
+							     else_b));
+			  OMP_CLAUSE_SIZE (node) = size;
+			}
+		      else
+			OMP_CLAUSE_SIZE (node)
+			  = gfc_full_array_size (block, decl,
+						 GFC_TYPE_ARRAY_RANK (type));
+		      tree elemsz
+			= TYPE_SIZE_UNIT (gfc_get_element_type (type));
+		      elemsz = fold_convert (gfc_array_index_type, elemsz);
+		      OMP_CLAUSE_SIZE (node)
+			= fold_build2 (MULT_EXPR, gfc_array_index_type,
+				       OMP_CLAUSE_SIZE (node), elemsz);
+		    }
+		  else
+		    OMP_CLAUSE_DECL (node) = decl;
+		}
+	      else
+		{
+		  tree ptr, ptr2;
+		  gfc_init_se (&se, NULL);
+		  if (n->expr->ref->u.ar.type == AR_ELEMENT)
+		    {
+		      gfc_conv_expr_reference (&se, n->expr);
+		      gfc_add_block_to_block (block, &se.pre);
+		      ptr = se.expr;
+		      OMP_CLAUSE_SIZE (node)
+			= TYPE_SIZE_UNIT (TREE_TYPE (ptr));
+		    }
+		  else
+		    {
+		      gfc_conv_expr_descriptor (&se, n->expr);
+		      ptr = gfc_conv_array_data (se.expr);
+		      tree type = TREE_TYPE (se.expr);
+		      gfc_add_block_to_block (block, &se.pre);
+		      OMP_CLAUSE_SIZE (node)
+			= gfc_full_array_size (block, se.expr,
+					       GFC_TYPE_ARRAY_RANK (type));
+		      tree elemsz
+			= TYPE_SIZE_UNIT (gfc_get_element_type (type));
+		      elemsz = fold_convert (gfc_array_index_type, elemsz);
+		      OMP_CLAUSE_SIZE (node)
+			= fold_build2 (MULT_EXPR, gfc_array_index_type,
+				       OMP_CLAUSE_SIZE (node), elemsz);
+		    }
+		  gfc_add_block_to_block (block, &se.post);
+		  ptr = fold_convert (build_pointer_type (char_type_node),
+				      ptr);
+		  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+
+		  if (POINTER_TYPE_P (TREE_TYPE (decl))
+		      && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
+		    {
+		      node4 = build_omp_clause (input_location,
+						OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_MAP_KIND (node4) = OMP_CLAUSE_MAP_POINTER;
+		      OMP_CLAUSE_DECL (node4) = decl;
+		      OMP_CLAUSE_SIZE (node4) = size_int (0);
+		      decl = build_fold_indirect_ref (decl);
+		    }
+		  ptr = fold_convert (sizetype, ptr);
+		  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+		    {
+		      tree type = TREE_TYPE (decl);
+		      ptr2 = gfc_conv_descriptor_data_get (decl);
+		      node2 = build_omp_clause (input_location,
+						OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_MAP_KIND (node2) = OMP_CLAUSE_MAP_TO_PSET;
+		      OMP_CLAUSE_DECL (node2) = decl;
+		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+		      node3 = build_omp_clause (input_location,
+						OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_MAP_KIND (node3) = OMP_CLAUSE_MAP_POINTER;
+		      OMP_CLAUSE_DECL (node3)
+			= gfc_conv_descriptor_data_get (decl);
+		    }
+		  else
+		    {
+		      if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+			ptr2 = build_fold_addr_expr (decl);
+		      else
+			{
+			  gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
+			  ptr2 = decl;
+			}
+		      node3 = build_omp_clause (input_location,
+						OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_MAP_KIND (node3) = OMP_CLAUSE_MAP_POINTER;
+		      OMP_CLAUSE_DECL (node3) = decl;
+		    }
+		  ptr2 = fold_convert (sizetype, ptr2);
+		  OMP_CLAUSE_SIZE (node3)
+		    = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+		}
+	      switch (n->u.map_op)
+		{
+		case OMP_MAP_ALLOC:
+		  OMP_CLAUSE_MAP_KIND (node) = OMP_CLAUSE_MAP_ALLOC;
+		  break;
+		case OMP_MAP_TO:
+		  OMP_CLAUSE_MAP_KIND (node) = OMP_CLAUSE_MAP_TO;
+		  break;
+		case OMP_MAP_FROM:
+		  OMP_CLAUSE_MAP_KIND (node) = OMP_CLAUSE_MAP_FROM;
+		  break;
+		case OMP_MAP_TOFROM:
+		  OMP_CLAUSE_MAP_KIND (node) = OMP_CLAUSE_MAP_TOFROM;
+		  break;
+		default:
+		  gcc_unreachable ();
+		}
+	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+	      if (node2)
+		omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
+	      if (node3)
+		omp_clauses = gfc_trans_add_clause (node3, omp_clauses);
+	      if (node4)
+		omp_clauses = gfc_trans_add_clause (node4, omp_clauses);
+	    }
+	  break;
+	case OMP_LIST_TO:
+	case OMP_LIST_FROM:
+	  for (; n != NULL; n = n->next)
+	    {
+	      if (!n->sym->attr.referenced)
+		continue;
+
+	      tree node = build_omp_clause (input_location,
+					    list == OMP_LIST_TO
+					    ? OMP_CLAUSE_TO : OMP_CLAUSE_FROM);
+	      if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
+		{
+		  tree decl = gfc_get_symbol_decl (n->sym);
+		  if (gfc_omp_privatize_by_reference (decl))
+		    decl = build_fold_indirect_ref (decl);
+		  if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+		    {
+		      tree type = TREE_TYPE (decl);
+		      tree ptr = gfc_conv_descriptor_data_get (decl);
+		      ptr = fold_convert (build_pointer_type (char_type_node),
+					  ptr);
+		      ptr = build_fold_indirect_ref (ptr);
+		      OMP_CLAUSE_DECL (node) = ptr;
+		      OMP_CLAUSE_SIZE (node)
+			= gfc_full_array_size (block, decl,
+					       GFC_TYPE_ARRAY_RANK (type));
+		      tree elemsz
+			= TYPE_SIZE_UNIT (gfc_get_element_type (type));
+		      elemsz = fold_convert (gfc_array_index_type, elemsz);
+		      OMP_CLAUSE_SIZE (node)
+			= fold_build2 (MULT_EXPR, gfc_array_index_type,
+				       OMP_CLAUSE_SIZE (node), elemsz);
+		    }
+		  else
+		    OMP_CLAUSE_DECL (node) = decl;
+		}
+	      else
+		{
+		  tree ptr;
+		  gfc_init_se (&se, NULL);
+		  if (n->expr->ref->u.ar.type == AR_ELEMENT)
+		    {
+		      gfc_conv_expr_reference (&se, n->expr);
+		      ptr = se.expr;
+		      gfc_add_block_to_block (block, &se.pre);
+		      OMP_CLAUSE_SIZE (node)
+			= TYPE_SIZE_UNIT (TREE_TYPE (ptr));
+		    }
+		  else
+		    {
+		      gfc_conv_expr_descriptor (&se, n->expr);
+		      ptr = gfc_conv_array_data (se.expr);
+		      tree type = TREE_TYPE (se.expr);
+		      gfc_add_block_to_block (block, &se.pre);
+		      OMP_CLAUSE_SIZE (node)
+			= gfc_full_array_size (block, se.expr,
+					       GFC_TYPE_ARRAY_RANK (type));
+		      tree elemsz
+			= TYPE_SIZE_UNIT (gfc_get_element_type (type));
+		      elemsz = fold_convert (gfc_array_index_type, elemsz);
+		      OMP_CLAUSE_SIZE (node)
+			= fold_build2 (MULT_EXPR, gfc_array_index_type,
+				       OMP_CLAUSE_SIZE (node), elemsz);
+		    }
+		  gfc_add_block_to_block (block, &se.post);
+		  ptr = fold_convert (build_pointer_type (char_type_node),
+				      ptr);
+		  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
 		}
-	      OMP_CLAUSE_DEPEND_KIND (node)
-		= ((list == OMP_LIST_DEPEND_IN)
-		   ? OMP_CLAUSE_DEPEND_IN : OMP_CLAUSE_DEPEND_OUT);
 	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
 	    }
 	  break;
@@ -1920,7 +2306,69 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
 
-  return omp_clauses;
+  if (clauses->num_teams)
+    {
+      tree num_teams;
+
+      gfc_init_se (&se, NULL);
+      gfc_conv_expr (&se, clauses->num_teams);
+      gfc_add_block_to_block (block, &se.pre);
+      num_teams = gfc_evaluate_now (se.expr, block);
+      gfc_add_block_to_block (block, &se.post);
+
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_NUM_TEAMS);
+      OMP_CLAUSE_NUM_TEAMS_EXPR (c) = num_teams;
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+    }
+
+  if (clauses->device)
+    {
+      tree device;
+
+      gfc_init_se (&se, NULL);
+      gfc_conv_expr (&se, clauses->device);
+      gfc_add_block_to_block (block, &se.pre);
+      device = gfc_evaluate_now (se.expr, block);
+      gfc_add_block_to_block (block, &se.post);
+
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_DEVICE);
+      OMP_CLAUSE_DEVICE_ID (c) = device;
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+    }
+
+  if (clauses->thread_limit)
+    {
+      tree thread_limit;
+
+      gfc_init_se (&se, NULL);
+      gfc_conv_expr (&se, clauses->thread_limit);
+      gfc_add_block_to_block (block, &se.pre);
+      thread_limit = gfc_evaluate_now (se.expr, block);
+      gfc_add_block_to_block (block, &se.post);
+
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_THREAD_LIMIT);
+      OMP_CLAUSE_THREAD_LIMIT_EXPR (c) = thread_limit;
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+    }
+
+  chunk_size = NULL_TREE;
+  if (clauses->dist_chunk_size)
+    {
+      gfc_init_se (&se, NULL);
+      gfc_conv_expr (&se, clauses->dist_chunk_size);
+      gfc_add_block_to_block (block, &se.pre);
+      chunk_size = gfc_evaluate_now (se.expr, block);
+      gfc_add_block_to_block (block, &se.post);
+    }
+
+  if (clauses->dist_sched_kind != OMP_SCHED_NONE)
+    {
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_DIST_SCHEDULE);
+      OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (c) = chunk_size;
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+    }
+
+  return nreverse (omp_clauses);
 }
 
 /* Like gfc_trans_code, but force creation of a BIND_EXPR around it.  */
@@ -2329,12 +2777,13 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock,
 
       if (clauses)
 	{
-	  gfc_omp_namelist *n;
-	  for (n = clauses->lists[(op == EXEC_OMP_SIMD && collapse == 1)
-				  ? OMP_LIST_LINEAR : OMP_LIST_LASTPRIVATE];
-	       n != NULL; n = n->next)
-	    if (code->ext.iterator->var->symtree->n.sym == n->sym)
-	      break;
+	  gfc_omp_namelist *n = NULL;
+	  if (op != EXEC_OMP_DISTRIBUTE)
+	    for (n = clauses->lists[(op == EXEC_OMP_SIMD && collapse == 1)
+				    ? OMP_LIST_LINEAR : OMP_LIST_LASTPRIVATE];
+		 n != NULL; n = n->next)
+	      if (code->ext.iterator->var->symtree->n.sym == n->sym)
+		break;
 	  if (n != NULL)
 	    dovar_found = 1;
 	  else if (n == NULL && op != EXEC_OMP_SIMD)
@@ -2554,7 +3003,13 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock,
     }
 
   /* End of loop body.  */
-  stmt = make_node (op == EXEC_OMP_SIMD ? OMP_SIMD : OMP_FOR);
+  switch (op)
+    {
+    case EXEC_OMP_SIMD: stmt = make_node (OMP_SIMD); break;
+    case EXEC_OMP_DO: stmt = make_node (OMP_FOR); break;
+    case EXEC_OMP_DISTRIBUTE: stmt = make_node (OMP_DISTRIBUTE); break;
+    default: gcc_unreachable ();
+    }
 
   TREE_TYPE (stmt) = void_type_node;
   OMP_FOR_BODY (stmt) = gfc_finish_block (&body);
@@ -2610,6 +3065,9 @@ enum
   GFC_OMP_SPLIT_SIMD,
   GFC_OMP_SPLIT_DO,
   GFC_OMP_SPLIT_PARALLEL,
+  GFC_OMP_SPLIT_DISTRIBUTE,
+  GFC_OMP_SPLIT_TEAMS,
+  GFC_OMP_SPLIT_TARGET,
   GFC_OMP_SPLIT_NUM
 };
 
@@ -2617,7 +3075,10 @@ enum
 {
   GFC_OMP_MASK_SIMD = (1 << GFC_OMP_SPLIT_SIMD),
   GFC_OMP_MASK_DO = (1 << GFC_OMP_SPLIT_DO),
-  GFC_OMP_MASK_PARALLEL = (1 << GFC_OMP_SPLIT_PARALLEL)
+  GFC_OMP_MASK_PARALLEL = (1 << GFC_OMP_SPLIT_PARALLEL),
+  GFC_OMP_MASK_DISTRIBUTE = (1 << GFC_OMP_SPLIT_DISTRIBUTE),
+  GFC_OMP_MASK_TEAMS = (1 << GFC_OMP_SPLIT_TEAMS),
+  GFC_OMP_MASK_TARGET = (1 << GFC_OMP_SPLIT_TARGET)
 };
 
 static void
@@ -2628,10 +3089,32 @@ gfc_split_omp_clauses (gfc_code *code,
   memset (clausesa, 0, GFC_OMP_SPLIT_NUM * sizeof (gfc_omp_clauses));
   switch (code->op)
     {
+    case EXEC_OMP_DISTRIBUTE:
+      innermost = GFC_OMP_SPLIT_DISTRIBUTE;
+      break;
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+      mask = GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO;
+      innermost = GFC_OMP_SPLIT_DO;
+      break;
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+      mask = GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_PARALLEL
+	     | GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
+      innermost = GFC_OMP_SPLIT_SIMD;
+      break;
+    case EXEC_OMP_DISTRIBUTE_SIMD:
+      mask = GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_SIMD;
+      innermost = GFC_OMP_SPLIT_SIMD;
+      break;
+    case EXEC_OMP_DO:
+      innermost = GFC_OMP_SPLIT_DO;
+      break;
     case EXEC_OMP_DO_SIMD:
       mask = GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
       innermost = GFC_OMP_SPLIT_SIMD;
       break;
+    case EXEC_OMP_PARALLEL:
+      innermost = GFC_OMP_SPLIT_PARALLEL;
+      break;
     case EXEC_OMP_PARALLEL_DO:
       mask = GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO;
       innermost = GFC_OMP_SPLIT_DO;
@@ -2640,11 +3123,99 @@ gfc_split_omp_clauses (gfc_code *code,
       mask = GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
       innermost = GFC_OMP_SPLIT_SIMD;
       break;
+    case EXEC_OMP_SIMD:
+      innermost = GFC_OMP_SPLIT_SIMD;
+      break;
+    case EXEC_OMP_TARGET:
+      innermost = GFC_OMP_SPLIT_TARGET;
+      break;
+    case EXEC_OMP_TARGET_TEAMS:
+      mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS;
+      innermost = GFC_OMP_SPLIT_TEAMS;
+      break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+      mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS
+	     | GFC_OMP_MASK_DISTRIBUTE;
+      innermost = GFC_OMP_SPLIT_DISTRIBUTE;
+      break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE
+	     | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO;
+      innermost = GFC_OMP_SPLIT_DO;
+      break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE
+	     | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
+      innermost = GFC_OMP_SPLIT_SIMD;
+      break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+      mask = GFC_OMP_MASK_TARGET | GFC_OMP_MASK_TEAMS
+	     | GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_SIMD;
+      innermost = GFC_OMP_SPLIT_SIMD;
+      break;
+    case EXEC_OMP_TEAMS:
+      innermost = GFC_OMP_SPLIT_TEAMS;
+      break;
+    case EXEC_OMP_TEAMS_DISTRIBUTE:
+      mask = GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE;
+      innermost = GFC_OMP_SPLIT_DISTRIBUTE;
+      break;
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      mask = GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE
+	     | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO;
+      innermost = GFC_OMP_SPLIT_DO;
+      break;
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      mask = GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE
+	     | GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO | GFC_OMP_MASK_SIMD;
+      innermost = GFC_OMP_SPLIT_SIMD;
+      break;
+    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+      mask = GFC_OMP_MASK_TEAMS | GFC_OMP_MASK_DISTRIBUTE | GFC_OMP_MASK_SIMD;
+      innermost = GFC_OMP_SPLIT_SIMD;
+      break;
     default:
       gcc_unreachable ();
     }
+  if (mask == 0)
+    {
+      clausesa[innermost] = *code->ext.omp_clauses;
+      return;
+    }
   if (code->ext.omp_clauses != NULL)
     {
+      if (mask & GFC_OMP_MASK_TARGET)
+	{
+	  /* First the clauses that are unique to some constructs.  */
+	  clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_MAP]
+	    = code->ext.omp_clauses->lists[OMP_LIST_MAP];
+	  clausesa[GFC_OMP_SPLIT_TARGET].device
+	    = code->ext.omp_clauses->device;
+	}
+      if (mask & GFC_OMP_MASK_TEAMS)
+	{
+	  /* First the clauses that are unique to some constructs.  */
+	  clausesa[GFC_OMP_SPLIT_TEAMS].num_teams
+	    = code->ext.omp_clauses->num_teams;
+	  clausesa[GFC_OMP_SPLIT_TEAMS].thread_limit
+	    = code->ext.omp_clauses->thread_limit;
+	  /* Shared and default clauses are allowed on parallel and teams.  */
+	  clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_SHARED]
+	    = code->ext.omp_clauses->lists[OMP_LIST_SHARED];
+	  clausesa[GFC_OMP_SPLIT_TEAMS].default_sharing
+	    = code->ext.omp_clauses->default_sharing;
+	}
+      if (mask & GFC_OMP_MASK_DISTRIBUTE)
+	{
+	  /* First the clauses that are unique to some constructs.  */
+	  clausesa[GFC_OMP_SPLIT_DISTRIBUTE].dist_sched_kind
+	    = code->ext.omp_clauses->dist_sched_kind;
+	  clausesa[GFC_OMP_SPLIT_DISTRIBUTE].dist_chunk_size
+	    = code->ext.omp_clauses->dist_chunk_size;
+	  /* Duplicate collapse.  */
+	  clausesa[GFC_OMP_SPLIT_DISTRIBUTE].collapse
+	    = code->ext.omp_clauses->collapse;
+	}
       if (mask & GFC_OMP_MASK_PARALLEL)
 	{
 	  /* First the clauses that are unique to some constructs.  */
@@ -2659,9 +3230,6 @@ gfc_split_omp_clauses (gfc_code *code,
 	    = code->ext.omp_clauses->lists[OMP_LIST_SHARED];
 	  clausesa[GFC_OMP_SPLIT_PARALLEL].default_sharing
 	    = code->ext.omp_clauses->default_sharing;
-	  /* FIXME: This is currently being discussed.  */
-	  clausesa[GFC_OMP_SPLIT_PARALLEL].if_expr
-	    = code->ext.omp_clauses->if_expr;
 	}
       if (mask & GFC_OMP_MASK_DO)
 	{
@@ -2701,6 +3269,12 @@ gfc_split_omp_clauses (gfc_code *code,
       /* Firstprivate clause is supported on all constructs but
 	 target and simd.  Put it on the outermost of those and
 	 duplicate on parallel.  */
+      if (mask & GFC_OMP_MASK_TEAMS)
+	clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_FIRSTPRIVATE]
+	  = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE];
+      else if (mask & GFC_OMP_MASK_DISTRIBUTE)
+	clausesa[GFC_OMP_SPLIT_DISTRIBUTE].lists[OMP_LIST_FIRSTPRIVATE]
+	  = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE];
       if (mask & GFC_OMP_MASK_PARALLEL)
 	clausesa[GFC_OMP_SPLIT_PARALLEL].lists[OMP_LIST_FIRSTPRIVATE]
 	  = code->ext.omp_clauses->lists[OMP_LIST_FIRSTPRIVATE];
@@ -2722,6 +3296,9 @@ gfc_split_omp_clauses (gfc_code *code,
       /* Reduction is allowed on simd, do, parallel and teams.
 	 Duplicate it on all of them, but omit on do if
 	 parallel is present.  */
+      if (mask & GFC_OMP_MASK_TEAMS)
+	clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_REDUCTION]
+	  = code->ext.omp_clauses->lists[OMP_LIST_REDUCTION];
       if (mask & GFC_OMP_MASK_PARALLEL)
 	clausesa[GFC_OMP_SPLIT_PARALLEL].lists[OMP_LIST_REDUCTION]
 	  = code->ext.omp_clauses->lists[OMP_LIST_REDUCTION];
@@ -2731,6 +3308,13 @@ gfc_split_omp_clauses (gfc_code *code,
       if (mask & GFC_OMP_MASK_SIMD)
 	clausesa[GFC_OMP_SPLIT_SIMD].lists[OMP_LIST_REDUCTION]
 	  = code->ext.omp_clauses->lists[OMP_LIST_REDUCTION];
+      /* FIXME: This is currently being discussed.  */
+      if (mask & GFC_OMP_MASK_PARALLEL)
+	clausesa[GFC_OMP_SPLIT_PARALLEL].if_expr
+	  = code->ext.omp_clauses->if_expr;
+      else
+	clausesa[GFC_OMP_SPLIT_TARGET].if_expr
+	  = code->ext.omp_clauses->if_expr;
     }
   if ((mask & (GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO))
       == (GFC_OMP_MASK_PARALLEL | GFC_OMP_MASK_DO))
@@ -2738,14 +3322,17 @@ gfc_split_omp_clauses (gfc_code *code,
 }
 
 static tree
-gfc_trans_omp_do_simd (gfc_code *code, gfc_omp_clauses *clausesa,
-		       tree omp_clauses)
+gfc_trans_omp_do_simd (gfc_code *code, stmtblock_t *pblock,
+		       gfc_omp_clauses *clausesa, tree omp_clauses)
 {
-  stmtblock_t block, *pblock = NULL;
+  stmtblock_t block;
   gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
   tree stmt, body, omp_do_clauses = NULL_TREE;
 
-  gfc_start_block (&block);
+  if (pblock == NULL)
+    gfc_start_block (&block);
+  else
+    gfc_init_block (&block);
 
   if (clausesa == NULL)
     {
@@ -2755,13 +3342,17 @@ gfc_trans_omp_do_simd (gfc_code *code, gfc_omp_clauses *clausesa,
   if (gfc_option.gfc_flag_openmp)
     omp_do_clauses
       = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_DO], code->loc);
-  pblock = &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, gfc_omp_clauses *clausesa,
 }
 
 static tree
-gfc_trans_omp_parallel_do (gfc_code *code)
+gfc_trans_omp_parallel_do (gfc_code *code, stmtblock_t *pblock,
+			   gfc_omp_clauses *clausesa)
 {
-  stmtblock_t block, *pblock = NULL;
-  gfc_omp_clauses clausesa[GFC_OMP_SPLIT_NUM];
+  stmtblock_t block, *new_pblock = pblock;
+  gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
   tree stmt, omp_clauses = NULL_TREE;
 
-  gfc_start_block (&block);
+  if (pblock == NULL)
+    gfc_start_block (&block);
+  else
+    gfc_init_block (&block);
 
-  gfc_split_omp_clauses (code, clausesa);
+  if (clausesa == NULL)
+    {
+      clausesa = clausesa_buf;
+      gfc_split_omp_clauses (code, clausesa);
+    }
   omp_clauses
     = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_PARALLEL],
 			     code->loc);
-  if (!clausesa[GFC_OMP_SPLIT_DO].ordered
-      && clausesa[GFC_OMP_SPLIT_DO].sched_kind != OMP_SCHED_STATIC)
-    pblock = &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 *code)
 }
 
 static tree
-gfc_trans_omp_parallel_do_simd (gfc_code *code)
+gfc_trans_omp_parallel_do_simd (gfc_code *code, stmtblock_t *pblock,
+				gfc_omp_clauses *clausesa)
 {
   stmtblock_t block;
-  gfc_omp_clauses clausesa[GFC_OMP_SPLIT_NUM];
+  gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
   tree stmt, omp_clauses = NULL_TREE;
 
-  gfc_start_block (&block);
+  if (pblock == NULL)
+    gfc_start_block (&block);
+  else
+    gfc_init_block (&block);
 
-  gfc_split_omp_clauses (code, clausesa);
+  if (clausesa == NULL)
+    {
+      clausesa = clausesa_buf;
+      gfc_split_omp_clauses (code, clausesa);
+    }
   if (gfc_option.gfc_flag_openmp)
     omp_clauses
       = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_PARALLEL],
 			       code->loc);
-  pushlevel ();
-  stmt = gfc_trans_omp_do_simd (code, clausesa, omp_clauses);
-  if (TREE_CODE (stmt) != BIND_EXPR)
-    stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
-  else
-    poplevel (0, 0);
+  if (pblock == NULL)
+    pushlevel ();
+  stmt = gfc_trans_omp_do_simd (code, pblock, clausesa, omp_clauses);
+  if (pblock == NULL)
+    {
+      if (TREE_CODE (stmt) != BIND_EXPR)
+	stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+      else
+	poplevel (0, 0);
+    }
+  else if (TREE_CODE (stmt) != BIND_EXPR)
+    stmt = build3_v (BIND_EXPR, NULL, stmt, NULL_TREE);
   if (gfc_option.gfc_flag_openmp)
     {
       stmt = build2_loc (input_location, OMP_PARALLEL, void_type_node, stmt,
@@ -2969,6 +3590,170 @@ gfc_trans_omp_taskyield (void)
 }
 
 static tree
+gfc_trans_omp_distribute (gfc_code *code, gfc_omp_clauses *clausesa)
+{
+  stmtblock_t block;
+  gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
+  tree stmt, omp_clauses = NULL_TREE;
+
+  gfc_start_block (&block);
+  if (clausesa == NULL)
+    {
+      clausesa = clausesa_buf;
+      gfc_split_omp_clauses (code, clausesa);
+    }
+  if (gfc_option.gfc_flag_openmp)
+    omp_clauses
+      = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_DISTRIBUTE],
+			       code->loc);
+  switch (code->op)
+    {
+    case EXEC_OMP_DISTRIBUTE:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TEAMS_DISTRIBUTE:
+      /* This is handled in gfc_trans_omp_do.  */
+      gcc_unreachable ();
+      break;
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+      stmt = gfc_trans_omp_parallel_do (code, &block, clausesa);
+      if (TREE_CODE (stmt) != BIND_EXPR)
+	stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+      else
+	poplevel (0, 0);
+      break;
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+      stmt = gfc_trans_omp_parallel_do_simd (code, &block, clausesa);
+      if (TREE_CODE (stmt) != BIND_EXPR)
+	stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+      else
+	poplevel (0, 0);
+      break;
+    case EXEC_OMP_DISTRIBUTE_SIMD:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+      stmt = gfc_trans_omp_do (code, EXEC_OMP_SIMD, &block,
+			       &clausesa[GFC_OMP_SPLIT_SIMD], NULL_TREE);
+      if (TREE_CODE (stmt) != BIND_EXPR)
+	stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
+      else
+	poplevel (0, 0);
+      break;
+    default:
+      gcc_unreachable ();
+    }
+  if (gfc_option.gfc_flag_openmp)
+    {
+      tree distribute = make_node (OMP_DISTRIBUTE);
+      TREE_TYPE (distribute) = void_type_node;
+      OMP_FOR_BODY (distribute) = stmt;
+      OMP_FOR_CLAUSES (distribute) = omp_clauses;
+      stmt = distribute;
+    }
+  gfc_add_expr_to_block (&block, stmt);
+  return gfc_finish_block (&block);
+}
+
+static tree
+gfc_trans_omp_teams (gfc_code *code, gfc_omp_clauses *clausesa)
+{
+  stmtblock_t block;
+  gfc_omp_clauses clausesa_buf[GFC_OMP_SPLIT_NUM];
+  tree stmt, omp_clauses = NULL_TREE;
+
+  gfc_start_block (&block);
+  if (clausesa == NULL)
+    {
+      clausesa = clausesa_buf;
+      gfc_split_omp_clauses (code, clausesa);
+    }
+  if (gfc_option.gfc_flag_openmp)
+    omp_clauses
+      = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_TEAMS],
+			       code->loc);
+  switch (code->op)
+    {
+    case EXEC_OMP_TARGET_TEAMS:
+    case EXEC_OMP_TEAMS:
+      stmt = gfc_trans_omp_code (code->block->next, true);
+      break;
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TEAMS_DISTRIBUTE:
+      stmt = gfc_trans_omp_do (code, EXEC_OMP_DISTRIBUTE, NULL,
+			       &clausesa[GFC_OMP_SPLIT_DISTRIBUTE],
+			       NULL);
+      break;
+    default:
+      stmt = gfc_trans_omp_distribute (code, clausesa);
+      break;
+    }
+  stmt = build2_loc (input_location, OMP_TEAMS, void_type_node, stmt,
+		     omp_clauses);
+  gfc_add_expr_to_block (&block, stmt);
+  return gfc_finish_block (&block);
+}
+
+static tree
+gfc_trans_omp_target (gfc_code *code)
+{
+  stmtblock_t block;
+  gfc_omp_clauses clausesa[GFC_OMP_SPLIT_NUM];
+  tree stmt, omp_clauses = NULL_TREE;
+
+  gfc_start_block (&block);
+  gfc_split_omp_clauses (code, clausesa);
+  if (gfc_option.gfc_flag_openmp)
+    omp_clauses
+      = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_TARGET],
+			       code->loc);
+  if (code->op == EXEC_OMP_TARGET)
+    stmt = gfc_trans_omp_code (code->block->next, true);
+  else
+    stmt = gfc_trans_omp_teams (code, clausesa);
+  if (TREE_CODE (stmt) != BIND_EXPR)
+    stmt = build3_v (BIND_EXPR, NULL, stmt, NULL_TREE);
+  if (gfc_option.gfc_flag_openmp)
+    stmt = build2_loc (input_location, OMP_TARGET, void_type_node, stmt,
+		       omp_clauses);
+  gfc_add_expr_to_block (&block, stmt);
+  return gfc_finish_block (&block);
+}
+
+static tree
+gfc_trans_omp_target_data (gfc_code *code)
+{
+  stmtblock_t block;
+  tree stmt, omp_clauses;
+
+  gfc_start_block (&block);
+  omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+				       code->loc);
+  stmt = gfc_trans_omp_code (code->block->next, true);
+  stmt = build2_loc (input_location, OMP_TARGET_DATA, void_type_node, stmt,
+		     omp_clauses);
+  gfc_add_expr_to_block (&block, stmt);
+  return gfc_finish_block (&block);
+}
+
+static tree
+gfc_trans_omp_target_update (gfc_code *code)
+{
+  stmtblock_t block;
+  tree stmt, omp_clauses;
+
+  gfc_start_block (&block);
+  omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+				       code->loc);
+  stmt = build1_loc (input_location, OMP_TARGET_UPDATE, void_type_node,
+		     omp_clauses);
+  gfc_add_expr_to_block (&block, stmt);
+  return gfc_finish_block (&block);
+}
+
+static tree
 gfc_trans_omp_workshare (gfc_code *code, gfc_omp_clauses *clauses)
 {
   tree res, tmp, stmt;
@@ -3141,12 +3926,17 @@ gfc_trans_omp_directive (gfc_code *code)
       return gfc_trans_omp_cancellation_point (code);
     case EXEC_OMP_CRITICAL:
       return gfc_trans_omp_critical (code);
+    case EXEC_OMP_DISTRIBUTE:
     case EXEC_OMP_DO:
     case EXEC_OMP_SIMD:
       return gfc_trans_omp_do (code, code->op, NULL, code->ext.omp_clauses,
 			       NULL);
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_DISTRIBUTE_SIMD:
+      return gfc_trans_omp_distribute (code, NULL);
     case EXEC_OMP_DO_SIMD:
-      return gfc_trans_omp_do_simd (code, NULL, NULL_TREE);
+      return gfc_trans_omp_do_simd (code, NULL, NULL, NULL_TREE);
     case EXEC_OMP_FLUSH:
       return gfc_trans_omp_flush ();
     case EXEC_OMP_MASTER:
@@ -3156,9 +3946,9 @@ gfc_trans_omp_directive (gfc_code *code)
     case EXEC_OMP_PARALLEL:
       return gfc_trans_omp_parallel (code);
     case EXEC_OMP_PARALLEL_DO:
-      return gfc_trans_omp_parallel_do (code);
+      return gfc_trans_omp_parallel_do (code, NULL, NULL);
     case EXEC_OMP_PARALLEL_DO_SIMD:
-      return gfc_trans_omp_parallel_do_simd (code);
+      return gfc_trans_omp_parallel_do_simd (code, NULL, NULL);
     case EXEC_OMP_PARALLEL_SECTIONS:
       return gfc_trans_omp_parallel_sections (code);
     case EXEC_OMP_PARALLEL_WORKSHARE:
@@ -3167,6 +3957,17 @@ gfc_trans_omp_directive (gfc_code *code)
       return gfc_trans_omp_sections (code, code->ext.omp_clauses);
     case EXEC_OMP_SINGLE:
       return gfc_trans_omp_single (code, code->ext.omp_clauses);
+    case EXEC_OMP_TARGET:
+    case EXEC_OMP_TARGET_TEAMS:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+      return gfc_trans_omp_target (code);
+    case EXEC_OMP_TARGET_DATA:
+      return gfc_trans_omp_target_data (code);
+    case EXEC_OMP_TARGET_UPDATE:
+      return gfc_trans_omp_target_update (code);
     case EXEC_OMP_TASK:
       return gfc_trans_omp_task (code);
     case EXEC_OMP_TASKGROUP:
@@ -3175,6 +3976,12 @@ gfc_trans_omp_directive (gfc_code *code)
       return gfc_trans_omp_taskwait ();
     case EXEC_OMP_TASKYIELD:
       return gfc_trans_omp_taskyield ();
+    case EXEC_OMP_TEAMS:
+    case EXEC_OMP_TEAMS_DISTRIBUTE:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+    case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+      return gfc_trans_omp_teams (code, NULL);
     case EXEC_OMP_WORKSHARE:
       return gfc_trans_omp_workshare (code, code->ext.omp_clauses);
     default:
diff --git a/gcc/fortran/trans.c b/gcc/fortran/trans.c
index cfb8038440b..1925506594c 100644
--- a/gcc/fortran/trans.c
+++ b/gcc/fortran/trans.c
@@ -1851,6 +1851,10 @@ trans_code (gfc_code * code, tree cond)
 	case EXEC_OMP_CANCEL:
 	case EXEC_OMP_CANCELLATION_POINT:
 	case EXEC_OMP_CRITICAL:
+	case EXEC_OMP_DISTRIBUTE:
+	case EXEC_OMP_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_DISTRIBUTE_SIMD:
 	case EXEC_OMP_DO:
 	case EXEC_OMP_DO_SIMD:
 	case EXEC_OMP_FLUSH:
@@ -1864,10 +1868,23 @@ trans_code (gfc_code * code, tree cond)
 	case EXEC_OMP_SECTIONS:
 	case EXEC_OMP_SIMD:
 	case EXEC_OMP_SINGLE:
+	case EXEC_OMP_TARGET:
+	case EXEC_OMP_TARGET_DATA:
+	case EXEC_OMP_TARGET_TEAMS:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+	case EXEC_OMP_TARGET_UPDATE:
 	case EXEC_OMP_TASK:
 	case EXEC_OMP_TASKGROUP:
 	case EXEC_OMP_TASKWAIT:
 	case EXEC_OMP_TASKYIELD:
+	case EXEC_OMP_TEAMS:
+	case EXEC_OMP_TEAMS_DISTRIBUTE:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+	case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
 	case EXEC_OMP_WORKSHARE:
 	  res = gfc_trans_omp_directive (code);
 	  break;
diff --git a/gcc/fortran/trans.h b/gcc/fortran/trans.h
index 7ab9dd4feed..c272c0d819e 100644
--- a/gcc/fortran/trans.h
+++ b/gcc/fortran/trans.h
@@ -671,6 +671,7 @@ tree gfc_omp_clause_default_ctor (tree, tree, tree);
 tree gfc_omp_clause_copy_ctor (tree, tree, tree);
 tree gfc_omp_clause_assign_op (tree, tree, tree);
 tree gfc_omp_clause_dtor (tree, tree);
+void gfc_omp_finish_clause (tree, gimple_seq *);
 bool gfc_omp_disregard_value_expr (tree, bool);
 bool gfc_omp_private_debug_clause (tree, bool);
 bool gfc_omp_private_outer_ref (tree);
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 10f8ac6d0e0..2efc8992008 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -5650,6 +5650,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
   n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
   if (ctx->region_type == ORT_TARGET)
     {
+      ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
       if (n == NULL)
 	{
 	  if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
@@ -5662,8 +5663,12 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	    omp_add_variable (ctx, decl, GOVD_MAP | flags);
 	}
       else
-	n->value |= flags;
-      ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
+	{
+	  /* If nothing changed, there's nothing left to do.  */
+	  if ((n->value & flags) == flags)
+	    return ret;
+	  n->value |= flags;
+	}
       goto do_outer;
     }
 
@@ -6201,13 +6206,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   gimplify_omp_ctxp = ctx;
 }
 
+struct gimplify_adjust_omp_clauses_data
+{
+  tree *list_p;
+  gimple_seq *pre_p;
+};
+
 /* For all variables that were not actually used within the context,
    remove PRIVATE, SHARED, and FIRSTPRIVATE clauses.  */
 
 static int
 gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 {
-  tree *list_p = (tree *) data;
+  tree *list_p = ((struct gimplify_adjust_omp_clauses_data *) data)->list_p;
+  gimple_seq *pre_p
+    = ((struct gimplify_adjust_omp_clauses_data *) data)->pre_p;
   tree decl = (tree) n->key;
   unsigned flags = n->value;
   enum omp_clause_code code;
@@ -6308,15 +6321,21 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
       OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (nc) = 1;
       OMP_CLAUSE_CHAIN (nc) = *list_p;
       OMP_CLAUSE_CHAIN (clause) = nc;
-      lang_hooks.decls.omp_finish_clause (nc);
+      struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+      gimplify_omp_ctxp = ctx->outer_context;
+      lang_hooks.decls.omp_finish_clause (nc, pre_p);
+      gimplify_omp_ctxp = ctx;
     }
   *list_p = clause;
-  lang_hooks.decls.omp_finish_clause (clause);
+  struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+  gimplify_omp_ctxp = ctx->outer_context;
+  lang_hooks.decls.omp_finish_clause (clause, pre_p);
+  gimplify_omp_ctxp = ctx;
   return 0;
 }
 
 static void
-gimplify_adjust_omp_clauses (tree *list_p)
+gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
 {
   struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
   tree c, decl;
@@ -6521,7 +6540,10 @@ gimplify_adjust_omp_clauses (tree *list_p)
     }
 
   /* Add in any implicit data sharing.  */
-  splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, list_p);
+  struct gimplify_adjust_omp_clauses_data data;
+  data.list_p = list_p;
+  data.pre_p = pre_p;
+  splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
 
   gimplify_omp_ctxp = ctx->outer_context;
   delete_omp_context (ctx);
@@ -6552,7 +6574,7 @@ gimplify_omp_parallel (tree *expr_p, gimple_seq *pre_p)
   else
     pop_gimplify_context (NULL);
 
-  gimplify_adjust_omp_clauses (&OMP_PARALLEL_CLAUSES (expr));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_PARALLEL_CLAUSES (expr));
 
   g = gimple_build_omp_parallel (body,
 				 OMP_PARALLEL_CLAUSES (expr),
@@ -6588,7 +6610,7 @@ gimplify_omp_task (tree *expr_p, gimple_seq *pre_p)
   else
     pop_gimplify_context (NULL);
 
-  gimplify_adjust_omp_clauses (&OMP_TASK_CLAUSES (expr));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_TASK_CLAUSES (expr));
 
   g = gimple_build_omp_task (body,
 			     OMP_TASK_CLAUSES (expr),
@@ -6934,7 +6956,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	TREE_OPERAND (TREE_OPERAND (t, 1), 0) = var;
       }
 
-  gimplify_adjust_omp_clauses (&OMP_FOR_CLAUSES (orig_for_stmt));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_FOR_CLAUSES (orig_for_stmt));
 
   int kind;
   switch (TREE_CODE (orig_for_stmt))
@@ -7034,7 +7056,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
     }
   else
     gimplify_and_add (OMP_BODY (expr), &body);
-  gimplify_adjust_omp_clauses (&OMP_CLAUSES (expr));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_CLAUSES (expr));
 
   switch (TREE_CODE (expr))
     {
@@ -7073,7 +7095,7 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 
   gimplify_scan_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr), pre_p,
 			     ORT_WORKSHARE);
-  gimplify_adjust_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr));
+  gimplify_adjust_omp_clauses (pre_p, &OMP_TARGET_UPDATE_CLAUSES (expr));
   stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_UPDATE,
 				  OMP_TARGET_UPDATE_CLAUSES (expr));
 
diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h
index d9a1dfdc194..76bb907151b 100644
--- a/gcc/langhooks-def.h
+++ b/gcc/langhooks-def.h
@@ -75,6 +75,7 @@ extern bool lhd_handle_option (size_t, const char *, int, int, location_t,
 extern int lhd_gimplify_expr (tree *, gimple_seq *, gimple_seq *);
 extern enum omp_clause_default_kind lhd_omp_predetermined_sharing (tree);
 extern tree lhd_omp_assignment (tree, tree, tree);
+extern void lhd_omp_finish_clause (tree, gimple_seq *);
 struct gimplify_omp_ctx;
 extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
 					       tree);
@@ -215,7 +216,7 @@ extern tree lhd_make_node (enum tree_code);
 #define LANG_HOOKS_OMP_CLAUSE_COPY_CTOR lhd_omp_assignment
 #define LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP lhd_omp_assignment
 #define LANG_HOOKS_OMP_CLAUSE_DTOR hook_tree_tree_tree_null
-#define LANG_HOOKS_OMP_FINISH_CLAUSE hook_void_tree
+#define LANG_HOOKS_OMP_FINISH_CLAUSE lhd_omp_finish_clause
 
 #define LANG_HOOKS_DECLS { \
   LANG_HOOKS_GLOBAL_BINDINGS_P, \
diff --git a/gcc/langhooks.c b/gcc/langhooks.c
index 8f65c6860a5..add08566cbd 100644
--- a/gcc/langhooks.c
+++ b/gcc/langhooks.c
@@ -515,6 +515,13 @@ lhd_omp_assignment (tree clause ATTRIBUTE_UNUSED, tree dst, tree src)
   return build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, src);
 }
 
+/* Finalize clause C.  */
+
+void
+lhd_omp_finish_clause (tree, gimple_seq *)
+{
+}
+
 /* Register language specific type size variables as potentially OpenMP
    firstprivate variables.  */
 
diff --git a/gcc/langhooks.h b/gcc/langhooks.h
index 35b47bc6574..33aa55833b3 100644
--- a/gcc/langhooks.h
+++ b/gcc/langhooks.h
@@ -230,7 +230,7 @@ struct lang_hooks_for_decls
   tree (*omp_clause_dtor) (tree clause, tree decl);
 
   /* Do language specific checking on an implicitly determined clause.  */
-  void (*omp_finish_clause) (tree clause);
+  void (*omp_finish_clause) (tree clause, gimple_seq *pre_p);
 };
 
 /* Language hooks related to LTO serialization.  */
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 67254cc5bdf..a30ce5aa135 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1678,6 +1678,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 		}
 	      else
 		{
+		  if (ctx->outer)
+		    {
+		      scan_omp_op (&OMP_CLAUSE_DECL (c), ctx->outer);
+		      decl = OMP_CLAUSE_DECL (c);
+		    }
 		  gcc_assert (!splay_tree_lookup (ctx->field_map,
 						  (splay_tree_key) decl));
 		  tree field
@@ -2011,6 +2016,7 @@ scan_omp_parallel (gimple_stmt_iterator *gsi, omp_context *outer_ctx)
 	      tree temp = create_tmp_var (type, NULL);
 	      tree c = build_omp_clause (UNKNOWN_LOCATION,
 					 OMP_CLAUSE__LOOPTEMP_);
+	      insert_decl_map (&outer_ctx->cb, temp, temp);
 	      OMP_CLAUSE_DECL (c) = temp;
 	      OMP_CLAUSE_CHAIN (c) = gimple_omp_parallel_clauses (stmt);
 	      gimple_omp_parallel_set_clauses (stmt, c);
@@ -2508,6 +2514,23 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 	  return false;
 	}
       break;
+    case GIMPLE_OMP_TARGET:
+      for (; ctx != NULL; ctx = ctx->outer)
+	if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+	    && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION)
+	  {
+	    const char *name;
+	    switch (gimple_omp_target_kind (stmt))
+	      {
+	      case GF_OMP_TARGET_KIND_REGION: name = "target"; break;
+	      case GF_OMP_TARGET_KIND_DATA: name = "target data"; break;
+	      case GF_OMP_TARGET_KIND_UPDATE: name = "target update"; break;
+	      default: gcc_unreachable ();
+	      }
+	    warning_at (gimple_location (stmt), 0,
+			"%s construct inside of target region", name);
+	  }
+      break;
     default:
       break;
     }
@@ -9041,7 +9064,10 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 					OMP_CLAUSE__LOOPTEMP_);
 	    }
 	  else
-	    temp = create_tmp_var (type, NULL);
+	    {
+	      temp = create_tmp_var (type, NULL);
+	      insert_decl_map (&ctx->outer->cb, temp, temp);
+	    }
 	  *pc = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE__LOOPTEMP_);
 	  OMP_CLAUSE_DECL (*pc) = temp;
 	  pc = &OMP_CLAUSE_CHAIN (*pc);
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 80c4e2afbd0..9b76f2394ab 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,14 @@
+2014-06-18  Jakub Jelinek  <jakub@redhat.com>
+
+	* gfortran.dg/gomp/declare-simd-1.f90: New test.
+	* gfortran.dg/gomp/depend-1.f90: New test.
+	* gfortran.dg/gomp/target1.f90: New test.
+	* gfortran.dg/gomp/target2.f90: New test.
+	* gfortran.dg/gomp/target3.f90: New test.
+	* gfortran.dg/gomp/udr4.f90: Adjust expected diagnostics.
+	* gfortran.dg/openmp-define-3.f90: Expect _OPENMP 201307 instead of
+	201107.
+
 2014-06-18  Dominique d'Humieres <dominiq@lps.ens.fr>
 
 	PR fortran/61126
diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-simd-1.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-simd-1.f90
new file mode 100644
index 00000000000..d6ae7c9c812
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/declare-simd-1.f90
@@ -0,0 +1,9 @@
+! { dg-do compile }
+
+subroutine fn1 (x)
+  integer :: x
+!$omp declare simd (fn1) inbranch notinbranch uniform (x) ! { dg-error "Unclassifiable OpenMP directive" }
+end subroutine fn1
+subroutine fn2 (x)
+!$omp declare simd (fn100)	! { dg-error "should refer to containing procedure" }
+end subroutine fn2
diff --git a/gcc/testsuite/gfortran.dg/gomp/depend-1.f90 b/gcc/testsuite/gfortran.dg/gomp/depend-1.f90
new file mode 100644
index 00000000000..bd6d26a3830
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/depend-1.f90
@@ -0,0 +1,13 @@
+! { dg-do compile }
+
+subroutine foo (x)
+  integer :: x(5, *)
+!$omp parallel
+!$omp single
+!$omp task depend(in:x(:,5))
+!$omp end task
+!$omp task depend(in:x(5,:))	! { dg-error "Rightmost upper bound of assumed size array section|proper array section" }
+!$omp end task
+!$omp end single
+!$omp end parallel
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target1.f90 b/gcc/testsuite/gfortran.dg/gomp/target1.f90
new file mode 100644
index 00000000000..14db4970bdc
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target1.f90
@@ -0,0 +1,520 @@
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+module target1
+  interface
+    subroutine dosomething (a, n, m)
+      integer :: a (:), n, m
+      !$omp declare target
+    end subroutine dosomething
+  end interface
+contains
+  subroutine foo (n, o, p, q, r, pp)
+    integer :: n, o, p, q, r, s, i, j
+    integer :: a (2:o)
+    integer, pointer :: pp
+  !$omp target data device (n + 1) if (n .ne. 6) map (tofrom: n, r)
+    !$omp target device (n + 1) if (n .ne. 6) map (from: n) map (alloc: a(2:o))
+      call dosomething (a, n, 0)
+    !$omp end target
+    !$omp target teams device (n + 1) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r)
+      r = r + 1
+      p = q
+      call dosomething (a, n, p + q)
+    !$omp end target teams
+    !$omp target teams distribute device (n + 1) num_teams (n + 4) collapse (2) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+      end do
+    !$omp target teams distribute device (n + 1) num_teams (n + 4) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+      end do
+    !$omp end target teams distribute
+    !$omp target teams distribute parallel do device (n + 1) num_teams (n + 4) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+    !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+    !$omp & ordered schedule (static, 8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+	  !$omp ordered
+	    p = q
+	  !$omp end ordered
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp target teams distribute parallel do device (n + 1) num_teams (n + 4) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) num_threads (n + 4) &
+    !$omp & proc_bind (master) lastprivate (s) ordered schedule (static, 8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+        !$omp ordered
+          p = q
+        !$omp end ordered
+	s = i * 10
+      end do
+    !$omp end target teams distribute parallel do
+    !$omp target teams distribute parallel do simd device (n + 1) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+    !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+    !$omp & schedule (static, 8) num_teams (n + 4) safelen(8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          a(2+i*10+j) = p + q
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp target teams distribute parallel do simd device (n + 1) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) num_threads (n + 4) &
+    !$omp & proc_bind (master) lastprivate (s) schedule (static, 8) &
+    !$omp & num_teams (n + 4) safelen(16) linear(i:1) aligned (pp:4)
+      do i = 1, 10
+        r = r + 1
+        p = q
+        a(1+i) = p + q
+	s = i * 10
+      end do
+    !$omp end target teams distribute parallel do simd
+    !$omp target teams distribute simd device (n + 1) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+    !$omp & lastprivate (s) num_teams (n + 4) safelen(8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          a(2+i*10+j) = p + q
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp target teams distribute simd device (n + 1) &
+    !$omp & if (n .ne. 6)map (from: n) map (alloc: a(2:o)) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) lastprivate (s) &
+    !$omp & num_teams (n + 4) safelen(16) linear(i:1) aligned (pp:4)
+      do i = 1, 10
+        r = r + 1
+        p = q
+        a(1+i) = p + q
+	s = i * 10
+      end do
+    !$omp end target teams distribute simd
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams num_teams (n + 4) thread_limit (n * 2) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r)
+      r = r + 1
+      p = q
+      call dosomething (a, n, p + q)
+    !$omp end teams
+    !$omp end target
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams distribute num_teams (n + 4) collapse (2) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+      end do
+    !$omp end target
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams distribute num_teams (n + 4) default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+      end do
+    !$omp end teams distribute
+    !$omp end target
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams distribute parallel do num_teams (n + 4) &
+    !$omp & if (n .ne. 6) default(shared) ordered schedule (static, 8) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+    !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+	  !$omp ordered
+	    p = q
+	  !$omp end ordered
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams distribute parallel do num_teams (n + 4)if(n.ne.6)default(shared)&
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) num_threads (n + 4) &
+    !$omp & proc_bind (master) lastprivate (s) ordered schedule (static, 8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+        !$omp ordered
+          p = q
+        !$omp end ordered
+	s = i * 10
+      end do
+    !$omp end teams distribute parallel do
+    !$omp end target
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams distribute parallel do simd if(n.ne.6)default(shared)&
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+    !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+    !$omp & schedule (static, 8) num_teams (n + 4) safelen(8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          a(2+i*10+j) = p + q
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams distribute parallel do simd if (n .ne. 6)default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) num_threads (n + 4) &
+    !$omp & proc_bind (master) lastprivate (s) schedule (static, 8) &
+    !$omp & num_teams (n + 4) safelen(16) linear(i:1) aligned (pp:4)
+      do i = 1, 10
+        r = r + 1
+        p = q
+        a(1+i) = p + q
+	s = i * 10
+      end do
+    !$omp end teams distribute parallel do simd
+    !$omp end target
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams distribute simd default(shared) safelen(8) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) collapse (2) &
+    !$omp & lastprivate (s) num_teams (n + 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          a(2+i*10+j) = p + q
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target
+    !$omp target device (n + 1) if (n .ne. 6)map (from: n) map (alloc: a(2:o))
+    !$omp teams distribute simd default(shared) aligned (pp:4) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & thread_limit (n * 2) dist_schedule (static, 4) lastprivate (s)
+      do i = 1, 10
+        r = r + 1
+        p = q
+        a(1+i) = p + q
+	s = i * 10
+      end do
+    !$omp end teams distribute simd
+    !$omp end target
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction ( + : r )
+    !$omp distribute collapse (2) firstprivate (q) dist_schedule (static, 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+      end do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute firstprivate (q) dist_schedule (static, 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+      end do
+    !$omp end distribute
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute parallel do if (n .ne. 6) default(shared) &
+    !$omp & ordered schedule (static, 8) private (p) firstprivate (q) &
+    !$omp & shared(n)reduction(+:r)dist_schedule(static,4)collapse(2)&
+    !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+	  !$omp ordered
+	    p = q
+	  !$omp end ordered
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute parallel do if(n.ne.6)default(shared)&
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) num_threads (n + 4) &
+    !$omp & proc_bind (master) lastprivate (s) ordered schedule (static, 8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+        !$omp ordered
+          p = q
+        !$omp end ordered
+	s = i * 10
+      end do
+    !$omp end distribute parallel do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute parallel do simd if(n.ne.6)default(shared)&
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) collapse (2) safelen(8) &
+    !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+    !$omp & schedule (static, 8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          a(2+i*10+j) = p + q
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute parallel do simd if (n .ne. 6)default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) num_threads (n + 4) &
+    !$omp & proc_bind (master) lastprivate (s) schedule (static, 8) &
+    !$omp & safelen(16) linear(i:1) aligned (pp:4)
+      do i = 1, 10
+        r = r + 1
+        p = q
+        a(1+i) = p + q
+	s = i * 10
+      end do
+    !$omp end distribute parallel do simd
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute simd safelen(8) lastprivate(s) &
+    !$omp & private (p) firstprivate (q) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) collapse (2)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          a(2+i*10+j) = p + q
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute simd aligned (pp:4) &
+    !$omp & private (p) firstprivate (q) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) lastprivate (s)
+      do i = 1, 10
+        r = r + 1
+        p = q
+        a(1+i) = p + q
+	s = i * 10
+      end do
+    !$omp end distribute simd
+    !$omp end target teams
+  !$omp end target data
+  end subroutine
+  subroutine bar (n, o, p, r, pp)
+    integer :: n, o, p, q, r, s, i, j
+    integer :: a (2:o)
+    integer, pointer :: pp
+    common /blk/ i, j, q
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction ( + : r )
+    !$omp distribute collapse (2) firstprivate (q) dist_schedule (static, 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+      end do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute firstprivate (q) dist_schedule (static, 4)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+      end do
+    !$omp end distribute
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute parallel do if (n .ne. 6) default(shared) &
+    !$omp & ordered schedule (static, 8) private (p) firstprivate (q) &
+    !$omp & shared(n)reduction(+:r)dist_schedule(static,4)collapse(2)&
+    !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+	  !$omp ordered
+	    p = q
+	  !$omp end ordered
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute parallel do if(n.ne.6)default(shared)&
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) num_threads (n + 4) &
+    !$omp & proc_bind (master) lastprivate (s) ordered schedule (static, 8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          call dosomething (a, n, p + q)
+        end do
+        !$omp ordered
+          p = q
+        !$omp end ordered
+	s = i * 10
+      end do
+    !$omp end distribute parallel do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute parallel do simd if(n.ne.6)default(shared)&
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) collapse (2) safelen(8) &
+    !$omp & num_threads (n + 4) proc_bind (spread) lastprivate (s) &
+    !$omp & schedule (static, 8)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          a(2+i*10+j) = p + q
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute parallel do simd if (n .ne. 6)default(shared) &
+    !$omp & private (p) firstprivate (q) shared (n) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) num_threads (n + 4) &
+    !$omp & proc_bind (master) lastprivate (s) schedule (static, 8) &
+    !$omp & safelen(16) linear(i:1) aligned (pp:4)
+      do i = 1, 10
+        r = r + 1
+        p = q
+        a(1+i) = p + q
+	s = i * 10
+      end do
+    !$omp end distribute parallel do simd
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute simd safelen(8) lastprivate(s) &
+    !$omp & private (p) firstprivate (q) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) collapse (2)
+      do i = 1, 10
+        do j = 1, 10
+          r = r + 1
+          p = q
+          a(2+i*10+j) = p + q
+	  s = i * 10 + j
+        end do
+      end do
+    !$omp end target teams
+    !$omp target teams device (n + 1) if (n .ne. 6)map (from: n) &
+    !$omp & map (alloc: a(2:o)) num_teams (n + 4) thread_limit (n * 2) &
+    !$omp & default(shared) shared(n) private (p) reduction(+:r)
+    !$omp distribute simd aligned (pp:4) &
+    !$omp & private (p) firstprivate (q) reduction (+: r) &
+    !$omp & dist_schedule (static, 4) lastprivate (s)
+      do i = 1, 10
+        r = r + 1
+        p = q
+        a(1+i) = p + q
+	s = i * 10
+      end do
+    !$omp end distribute simd
+    !$omp end target teams
+  end subroutine
+end module
diff --git a/gcc/testsuite/gfortran.dg/gomp/target2.f90 b/gcc/testsuite/gfortran.dg/gomp/target2.f90
new file mode 100644
index 00000000000..7521331fcb1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target2.f90
@@ -0,0 +1,74 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -ffree-line-length-160" }
+
+subroutine foo (n, s, t, u, v, w)
+  integer :: n, i, s, t, u, v, w
+  common /bar/ i
+  !$omp simd safelen(s + 1)
+  do i = 1, n
+  end do
+  !$omp do schedule (static, t * 2)
+  do i = 1, n
+  end do
+  !$omp do simd safelen(s + 1) schedule (static, t * 2)
+  do i = 1, n
+  end do
+  !$omp parallel do schedule (static, t * 2) num_threads (u - 1)
+  do i = 1, n
+  end do
+  !$omp parallel do simd safelen(s + 1) schedule (static, t * 2) num_threads (u - 1)
+  do i = 1, n
+  end do
+  !$omp distribute dist_schedule (static, v + 8)
+  do i = 1, n
+  end do
+  !$omp distribute simd dist_schedule (static, v + 8) safelen(s + 1)
+  do i = 1, n
+  end do
+  !$omp distribute parallel do simd dist_schedule (static, v + 8) safelen(s + 1) &
+  !$omp & schedule (static, t * 2) num_threads (u - 1)
+  do i = 1, n
+  end do
+  !$omp distribute parallel do dist_schedule (static, v + 8) num_threads (u - 1) &
+  !$omp & schedule (static, t * 2)
+  do i = 1, n
+  end do
+  !$omp target
+  !$omp teams distribute dist_schedule (static, v + 8) num_teams (w + 8)
+  do i = 1, n
+  end do
+  !$omp end target
+  !$omp target
+  !$omp teams distribute simd dist_schedule (static, v + 8) safelen(s + 1) &
+  !$omp & num_teams (w + 8)
+  do i = 1, n
+  end do
+  !$omp end target
+  !$omp target
+  !$omp teams distribute parallel do simd dist_schedule (static, v + 8) safelen(s + 1) &
+  !$omp & schedule (static, t * 2) num_threads (u - 1) num_teams (w + 8)
+  do i = 1, n
+  end do
+  !$omp end target
+  !$omp target
+  !$omp teams distribute parallel do dist_schedule (static, v + 8) num_threads (u - 1) &
+  !$omp & schedule (static, t * 2) num_teams (w + 8)
+  do i = 1, n
+  end do
+  !$omp end target
+  !$omp target teams distribute dist_schedule (static, v + 8) num_teams (w + 8)
+  do i = 1, n
+  end do
+  !$omp target teams distribute simd dist_schedule (static, v + 8) safelen(s + 1) &
+  !$omp & num_teams (w + 8)
+  do i = 1, n
+  end do
+  !$omp target teams distribute parallel do simd dist_schedule (static, v + 8) safelen(s + 1) &
+  !$omp & schedule (static, t * 2) num_threads (u - 1) num_teams (w + 8)
+  do i = 1, n
+  end do
+  !$omp target teams distribute parallel do dist_schedule (static, v + 8) num_threads (u - 1) &
+  !$omp & schedule (static, t * 2) num_teams (w + 8)
+  do i = 1, n
+  end do
+end subroutine
diff --git a/gcc/testsuite/gfortran.dg/gomp/target3.f90 b/gcc/testsuite/gfortran.dg/gomp/target3.f90
new file mode 100644
index 00000000000..53a9682bf96
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target3.f90
@@ -0,0 +1,12 @@
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+subroutine foo (r)
+  integer :: i, r
+  !$omp target
+  !$omp target teams distribute parallel do reduction (+: r) ! { dg-warning "target construct inside of target region" }
+    do i = 1, 10
+      r = r + 1
+    end do
+  !$omp end target
+end subroutine
diff --git a/gcc/testsuite/gfortran.dg/gomp/udr4.f90 b/gcc/testsuite/gfortran.dg/gomp/udr4.f90
index 223dfd04cd2..7e86a757214 100644
--- a/gcc/testsuite/gfortran.dg/gomp/udr4.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/udr4.f90
@@ -6,7 +6,7 @@ subroutine f3
 !$omp declare reduction (foo) ! { dg-error "Unclassifiable OpenMP directive" }
 !$omp declare reduction (foo:integer) ! { dg-error "Unclassifiable OpenMP directive" }
 !$omp declare reduction (foo:integer:omp_out=omp_out+omp_in) &
-!$omp & initializer(omp_priv=0) initializer(omp_priv=0) ! { dg-error "Unclassifiable statement" }
+!$omp & initializer(omp_priv=0) initializer(omp_priv=0) ! { dg-error "Unexpected junk after" }
 end subroutine f3
 subroutine f4
   implicit integer (o)
diff --git a/gcc/testsuite/gfortran.dg/openmp-define-3.f90 b/gcc/testsuite/gfortran.dg/openmp-define-3.f90
index 3d559864faf..44d5c9de49b 100644
--- a/gcc/testsuite/gfortran.dg/openmp-define-3.f90
+++ b/gcc/testsuite/gfortran.dg/openmp-define-3.f90
@@ -6,6 +6,6 @@
 # error _OPENMP not defined
 #endif
 
-#if _OPENMP != 201107
+#if _OPENMP != 201307
 # error _OPENMP defined to wrong value
 #endif
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index aa7498b90e0..a17655389b4 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -1152,6 +1152,11 @@ enum omp_clause_map_kind
      array sections.  OMP_CLAUSE_SIZE for these is not the pointer size,
      which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias.  */
   OMP_CLAUSE_MAP_POINTER,
+  /* Also internal, behaves like OMP_CLAUS_MAP_TO, but additionally any
+     OMP_CLAUSE_MAP_POINTER records consecutive after it which have addresses
+     falling into that range will not be ignored if OMP_CLAUSE_MAP_TO_PSET
+     wasn't mapped already.  */
+  OMP_CLAUSE_MAP_TO_PSET,
   OMP_CLAUSE_MAP_LAST
 };
 
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index 85c6a03f737..ea2fb722795 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1085,6 +1085,10 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_LINEAR:
 	  if (OMP_CLAUSE_LINEAR_GIMPLE_SEQ (clause))
 	    need_stmts = true;
+	  wi->val_only = true;
+	  wi->is_lhs = false;
+	  convert_nonlocal_reference_op (&OMP_CLAUSE_LINEAR_STEP (clause),
+					 &dummy, wi);
 	  goto do_decl_clause;
 
 	case OMP_CLAUSE_PRIVATE:
@@ -1113,10 +1117,42 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_DEPEND:
+	case OMP_CLAUSE_DEVICE:
+	case OMP_CLAUSE_NUM_TEAMS:
+	case OMP_CLAUSE_THREAD_LIMIT:
+	case OMP_CLAUSE_SAFELEN:
 	  wi->val_only = true;
 	  wi->is_lhs = false;
 	  convert_nonlocal_reference_op (&OMP_CLAUSE_OPERAND (clause, 0),
-	                                 &dummy, wi);
+					 &dummy, wi);
+	  break;
+
+	case OMP_CLAUSE_DIST_SCHEDULE:
+	  if (OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (clause) != NULL)
+	    {
+	      wi->val_only = true;
+	      wi->is_lhs = false;
+	      convert_nonlocal_reference_op (&OMP_CLAUSE_OPERAND (clause, 0),
+					     &dummy, wi);
+	    }
+	  break;
+
+	case OMP_CLAUSE_MAP:
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
+	  if (OMP_CLAUSE_SIZE (clause))
+	    {
+	      wi->val_only = true;
+	      wi->is_lhs = false;
+	      convert_nonlocal_reference_op (&OMP_CLAUSE_SIZE (clause),
+					     &dummy, wi);
+	    }
+	  if (DECL_P (OMP_CLAUSE_DECL (clause)))
+	    goto do_decl_clause;
+	  wi->val_only = true;
+	  wi->is_lhs = false;
+	  convert_nonlocal_reference_op (&OMP_CLAUSE_DECL (clause),
+					 &dummy, wi);
 	  break;
 
 	case OMP_CLAUSE_NOWAIT:
@@ -1126,6 +1162,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_COLLAPSE:
 	case OMP_CLAUSE_UNTIED:
 	case OMP_CLAUSE_MERGEABLE:
+	case OMP_CLAUSE_PROC_BIND:
 	  break;
 
 	default:
@@ -1620,6 +1657,10 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_LINEAR:
 	  if (OMP_CLAUSE_LINEAR_GIMPLE_SEQ (clause))
 	    need_stmts = true;
+	  wi->val_only = true;
+	  wi->is_lhs = false;
+	  convert_local_reference_op (&OMP_CLAUSE_LINEAR_STEP (clause), &dummy,
+				      wi);
 	  goto do_decl_clause;
 
 	case OMP_CLAUSE_PRIVATE:
@@ -1653,12 +1694,45 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_DEPEND:
+	case OMP_CLAUSE_DEVICE:
+	case OMP_CLAUSE_NUM_TEAMS:
+	case OMP_CLAUSE_THREAD_LIMIT:
+	case OMP_CLAUSE_SAFELEN:
 	  wi->val_only = true;
 	  wi->is_lhs = false;
 	  convert_local_reference_op (&OMP_CLAUSE_OPERAND (clause, 0), &dummy,
 				      wi);
 	  break;
 
+	case OMP_CLAUSE_DIST_SCHEDULE:
+	  if (OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (clause) != NULL)
+	    {
+	      wi->val_only = true;
+	      wi->is_lhs = false;
+	      convert_local_reference_op (&OMP_CLAUSE_OPERAND (clause, 0),
+					  &dummy, wi);
+	    }
+	  break;
+
+	case OMP_CLAUSE_MAP:
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
+	  if (OMP_CLAUSE_SIZE (clause))
+	    {
+	      wi->val_only = true;
+	      wi->is_lhs = false;
+	      convert_local_reference_op (&OMP_CLAUSE_SIZE (clause),
+					  &dummy, wi);
+	    }
+	  if (DECL_P (OMP_CLAUSE_DECL (clause)))
+	    goto do_decl_clause;
+	  wi->val_only = true;
+	  wi->is_lhs = false;
+	  convert_local_reference_op (&OMP_CLAUSE_DECL (clause),
+				      &dummy, wi);
+	  break;
+
+
 	case OMP_CLAUSE_NOWAIT:
 	case OMP_CLAUSE_ORDERED:
 	case OMP_CLAUSE_DEFAULT:
@@ -1666,6 +1740,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_COLLAPSE:
 	case OMP_CLAUSE_UNTIED:
 	case OMP_CLAUSE_MERGEABLE:
+	case OMP_CLAUSE_PROC_BIND:
 	  break;
 
 	default:
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 3f6152fe8f2..59a825c7c8d 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -500,6 +500,7 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
 	  pp_string (buffer, "alloc");
 	  break;
 	case OMP_CLAUSE_MAP_TO:
+	case OMP_CLAUSE_MAP_TO_PSET:
 	  pp_string (buffer, "to");
 	  break;
 	case OMP_CLAUSE_MAP_FROM:
@@ -520,6 +521,9 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
 	  if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
 	      && OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_POINTER)
 	    pp_string (buffer, " [pointer assign, bias: ");
+	  else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+		   && OMP_CLAUSE_MAP_KIND (clause) == OMP_CLAUSE_MAP_TO_PSET)
+	    pp_string (buffer, " [pointer set, len: ");
 	  else
 	    pp_string (buffer, " [len: ");
 	  dump_generic_node (buffer, OMP_CLAUSE_SIZE (clause),
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 8e6d37ae961..e3fdb625ef6 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,22 @@
+2014-06-18  Jakub Jelinek  <jakub@redhat.com>
+
+	* omp_lib.f90.in (openmp_version): Set to 201307.
+	* omp_lib.h.in (openmp_version): Likewise.
+	* testsuite/libgomp.c/target-8.c: New test.
+	* testsuite/libgomp.fortran/declare-simd-1.f90: Add notinbranch
+	and inbranch clauses.
+	* testsuite/libgomp.fortran/depend-3.f90: New test.
+	* testsuite/libgomp.fortran/openmp_version-1.f: Adjust for new
+	openmp_version.
+	* testsuite/libgomp.fortran/openmp_version-2.f90: Likewise.
+	* testsuite/libgomp.fortran/target1.f90: New test.
+	* testsuite/libgomp.fortran/target2.f90: New test.
+	* testsuite/libgomp.fortran/target3.f90: New test.
+	* testsuite/libgomp.fortran/target4.f90: New test.
+	* testsuite/libgomp.fortran/target5.f90: New test.
+	* testsuite/libgomp.fortran/target6.f90: New test.
+	* testsuite/libgomp.fortran/target7.f90: New test.
+
 2014-06-10  Jakub Jelinek  <jakub@redhat.com>
 
 	PR fortran/60928
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index dda297a1d4e..757053c9fbc 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -42,7 +42,7 @@
       module omp_lib
         use omp_lib_kinds
         implicit none
-        integer, parameter :: openmp_version = 201107
+        integer, parameter :: openmp_version = 201307
 
         interface
           subroutine omp_init_lock (svar)
diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index 7725396ac50..691adb8655f 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -45,7 +45,7 @@
       parameter (omp_proc_bind_master = 2)
       parameter (omp_proc_bind_close = 3)
       parameter (omp_proc_bind_spread = 4)
-      parameter (openmp_version = 201107)
+      parameter (openmp_version = 201307)
 
       external omp_init_lock, omp_init_nest_lock
       external omp_destroy_lock, omp_destroy_nest_lock
diff --git a/libgomp/testsuite/libgomp.c/target-8.c b/libgomp/testsuite/libgomp.c/target-8.c
new file mode 100644
index 00000000000..35084575324
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-8.c
@@ -0,0 +1,26 @@
+/* { dg-do run } */
+/* { dg-options "-fopenmp" } */
+
+void
+foo (int *p)
+{
+  int i;
+  #pragma omp parallel
+  #pragma omp single
+  #pragma omp target teams distribute parallel for map(p[0:24])
+  for (i = 0; i < 24; i++)
+    p[i] = p[i] + 1;
+}
+
+int
+main ()
+{
+  int p[24], i;
+  for (i = 0; i < 24; i++)
+    p[i] = i;
+  foo (p);
+  for (i = 0; i < 24; i++)
+    if (p[i] != i + 1)
+      __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90 b/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90
index b141a3b3eda..5cd592c09db 100644
--- a/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/declare-simd-1.f90
@@ -6,7 +6,8 @@
 module declare_simd_1_mod
   contains
     real function foo (a, b, c)
-      !$omp declare simd (foo) simdlen (4) uniform (a) linear (b : 5)
+      !$omp declare simd (foo) simdlen (4) uniform (a) linear (b : 5) &
+      !$omp & notinbranch
       double precision, value :: a
       real, value :: c
       !$omp declare simd (foo)
@@ -22,6 +23,7 @@ end module declare_simd_1_mod
       real, value :: c
       real :: bar
       !$omp declare simd (bar) simdlen (4) linear (b : 2)
+      !$omp declare simd (bar) simdlen (16) inbranch
       double precision, value :: a
     end function bar
   end interface
diff --git a/libgomp/testsuite/libgomp.fortran/depend-3.f90 b/libgomp/testsuite/libgomp.fortran/depend-3.f90
new file mode 100644
index 00000000000..11be6410692
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/depend-3.f90
@@ -0,0 +1,42 @@
+! { dg-do run }
+
+  integer :: x(2, 3)
+  integer, allocatable :: z(:, :)
+  allocate (z(-2:3, 2:4))
+  call foo (x, z)
+contains
+  subroutine foo (x, z)
+    integer :: x(:, :), y
+    integer, allocatable :: z(:, :)
+    y = 1
+    !$omp parallel shared (x, y, z)
+      !$omp single
+        !$omp taskgroup
+          !$omp task depend(in: x)
+  	  if (y.ne.1) call abort
+          !$omp end task
+          !$omp task depend(out: x(1:2, 1:3))
+  	  y = 2
+          !$omp end task
+        !$omp end taskgroup
+        !$omp taskgroup
+          !$omp task depend(in: z)
+  	  if (y.ne.2) call abort
+          !$omp end task
+          !$omp task depend(out: z(-2:3, 2:4))
+  	  y = 3
+          !$omp end task
+        !$omp end taskgroup
+        !$omp taskgroup
+          !$omp task depend(in: x)
+  	  if (y.ne.3) call abort
+          !$omp end task
+          !$omp task depend(out: x(1:, 1:))
+  	  y = 4
+          !$omp end task
+        !$omp end taskgroup
+      !$omp end single
+    !$omp end parallel
+    if (y.ne.4) call abort
+  end subroutine
+end
diff --git a/libgomp/testsuite/libgomp.fortran/openmp_version-1.f b/libgomp/testsuite/libgomp.fortran/openmp_version-1.f
index aaa888189b1..be24adcca0c 100644
--- a/libgomp/testsuite/libgomp.fortran/openmp_version-1.f
+++ b/libgomp/testsuite/libgomp.fortran/openmp_version-1.f
@@ -4,6 +4,6 @@
       implicit none
       include "omp_lib.h"
 
-      if (openmp_version .ne. 201107) call abort;
+      if (openmp_version .ne. 201307) call abort;
 
       end program main
diff --git a/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90 b/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90
index b2d1d261f27..62712c7d206 100644
--- a/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90
+++ b/libgomp/testsuite/libgomp.fortran/openmp_version-2.f90
@@ -4,6 +4,6 @@ program main
   use omp_lib
   implicit none
 
-  if (openmp_version .ne. 201107) call abort;
+  if (openmp_version .ne. 201307) call abort;
 
 end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target1.f90 b/libgomp/testsuite/libgomp.fortran/target1.f90
new file mode 100644
index 00000000000..c70daace497
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target1.f90
@@ -0,0 +1,58 @@
+! { dg-do run }
+
+module target1
+contains
+  subroutine foo (p, v, w, n)
+    double precision, pointer :: p(:), v(:), w(:)
+    double precision :: q(n)
+    integer :: i, n
+    !$omp target if (n > 256) map (to: v(1:n), w(:n)) map (from: p(1:n), q)
+    !$omp parallel do simd
+      do i = 1, n
+        p(i) = v(i) * w(i)
+        q(i) = p(i)
+      end do
+    !$omp end target
+    if (any (p /= q)) call abort
+    do i = 1, n
+      if (p(i) /= i * iand (i, 63)) call abort
+    end do
+    !$omp target data if (n > 256) map (to: v(1:n), w) map (from: p, q)
+    !$omp target if (n > 256)
+      do i = 1, n
+        p(i) = 1.0
+        q(i) = 2.0
+      end do
+    !$omp end target
+    !$omp target if (n > 256)
+      do i = 1, n
+        p(i) = p(i) + v(i) * w(i)
+        q(i) = q(i) + v(i) * w(i)
+      end do
+    !$omp end target
+    !$omp target if (n > 256)
+      !$omp teams distribute parallel do simd linear(i:1)
+      do i = 1, n
+        p(i) = p(i) + 2.0
+        q(i) = q(i) + 3.0
+      end do
+    !$omp end target
+    !$omp end target data
+    if (any (p + 2.0 /= q)) call abort
+  end subroutine
+end module target1
+  use target1, only : foo
+  integer :: n, i
+  double precision, pointer :: p(:), v(:), w(:)
+  n = 10000
+  allocate (p(n), v(n), w(n))
+  do i = 1, n
+    v(i) = i
+    w(i) = iand (i, 63)
+  end do
+  call foo (p, v, w, n)
+  do i = 1, n
+    if (p(i) /= i * iand (i, 63) + 3) call abort
+  end do
+  deallocate (p, v, w)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target2.f90 b/libgomp/testsuite/libgomp.fortran/target2.f90
new file mode 100644
index 00000000000..42f704f2bb3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target2.f90
@@ -0,0 +1,96 @@
+! { dg-do run }
+! { dg-options "-fopenmp -ffree-line-length-160" }
+
+module target2
+contains
+  subroutine foo (a, b, c, d, e, f, g, n, q)
+    integer :: n, q
+    integer :: a, b(3:n), c(5:), d(2:*), e(:,:)
+    integer, pointer :: f, g(:)
+    integer :: h, i(3:n)
+    integer, pointer :: j, k(:)
+    logical :: r
+    allocate (j, k(4:n))
+    h = 14
+    i = 15
+    j = 16
+    k = 17
+    !$omp target map (to: a, b, c, d(2:n+1), e, f, g, h, i, j, k, n) map (from: r)
+      r = a /= 7
+      r = r .or. (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (f /= 12)
+      r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (h /= 14)
+      r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (j /= 16)
+      r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    if (r) call abort
+    !$omp target map (to: b(3:n), c(5:n+4), d(2:n+1), e(1:,:2), g(3:n), i(3:n), k(4:n), n) map (from: r)
+      r = (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    if (r) call abort
+    !$omp target map (to: b(5:n-2), c(7:n), d(4:n-2), e(1:,2:), g(5:n-3), i(6:n-4), k(5:n-5), n) map (from: r)
+      r = (any (b(5:n-2) /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c(7:n) /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(4:n-2) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e(1:,2:) /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (any (g(5:n-3) /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (any (i(6:n-4) /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (any (k(5:n-5) /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    !$omp target map (to: b(q+5:n-2+q), c(q+7:q+n), d(q+4:q+n-2), e(1:q+2,2:q+2), g(5+q:n-3+q), &
+    !$omp & i(6+q:n-4+q), k(5+q:n-5+q), n) map (from: r)
+      r = (any (b(5:n-2) /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c(7:n) /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(4:n-2) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e(1:,2:) /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (any (g(5:n-3) /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (any (i(6:n-4) /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (any (k(5:n-5) /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    if (r) call abort
+    !$omp target map (to: d(2:n+1), n)
+      r = a /= 7
+      r = r .or. (any (b /= 8)) .or. (lbound (b, 1) /= 3) .or. (ubound (b, 1) /= n)
+      r = r .or. (any (c /= 9)) .or. (lbound (c, 1) /= 5) .or. (ubound (c, 1) /= n + 4)
+      r = r .or. (any (d(2:n+1) /= 10)) .or. (lbound (d, 1) /= 2)
+      r = r .or. (any (e /= 11)) .or. (lbound (e, 1) /= 1) .or. (ubound (e, 1) /= 2)
+      r = r .or. (lbound (e, 2) /= 1) .or. (ubound (e, 2) /= 2)
+      r = r .or. (f /= 12)
+      r = r .or. (any (g /= 13)) .or. (lbound (g, 1) /= 3) .or. (ubound (g, 1) /= n)
+      r = r .or. (h /= 14)
+      r = r .or. (any (i /= 15)) .or. (lbound (i, 1) /= 3) .or. (ubound (i, 1) /= n)
+      r = r .or. (j /= 16)
+      r = r .or. (any (k /= 17)) .or. (lbound (k, 1) /= 4) .or. (ubound (k, 1) /= n)
+    !$omp end target
+    if (r) call abort
+  end subroutine foo
+end module target2
+  use target2, only : foo
+  integer, parameter :: n = 15, q = 0
+  integer :: a, b(2:n-1), c(n), d(n), e(3:4, 3:4)
+  integer, pointer :: f, g(:)
+  allocate (f, g(3:n))
+  a = 7
+  b = 8
+  c = 9
+  d = 10
+  e = 11
+  f = 12
+  g = 13
+  call foo (a, b, c, d, e, f, g, n, q)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target3.f90 b/libgomp/testsuite/libgomp.fortran/target3.f90
new file mode 100644
index 00000000000..1f197acdef7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target3.f90
@@ -0,0 +1,29 @@
+! { dg-do run }
+
+module target3
+contains
+  subroutine foo (f, g)
+    integer :: n
+    integer, pointer :: f, g(:)
+    integer, pointer :: j, k(:)
+    logical :: r
+    nullify (j)
+    k => null ()
+    !$omp target map (tofrom: f, g, j, k) map (from: r)
+      r = associated (f) .or. associated (g)
+      r = r .or. associated (j) .or. associated (k)
+    !$omp end target
+    if (r) call abort
+    !$omp target
+      r = associated (f) .or. associated (g)
+      r = r .or. associated (j) .or. associated (k)
+    !$omp end target
+    if (r) call abort
+  end subroutine foo
+end module target3
+  use target3, only : foo
+  integer, pointer :: f, g(:)
+  f => null ()
+  nullify (g)
+  call foo (f, g)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target4.f90 b/libgomp/testsuite/libgomp.fortran/target4.f90
new file mode 100644
index 00000000000..aa2f0a5ac19
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target4.f90
@@ -0,0 +1,48 @@
+! { dg-do run }
+
+module target4
+contains
+  subroutine foo (a,m,n)
+    integer :: m,n,i,j
+    double precision :: a(m, n), t
+    !$omp target data map(a) map(to: m, n)
+    do i=1,n
+      t = 0.0d0
+      !$omp target
+        !$omp parallel do reduction(+:t)
+          do j=1,m
+            t = t + a(j,i) * a(j,i)
+          end do
+      !$omp end target
+      t = 2.0d0 * t
+      !$omp target
+        !$omp parallel do
+          do j=1,m
+            a(j,i) = a(j,i) * t
+          end do
+      !$omp end target
+    end do
+    !$omp end target data
+  end subroutine foo
+end module target4
+  use target4, only : foo
+  integer :: i, j
+  double precision :: a(8, 9), res(8, 9)
+  do i = 1, 8
+    do j = 1, 9
+      a(i, j) = i + j
+    end do
+  end do
+  call foo (a, 8, 9)
+  res = reshape ((/ 1136.0d0, 1704.0d0, 2272.0d0, 2840.0d0, 3408.0d0, 3976.0d0, &
+&   4544.0d0, 5112.0d0, 2280.0d0, 3040.0d0, 3800.0d0, 4560.0d0, 5320.0d0, 6080.0d0, &
+&   6840.0d0, 7600.0d0, 3936.0d0, 4920.0d0, 5904.0d0, 6888.0d0, 7872.0d0, 8856.0d0, &
+&   9840.0d0, 10824.0d0, 6200.0d0, 7440.0d0, 8680.0d0, 9920.0d0, 11160.0d0, 12400.0d0, &
+&   13640.0d0, 14880.0d0, 9168.0d0, 10696.0d0, 12224.0d0, 13752.0d0, 15280.0d0, 16808.0d0, &
+&   18336.0d0, 19864.0d0, 12936.0d0, 14784.0d0, 16632.0d0, 18480.0d0, 20328.0d0, 22176.0d0, &
+&   24024.0d0, 25872.0d0, 17600.0d0, 19800.0d0, 22000.0d0, 24200.0d0, 26400.0d0, 28600.0d0, &
+&   30800.0d0, 33000.0d0, 23256.0d0, 25840.0d0, 28424.0d0, 31008.0d0, 33592.0d0, 36176.0d0, &
+&   38760.0d0, 41344.0d0, 30000.0d0, 33000.0d0, 36000.0d0, 39000.0d0, 42000.0d0, 45000.0d0, &
+&   48000.0d0, 51000.0d0 /), (/ 8, 9 /))
+  if (any (a /= res)) call abort
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target5.f90 b/libgomp/testsuite/libgomp.fortran/target5.f90
new file mode 100644
index 00000000000..c46faf226f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target5.f90
@@ -0,0 +1,21 @@
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+  integer :: r
+  r = 0
+  call foo (r)
+  if (r /= 11) call abort
+contains
+  subroutine foo (r)
+    integer :: i, r
+    !$omp parallel
+    !$omp single
+    !$omp target teams distribute parallel do reduction (+: r)
+      do i = 1, 10
+        r = r + 1
+      end do
+      r = r + 1
+    !$omp end single
+    !$omp end parallel
+  end subroutine
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target6.f90 b/libgomp/testsuite/libgomp.fortran/target6.f90
new file mode 100644
index 00000000000..13f5a52edd2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target6.f90
@@ -0,0 +1,50 @@
+! { dg-do run }
+
+module target6
+contains
+  subroutine foo (p, v, w, n)
+    double precision, pointer :: p(:), v(:), w(:)
+    double precision :: q(n)
+    integer :: i, n
+    !$omp target data if (n > 256) map (to: v(1:n), w(:n)) map (from: p(1:n), q)
+    !$omp target if (n > 256)
+    !$omp parallel do simd
+      do i = 1, n
+        p(i) = v(i) * w(i)
+        q(i) = p(i)
+      end do
+    !$omp end target
+    !$omp target update if (n > 256) from (p)
+    do i = 1, n
+      if (p(i) /= i * iand (i, 63)) call abort
+      v(i) = v(i) + 1
+    end do
+    !$omp target update if (n > 256) to (v(1:n))
+    !$omp target if (n > 256)
+    !$omp parallel do simd
+      do i = 1, n
+        p(i) = v(i) * w(i)
+      end do
+    !$omp end target
+    !$omp end target data
+    do i = 1, n
+      if (q(i) /= (v(i) - 1) * w(i)) call abort
+      if (p(i) /= q(i) + w(i)) call abort
+    end do
+  end subroutine
+end module target6
+  use target6, only : foo
+  integer :: n, i
+  double precision, pointer :: p(:), v(:), w(:)
+  n = 10000
+  allocate (p(n), v(n), w(n))
+  do i = 1, n
+    v(i) = i
+    w(i) = iand (i, 63)
+  end do
+  call foo (p, v, w, n)
+  do i = 1, n
+    if (p(i) /= (i + 1) * iand (i, 63)) call abort
+  end do
+  deallocate (p, v, w)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target7.f90 b/libgomp/testsuite/libgomp.fortran/target7.f90
new file mode 100644
index 00000000000..4af0ee371bd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target7.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+
+  interface
+    real function foo (x)
+      !$omp declare target
+      real, intent(in) :: x
+    end function foo
+  end interface
+  integer, parameter :: n = 1000
+  integer, parameter :: c = 100
+  integer :: i, j
+  real :: a(n)
+  do i = 1, n
+    a(i) = i
+  end do
+  do i = 1, n, c
+    !$omp task shared(a)
+      !$omp target map(a(i:i+c-1))
+        !$omp parallel do
+          do j = i, i + c - 1
+            a(j) = foo (a(j))
+          end do
+      !$omp end target
+    !$omp end task
+  end do
+  do i = 1, n
+    if (a(i) /= i + 1) call abort
+  end do
+end
+real function foo (x)
+  !$omp declare target
+  real, intent(in) :: x
+  foo = x + 1
+end function foo
author	jakub <jakub@138bc75d-0d04-0410-961f-82ee72b054a4>	2014-06-18 07:16:12 +0000
committer	jakub <jakub@138bc75d-0d04-0410-961f-82ee72b054a4>	2014-06-18 07:16:12 +0000
commit	691447aba6b3f8f3ea8402a35e7c53e2dbb9a6bd (patch)
tree	20e7887bc99ba095e639c296ace7f2c40b3b9f11
parent	23096bb05120b4fdca8164e7fa512ae25a369f30 (diff)
download	gcc-691447aba6b3f8f3ea8402a35e7c53e2dbb9a6bd.tar.gz