* tree.def (OMP_TEAMS, OMP_TARGET_DATA, OMP_TARGET,

	OMP_TARGET_UPDATE): New tree codes.
	* tree-cfg.c (make_edges): Handle GIMPLE_OMP_TARGET
	and GIMPLE_OMP_TEAMS.
	* omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_DIST_SCHEDULE.
	* gimple-low.c (lower_stmt): Handle GIMPLE_OMP_TARGET
	and GIMPLE_OMP_TEAMS.
	* tree.h (OMP_TEAMS_BODY, OMP_TEAMS_CLAUSES, OMP_TARGET_DATA_BODY,
	OMP_TARGET_DATA_CLAUSES, OMP_TARGET_BODY, OMP_TARGET_CLAUSES,
	OMP_TARGET_UPDATE_CLAUSES): Define.
 	* tree-nested.c (convert_nonlocal_reference_stmt,
	convert_local_reference_stmt, convert_gimple_call): Handle
	GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	* tree-inline.c (estimate_num_insns): Likewise.
	(remap_gimple_stmt): Likewise.  Adjust gimple_build_omp_for
	caller.
	* gimple.def: Adjust comments describing OMP_CLAUSEs.
	(GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS): New GIMPLE stmts.
	* tree-parloops.c (create_parallel_loop): Adjust gimple_build_omp_for
	caller.
	* tree-pretty-print.c (dump_generic_node): Handle OMP_TEAMS,
	OMP_TARGET, OMP_TARGET_DATA and OMP_TARGET_UPDATE.
	* gimple.h (GF_OMP_TARGET_KIND_MASK, GF_OMP_TARGET_KIND_REGION,
	GF_OMP_TARGET_KIND_DATA, GF_OMP_TARGET_KIND_UPDATE): New.
	(gimple_build_omp_for): Add kind argument to prototype.
	(gimple_build_omp_target, gimple_build_omp_teams): New prototypes.
	(gimple_has_substatements): Handle GIMPLE_OMP_TARGET and
	GIMPLE_OMP_TEAMS.
	(gimple_omp_subcode): Change GIMPLE_OMP_SINGLE to GIMPLE_OMP_TEAMS.
	(gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
	gimple_omp_target_set_clauses, gimple_omp_target_kind,
	gimple_omp_target_set_kind, gimple_omp_teams_clauses,
	gimple_omp_teams_clauses_ptr, gimple_omp_teams_set_clauses): New
	inline functions.
	(gimple_return_set_retval): Handle GIMPLE_OMP_TARGET and
	GIMPLE_OMP_TEAMS.
	* gimple.c (gimple_build_omp_for): Add kind argument, call
	gimple_omp_for_set_kind.
	(gimple_build_omp_target, gimple_build_omp_teams): New functions.
	(walk_gimple_op, walk_gimple_stmt, gimple_copy): Handle
	GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	* gimple-pretty-print.c (dump_gimple_omp_target,
	dump_gimple_omp_teams): New functions.
	(pp_gimple_stmt_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP.
	(enum omp_region_type): Add ORT_TEAMS, ORT_TARGET and ORT_TARGET_DATA.
	(omp_add_variable): Add temporary assertions.
	(omp_notice_threadprivate_variable): Complain if threadprivate vars
	appear in target region.
	(omp_notice_variable): ORT_TARGET, ORT_TARGET_DATA and ORT_TEAMS
	handling.
	(omp_check_private): Ignore ORT_TARGET and ORT_TARGET_DATA regions.
	(gimplify_scan_omp_clauses): Handle OMP_CLAUSE_MAP, OMP_CLAUSE_TO,
	OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_DIST_SCHEDULE
	and OMP_CLAUSE_DEVICE.
	(gimplify_adjust_omp_clauses): Likewise.
	(gimplify_adjust_omp_clauses_1): Handle GOVD_MAP.  Fix up
	check for privatization by also testing for GOVD_LINEAR.
	(gimplify_omp_for): Adjust gimple_build_omp_for caller.
	Clear *expr_p.
	(gimplify_omp_workshare): Handle also OMP_TARGET, OMP_TARGET_DATA
	and OMP_TEAMS.  Clear *expr_p.
	(gimplify_omp_target_update): New function.
	(gimplify_expr): Handle OMP_TARGET, OMP_TARGET_DATA, OMP_TARGET_UPDATE
	and OMP_TEAMS.
cp/
	* parser.c (cp_parser_omp_clause_cancelkind): Remove diagnostics.
	(cp_parser_omp_all_clauses): Require that OMP_CLAUSE_{TO,FROM}
	and OMP_CLAUSE_{PARALLEL,FOR,SECTIONS,TASKGROUP} must be first in
	the list of clauses.
	(OMP_TEAMS_CLAUSE_MASK, OMP_TARGET_CLAUSE_MASK,
	OMP_TARGET_DATA_CLAUSE_MASK, OMP_TARGET_UPDATE_CLAUSE_MASK,
	OMP_DISTRIBUTE_CLAUSE_MASK): Define.
	(cp_parser_omp_teams, cp_parser_omp_target, cp_parser_omp_target_data,
	cp_parser_omp_target_update, cp_parser_omp_distribute): New functions.
	(cp_parser_omp_construct): Handle PRAGMA_OMP_DISTRIBUTE and
	PRAGMA_OMP_TEAMS.
	(cp_parser_pragma): Handle PRAGMA_OMP_DISTRIBUTE, PRAGMA_OMP_TEAMS
	and PRAGMA_OMP_TARGET.
	* pt.c (tsubst_expr): Handle OMP_TEAMS, OMP_TARGET, OMP_TARGET_DATA
	and OMP_TARGET_UPDATE.


git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@199349 138bc75d-0d04-0410-961f-82ee72b054a4

18 files changed, 894 insertions, 63 deletions

diff --git a/gcc/ChangeLog.gomp b/gcc/ChangeLog.gomp
index d9b120f7ddf..0db46c9aef5 100644
--- a/gcc/ChangeLog.gomp
+++ b/gcc/ChangeLog.gomp
@@ -1,3 +1,71 @@
+2013-05-27  Jakub Jelinek  <jakub@redhat.com>
+
+	* tree.def (OMP_TEAMS, OMP_TARGET_DATA, OMP_TARGET,
+	OMP_TARGET_UPDATE): New tree codes.
+	* tree-cfg.c (make_edges): Handle GIMPLE_OMP_TARGET
+	and GIMPLE_OMP_TEAMS.
+	* omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_DIST_SCHEDULE.
+	* gimple-low.c (lower_stmt): Handle GIMPLE_OMP_TARGET
+	and GIMPLE_OMP_TEAMS.
+	* tree.h (OMP_TEAMS_BODY, OMP_TEAMS_CLAUSES, OMP_TARGET_DATA_BODY,
+	OMP_TARGET_DATA_CLAUSES, OMP_TARGET_BODY, OMP_TARGET_CLAUSES,
+	OMP_TARGET_UPDATE_CLAUSES): Define.
+ 	* tree-nested.c (convert_nonlocal_reference_stmt,
+	convert_local_reference_stmt, convert_gimple_call): Handle
+	GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
+	* tree-inline.c (estimate_num_insns): Likewise.
+	(remap_gimple_stmt): Likewise.  Adjust gimple_build_omp_for
+	caller.
+	* gimple.def: Adjust comments describing OMP_CLAUSEs.
+	(GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS): New GIMPLE stmts.
+	* tree-parloops.c (create_parallel_loop): Adjust gimple_build_omp_for
+	caller.
+	* tree-pretty-print.c (dump_generic_node): Handle OMP_TEAMS,
+	OMP_TARGET, OMP_TARGET_DATA and OMP_TARGET_UPDATE.
+	* gimple.h (GF_OMP_TARGET_KIND_MASK, GF_OMP_TARGET_KIND_REGION,
+	GF_OMP_TARGET_KIND_DATA, GF_OMP_TARGET_KIND_UPDATE): New.
+	(gimple_build_omp_for): Add kind argument to prototype.
+	(gimple_build_omp_target, gimple_build_omp_teams): New prototypes.
+	(gimple_has_substatements): Handle GIMPLE_OMP_TARGET and
+	GIMPLE_OMP_TEAMS.
+	(gimple_omp_subcode): Change GIMPLE_OMP_SINGLE to GIMPLE_OMP_TEAMS.
+	(gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
+	gimple_omp_target_set_clauses, gimple_omp_target_kind,
+	gimple_omp_target_set_kind, gimple_omp_teams_clauses,
+	gimple_omp_teams_clauses_ptr, gimple_omp_teams_set_clauses): New
+	inline functions.
+	(gimple_return_set_retval): Handle GIMPLE_OMP_TARGET and
+	GIMPLE_OMP_TEAMS.
+	* gimple.c (gimple_build_omp_for): Add kind argument, call
+	gimple_omp_for_set_kind.
+	(gimple_build_omp_target, gimple_build_omp_teams): New functions.
+	(walk_gimple_op, walk_gimple_stmt, gimple_copy): Handle
+	GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
+	* gimple-pretty-print.c (dump_gimple_omp_target,
+	dump_gimple_omp_teams): New functions.
+	(pp_gimple_stmt_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
+	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP.
+	(enum omp_region_type): Add ORT_TEAMS, ORT_TARGET and ORT_TARGET_DATA.
+	(omp_add_variable): Add temporary assertions.
+	(omp_notice_threadprivate_variable): Complain if threadprivate vars
+	appear in target region.
+	(omp_notice_variable): ORT_TARGET, ORT_TARGET_DATA and ORT_TEAMS
+	handling.
+	(omp_check_private): Ignore ORT_TARGET and ORT_TARGET_DATA regions.
+	(gimplify_scan_omp_clauses): Handle OMP_CLAUSE_MAP, OMP_CLAUSE_TO,
+	OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_DIST_SCHEDULE
+	and OMP_CLAUSE_DEVICE.
+	(gimplify_adjust_omp_clauses): Likewise.
+	(gimplify_adjust_omp_clauses_1): Handle GOVD_MAP.  Fix up
+	check for privatization by also testing for GOVD_LINEAR.
+	(gimplify_omp_for): Adjust gimple_build_omp_for caller.
+	Clear *expr_p.
+	(gimplify_omp_workshare): Handle also OMP_TARGET, OMP_TARGET_DATA
+	and OMP_TEAMS.  Clear *expr_p.
+	(gimplify_omp_target_update): New function.
+	(gimplify_expr): Handle OMP_TARGET, OMP_TARGET_DATA, OMP_TARGET_UPDATE
+	and OMP_TEAMS.
+
 2013-05-20  Jakub Jelinek  <jakub@redhat.com>
 
 	* omp-low.c (expand_omp_simd): For collapse > 1 loops,
diff --git a/gcc/cp/ChangeLog.gomp b/gcc/cp/ChangeLog.gomp
index 69be3ded378..9f3f38c18c6 100644
--- a/gcc/cp/ChangeLog.gomp
+++ b/gcc/cp/ChangeLog.gomp
@@ -1,3 +1,21 @@
+2013-05-27  Jakub Jelinek  <jakub@redhat.com>
+
+	* parser.c (cp_parser_omp_clause_cancelkind): Remove diagnostics.
+	(cp_parser_omp_all_clauses): Require that OMP_CLAUSE_{TO,FROM}
+	and OMP_CLAUSE_{PARALLEL,FOR,SECTIONS,TASKGROUP} must be first in
+	the list of clauses.
+	(OMP_TEAMS_CLAUSE_MASK, OMP_TARGET_CLAUSE_MASK,
+	OMP_TARGET_DATA_CLAUSE_MASK, OMP_TARGET_UPDATE_CLAUSE_MASK,
+	OMP_DISTRIBUTE_CLAUSE_MASK): Define.
+	(cp_parser_omp_teams, cp_parser_omp_target, cp_parser_omp_target_data,
+	cp_parser_omp_target_update, cp_parser_omp_distribute): New functions.
+	(cp_parser_omp_construct): Handle PRAGMA_OMP_DISTRIBUTE and
+	PRAGMA_OMP_TEAMS.
+	(cp_parser_pragma): Handle PRAGMA_OMP_DISTRIBUTE, PRAGMA_OMP_TEAMS
+	and PRAGMA_OMP_TARGET.
+	* pt.c (tsubst_expr): Handle OMP_TEAMS, OMP_TARGET, OMP_TARGET_DATA
+	and OMP_TARGET_UPDATE.
+
 2013-05-09  Jakub Jelinek  <jakub@redhat.com>
 
 	* cp-tree.h (cp_decl_specifier_seq): Add omp_declare_simd_clauses
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index d7b041c0b5c..5bbb46b1c87 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -26823,22 +26823,7 @@ cp_parser_omp_clause_cancelkind (cp_parser * /*parser*/,
 				 enum omp_clause_code code,
 				 tree list, location_t location)
 {
-  tree c;
-
-  for (c = list; c; c = OMP_CLAUSE_CHAIN (c))
-    switch (OMP_CLAUSE_CODE (c))
-      {
-      case OMP_CLAUSE_PARALLEL:
-      case OMP_CLAUSE_FOR:
-      case OMP_CLAUSE_SECTIONS:
-      case OMP_CLAUSE_TASKGROUP:
-	error_at (location, "only one of %<parallel%>, %<for%>, %<sections%> "
-			    "and %<taskgroup%> clauses can be specified");
-	break;
-      default:
-	break;
-      }
-  c = build_omp_clause (location, code);
+  tree c = build_omp_clause (location, code);
   OMP_CLAUSE_CHAIN (c) = list;
   return c;
 }
@@ -27260,7 +27245,6 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 
       token = cp_lexer_peek_token (parser->lexer);
       c_kind = cp_parser_omp_clause_name (parser);
-      first = false;
 
       switch (c_kind)
 	{
@@ -27359,31 +27343,48 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_PARALLEL,
 						     clauses, token->location);
 	  c_name = "parallel";
+	  if (!first)
+	    {
+	     clause_not_first:
+	      error_at (token->location, "%qs must be the first clause of %qs",
+			c_name, where);
+	      clauses = prev;
+	    }
 	  break;
 	case PRAGMA_OMP_CLAUSE_FOR:
 	  clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_FOR,
 						     clauses, token->location);
 	  c_name = "for";
+	  if (!first)
+	    goto clause_not_first;
 	  break;
 	case PRAGMA_OMP_CLAUSE_SECTIONS:
 	  clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_SECTIONS,
 						     clauses, token->location);
 	  c_name = "sections";
+	  if (!first)
+	    goto clause_not_first;
 	  break;
 	case PRAGMA_OMP_CLAUSE_TASKGROUP:
 	  clauses = cp_parser_omp_clause_cancelkind (parser, OMP_CLAUSE_TASKGROUP,
 						     clauses, token->location);
 	  c_name = "taskgroup";
+	  if (!first)
+	    goto clause_not_first;
 	  break;
 	case PRAGMA_OMP_CLAUSE_TO:
 	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO,
 					    clauses);
 	  c_name = "to";
+	  if (!first)
+	    goto clause_not_first;
 	  break;
 	case PRAGMA_OMP_CLAUSE_FROM:
 	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM,
 					    clauses);
 	  c_name = "from";
+	  if (!first)
+	    goto clause_not_first;
 	  break;
 	case PRAGMA_OMP_CLAUSE_UNIFORM:
 	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_UNIFORM,
@@ -27441,6 +27442,8 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  goto saw_error;
 	}
 
+      first = false;
+
       if (((mask >> c_kind) & 1) == 0)
 	{
 	  /* Remove the invalid clause(s) from the list to avoid
@@ -29013,6 +29016,180 @@ cp_parser_omp_cancellation_point (cp_parser *parser, cp_token *pragma_tok)
 }
 
 /* OpenMP 4.0:
+   # pragma omp teams teams-clause[optseq] new-line
+     structured-block  */
+
+#define OMP_TEAMS_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SHARED)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_TEAMS)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULT))
+
+static tree
+cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok)
+{
+  tree stmt = make_node (OMP_TEAMS);
+  TREE_TYPE (stmt) = void_type_node;
+
+  OMP_TEAMS_CLAUSES (stmt)
+    = cp_parser_omp_all_clauses (parser, OMP_TEAMS_CLAUSE_MASK,
+				 "#pragma omp teams", pragma_tok);
+  OMP_TEAMS_BODY (stmt) = cp_parser_omp_structured_block (parser);
+
+  return add_stmt (stmt);
+}
+
+/* OpenMP 4.0:
+   # pragma omp target data target-data-clause[optseq] new-line
+     structured-block  */
+
+#define OMP_TARGET_DATA_CLAUSE_MASK				\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF))
+
+static tree
+cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok)
+{
+  tree stmt = make_node (OMP_TARGET_DATA);
+  TREE_TYPE (stmt) = void_type_node;
+
+  OMP_TARGET_DATA_CLAUSES (stmt)
+    = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
+				 "#pragma omp target data", pragma_tok);
+  OMP_TARGET_DATA_BODY (stmt) = cp_parser_omp_structured_block (parser);
+
+  SET_EXPR_LOCATION (stmt, pragma_tok->location);
+  return add_stmt (stmt);
+}
+
+/* OpenMP 4.0:
+   # pragma omp target update target-update-clause[optseq] new-line */
+
+#define OMP_TARGET_UPDATE_CLAUSE_MASK				\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FROM)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF))
+
+static bool
+cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
+			     enum pragma_context context)
+{
+  if (context == pragma_stmt)
+    {
+      error_at (pragma_tok->location,
+		"%<#pragma omp target update%> may only be "
+		"used in compound statements");
+      cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+      return false;
+    }
+
+  tree clauses
+    = cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK,
+				 "#pragma omp target update", pragma_tok);
+  if (find_omp_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE
+      && find_omp_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE)
+    {
+      error_at (pragma_tok->location,
+		"%<#pragma omp target update must contain either "
+		"%<from%> or %<to%> clauses");
+      return false;
+    }
+
+  tree stmt = make_node (OMP_TARGET_UPDATE);
+  TREE_TYPE (stmt) = void_type_node;
+  OMP_TARGET_UPDATE_CLAUSES (stmt) = clauses;
+  SET_EXPR_LOCATION (stmt, pragma_tok->location);
+  add_stmt (stmt);
+  return false;
+}
+
+/* OpenMP 4.0:
+   # pragma omp target target-clause[optseq] new-line
+     structured-block  */
+
+#define OMP_TARGET_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF))
+
+static bool
+cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
+		      enum pragma_context context)
+{
+  if (context != pragma_stmt && context != pragma_compound)
+    {
+      cp_parser_error (parser, "expected declaration specifiers");
+      cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+      return false;
+    }
+
+  if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+    {
+      tree id = cp_lexer_peek_token (parser->lexer)->u.value;
+      const char *p = IDENTIFIER_POINTER (id);
+
+      if (strcmp (p, "data") == 0)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  cp_parser_omp_target_data (parser, pragma_tok);
+	  return true;
+	}
+      else if (strcmp (p, "update") == 0)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  return cp_parser_omp_target_update (parser, pragma_tok, context);
+	}
+    }
+
+  tree stmt = make_node (OMP_TARGET);
+  TREE_TYPE (stmt) = void_type_node;
+
+  OMP_TARGET_CLAUSES (stmt)
+    = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+				 "#pragma omp target", pragma_tok);
+  OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser);
+
+  SET_EXPR_LOCATION (stmt, pragma_tok->location);
+  add_stmt (stmt);
+  return true;
+}
+
+/* OpenMP 4.0:
+   #pragma omp distribute distribute-clause[optseq] new-line
+     for-loop  */
+
+#define OMP_DISTRIBUTE_CLAUSE_MASK				\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE))
+
+static tree
+cp_parser_omp_distribute (cp_parser *parser, cp_token *pragma_tok)
+{
+  tree clauses, sb, ret;
+  unsigned int save;
+
+  clauses = cp_parser_omp_all_clauses (parser, OMP_DISTRIBUTE_CLAUSE_MASK,
+				       "#pragma omp distribute", pragma_tok);
+
+  sb = begin_omp_structured_block ();
+  save = cp_parser_begin_omp_structured_block (parser);
+
+  ret = cp_parser_omp_for_loop (parser, OMP_DISTRIBUTE, clauses, NULL);
+
+  cp_parser_end_omp_structured_block (parser, save);
+  add_stmt (finish_omp_structured_block (sb));
+
+  return ret;
+}
+
+/* OpenMP 4.0:
    # pragma omp declare simd declare-simd-clauses[optseq] new-line  */
 
 #define OMP_DECLARE_SIMD_CLAUSE_MASK				\
@@ -29112,6 +29289,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
     case PRAGMA_OMP_CRITICAL:
       stmt = cp_parser_omp_critical (parser, pragma_tok);
       break;
+    case PRAGMA_OMP_DISTRIBUTE:
+      stmt = cp_parser_omp_distribute (parser, pragma_tok);
+      break;
     case PRAGMA_OMP_FOR:
       stmt = cp_parser_omp_for (parser, pragma_tok);
       break;
@@ -29139,6 +29319,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
     case PRAGMA_OMP_TASKGROUP:
       cp_parser_omp_taskgroup (parser, pragma_tok);
       return;
+    case PRAGMA_OMP_TEAMS:
+      stmt = cp_parser_omp_teams (parser, pragma_tok);
+      break;
     default:
       gcc_unreachable ();
     }
@@ -29609,6 +29792,7 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
 
     case PRAGMA_OMP_ATOMIC:
     case PRAGMA_OMP_CRITICAL:
+    case PRAGMA_OMP_DISTRIBUTE:
     case PRAGMA_OMP_FOR:
     case PRAGMA_OMP_MASTER:
     case PRAGMA_OMP_ORDERED:
@@ -29618,11 +29802,15 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
     case PRAGMA_OMP_SINGLE:
     case PRAGMA_OMP_TASK:
     case PRAGMA_OMP_TASKGROUP:
+    case PRAGMA_OMP_TEAMS:
       if (context != pragma_stmt && context != pragma_compound)
 	goto bad_stmt;
       cp_parser_omp_construct (parser, pragma_tok);
       return true;
 
+    case PRAGMA_OMP_TARGET:
+      return cp_parser_omp_target (parser, pragma_tok, context);
+
     case PRAGMA_OMP_SECTION:
       error_at (pragma_tok->location, 
 		"%<#pragma omp section%> may only be used in "
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 284f6ab6539..9583afd51ce 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -13330,6 +13330,9 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
 
     case OMP_SECTIONS:
     case OMP_SINGLE:
+    case OMP_TEAMS:
+    case OMP_TARGET_DATA:
+    case OMP_TARGET:
       tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false,
 				args, complain, in_decl);
       stmt = push_stmt_list ();
@@ -13342,6 +13345,14 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       add_stmt (t);
       break;
 
+    case OMP_TARGET_UPDATE:
+      tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false,
+				args, complain, in_decl);
+      t = copy_node (t);
+      OMP_CLAUSES (t) = tmp;
+      add_stmt (t);
+      break;
+
     case OMP_SECTION:
     case OMP_CRITICAL:
     case OMP_MASTER:
diff --git a/gcc/gimple-low.c b/gcc/gimple-low.c
index b06d194da65..cf61ef0d646 100644
--- a/gcc/gimple-low.c
+++ b/gcc/gimple-low.c
@@ -444,6 +444,8 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data)
 
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
+    case GIMPLE_OMP_TARGET:
+    case GIMPLE_OMP_TEAMS:
       data->cannot_fallthru = false;
       lower_omp_directive (gsi, data);
       data->cannot_fallthru = false;
diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index 14872166946..ee264e10a88 100644
--- a/gcc/gimple-pretty-print.c
+++ b/gcc/gimple-pretty-print.c
@@ -1264,6 +1264,78 @@ dump_gimple_omp_single (pretty_printer *buffer, gimple gs, int spc, int flags)
     }
 }
 
+/* Dump a GIMPLE_OMP_TARGET tuple on the pretty_printer BUFFER.  */
+
+static void
+dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags)
+{
+  const char *kind;
+  switch (gimple_omp_target_kind (gs))
+    {
+    case GF_OMP_TARGET_KIND_REGION:
+      kind = "";
+      break;
+    case GF_OMP_TARGET_KIND_DATA:
+      kind = " data";
+      break;
+    case GF_OMP_TARGET_KIND_UPDATE:
+      kind = " update";
+      break;
+    default:
+      gcc_unreachable ();
+    }
+  if (flags & TDF_RAW)
+    {
+      dump_gimple_fmt (buffer, spc, flags, "%G%s <%+BODY <%S>%nCLAUSES <", gs,
+		       kind, gimple_omp_body (gs));
+      dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags);
+      dump_gimple_fmt (buffer, spc, flags, " >");
+    }
+  else
+    {
+      pp_string (buffer, "#pragma omp target");
+      pp_string (buffer, kind);
+      dump_omp_clauses (buffer, gimple_omp_target_clauses (gs), spc, flags);
+      if (!gimple_seq_empty_p (gimple_omp_body (gs)))
+	{
+	  newline_and_indent (buffer, spc + 2);
+	  pp_character (buffer, '{');
+	  pp_newline (buffer);
+	  dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags);
+	  newline_and_indent (buffer, spc + 2);
+	  pp_character (buffer, '}');
+	}
+    }
+}
+
+/* Dump a GIMPLE_OMP_TEAMS tuple on the pretty_printer BUFFER.  */
+
+static void
+dump_gimple_omp_teams (pretty_printer *buffer, gimple gs, int spc, int flags)
+{
+  if (flags & TDF_RAW)
+    {
+      dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S>%nCLAUSES <", gs,
+		       gimple_omp_body (gs));
+      dump_omp_clauses (buffer, gimple_omp_teams_clauses (gs), spc, flags);
+      dump_gimple_fmt (buffer, spc, flags, " >");
+    }
+  else
+    {
+      pp_string (buffer, "#pragma omp teams");
+      dump_omp_clauses (buffer, gimple_omp_teams_clauses (gs), spc, flags);
+      if (!gimple_seq_empty_p (gimple_omp_body (gs)))
+	{
+	  newline_and_indent (buffer, spc + 2);
+	  pp_character (buffer, '{');
+	  pp_newline (buffer);
+	  dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags);
+	  newline_and_indent (buffer, spc + 2);
+	  pp_character (buffer, '}');
+	}
+    }
+}
+
 /* Dump a GIMPLE_OMP_SECTIONS tuple on the pretty_printer BUFFER.  */
 
 static void
@@ -2038,6 +2110,14 @@ pp_gimple_stmt_1 (pretty_printer *buffer, gimple gs, int spc, int flags)
       dump_gimple_omp_single (buffer, gs, spc, flags);
       break;
 
+    case GIMPLE_OMP_TARGET:
+      dump_gimple_omp_target (buffer, gs, spc, flags);
+      break;
+
+    case GIMPLE_OMP_TEAMS:
+      dump_gimple_omp_teams (buffer, gs, spc, flags);
+      break;
+
     case GIMPLE_OMP_RETURN:
       dump_gimple_omp_return (buffer, gs, spc, flags);
       break;
diff --git a/gcc/gimple.c b/gcc/gimple.c
index f5074199381..7d137fdfc5a 100644
--- a/gcc/gimple.c
+++ b/gcc/gimple.c
@@ -908,13 +908,14 @@ gimple_build_omp_critical (gimple_seq body, tree name)
    PRE_BODY is the sequence of statements that are loop invariant.  */
 
 gimple
-gimple_build_omp_for (gimple_seq body, tree clauses, size_t collapse,
+gimple_build_omp_for (gimple_seq body, int kind, tree clauses, size_t collapse,
 		      gimple_seq pre_body)
 {
   gimple p = gimple_alloc (GIMPLE_OMP_FOR, 0);
   if (body)
     gimple_omp_set_body (p, body);
   gimple_omp_for_set_clauses (p, clauses);
+  gimple_omp_for_set_kind (p, kind);
   p->gimple_omp_for.collapse = collapse;
   p->gimple_omp_for.iter
       = ggc_alloc_cleared_vec_gimple_omp_for_iter (collapse);
@@ -1094,6 +1095,41 @@ gimple_build_omp_single (gimple_seq body, tree clauses)
 }
 
 
+/* Build a GIMPLE_OMP_TARGET statement.
+
+   BODY is the sequence of statements that will be executed.
+   CLAUSES are any of the OMP target construct's clauses.  */
+
+gimple
+gimple_build_omp_target (gimple_seq body, int kind, tree clauses)
+{
+  gimple p = gimple_alloc (GIMPLE_OMP_TARGET, 0);
+  if (body)
+    gimple_omp_set_body (p, body);
+  gimple_omp_target_set_clauses (p, clauses);
+  gimple_omp_target_set_kind (p, kind);
+
+  return p;
+}
+
+
+/* Build a GIMPLE_OMP_TEAMS statement.
+
+   BODY is the sequence of statements that will be executed.
+   CLAUSES are any of the OMP teams construct's clauses.  */
+
+gimple
+gimple_build_omp_teams (gimple_seq body, tree clauses)
+{
+  gimple p = gimple_alloc (GIMPLE_OMP_TEAMS, 0);
+  if (body)
+    gimple_omp_set_body (p, body);
+  gimple_omp_teams_set_clauses (p, clauses);
+
+  return p;
+}
+
+
 /* Build a GIMPLE_OMP_ATOMIC_LOAD statement.  */
 
 gimple
@@ -1610,6 +1646,20 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op,
 	return ret;
       break;
 
+    case GIMPLE_OMP_TARGET:
+      ret = walk_tree (gimple_omp_target_clauses_ptr (stmt), callback_op, wi,
+		       pset);
+      if (ret)
+	return ret;
+      break;
+
+    case GIMPLE_OMP_TEAMS:
+      ret = walk_tree (gimple_omp_teams_clauses_ptr (stmt), callback_op, wi,
+		       pset);
+      if (ret)
+	return ret;
+      break;
+
     case GIMPLE_OMP_ATOMIC_LOAD:
       ret = walk_tree (gimple_omp_atomic_load_lhs_ptr (stmt), callback_op, wi,
 		       pset);
@@ -1786,6 +1836,8 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt,
     case GIMPLE_OMP_TASK:
     case GIMPLE_OMP_SECTIONS:
     case GIMPLE_OMP_SINGLE:
+    case GIMPLE_OMP_TARGET:
+    case GIMPLE_OMP_TEAMS:
       ret = walk_gimple_seq_mod (gimple_omp_body_ptr (stmt), callback_stmt,
 			     callback_op, wi);
       if (ret)
@@ -2308,6 +2360,8 @@ gimple_copy (gimple stmt)
 	  /* FALLTHRU  */
 
 	case GIMPLE_OMP_SINGLE:
+	case GIMPLE_OMP_TARGET:
+	case GIMPLE_OMP_TEAMS:
 	case GIMPLE_OMP_SECTION:
 	case GIMPLE_OMP_MASTER:
 	case GIMPLE_OMP_ORDERED:
diff --git a/gcc/gimple.def b/gcc/gimple.def
index acad572e353..8d75ada3ec3 100644
--- a/gcc/gimple.def
+++ b/gcc/gimple.def
@@ -287,7 +287,7 @@ DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_omp_ordered", GSS_OMP)
 
    BODY is a the sequence of statements to be executed by all threads.
 
-   CLAUSES is a TREE_LIST node with all the clauses.
+   CLAUSES is an OMP_CLAUSE chain with all the clauses.
 
    CHILD_FN is set when outlining the body of the parallel region.
    All the statements in BODY are moved into this newly created
@@ -306,7 +306,7 @@ DEFGSCODE(GIMPLE_OMP_PARALLEL, "gimple_omp_parallel", GSS_OMP_PARALLEL)
 
    BODY is a the sequence of statements to be executed by all threads.
 
-   CLAUSES is a TREE_LIST node with all the clauses.
+   CLAUSES is an OMP_CLAUSE chain with all the clauses.
 
    CHILD_FN is set when outlining the body of the explicit task region.
    All the statements in BODY are moved into this newly created
@@ -334,7 +334,7 @@ DEFGSCODE(GIMPLE_OMP_SECTION, "gimple_omp_section", GSS_OMP)
 /* OMP_SECTIONS <BODY, CLAUSES, CONTROL> represents #pragma omp sections.
 
    BODY is the sequence of statements in the sections body.
-   CLAUSES is a TREE_LIST node holding the list of associated clauses.
+   CLAUSES is an OMP_CLAUSE chain holding the list of associated clauses.
    CONTROL is a VAR_DECL used for deciding which of the sections
    to execute.  */
 DEFGSCODE(GIMPLE_OMP_SECTIONS, "gimple_omp_sections", GSS_OMP_SECTIONS)
@@ -346,9 +346,21 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE)
 
 /* GIMPLE_OMP_SINGLE <BODY, CLAUSES> represents #pragma omp single
    BODY is the sequence of statements inside the single section.
-   CLAUSES is a TREE_LIST node holding the associated clauses.  */
+   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
 DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE)
 
+/* GIMPLE_OMP_TARGET <BODY, CLAUSES> represents
+   #pragma omp target {,data,update}
+   BODY is the sequence of statements inside the target construct
+   (NULL for target update).
+   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
+DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_SINGLE)
+
+/* GIMPLE_OMP_TEAMS <BODY, CLAUSES> represents #pragma omp teams
+   BODY is the sequence of statements inside the single section.
+   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
+DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_SINGLE)
+
 /* GIMPLE_PREDICT <PREDICT, OUTCOME> specifies a hint for branch prediction.
 
    PREDICT is one of the predictors from predict.def.
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 93e2c2732f8..5060b282afd 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -115,6 +115,10 @@ enum gf_mask {
     GF_OMP_FOR_KIND_SIMD	= 1 << 0,
     GF_OMP_FOR_KIND_FOR_SIMD	= 2 << 0,
     GF_OMP_FOR_KIND_DISTRIBUTE	= 3 << 0,
+    GF_OMP_TARGET_KIND_MASK	= 3 << 0,
+    GF_OMP_TARGET_KIND_REGION	= 0 << 0,
+    GF_OMP_TARGET_KIND_DATA	= 1 << 0,
+    GF_OMP_TARGET_KIND_UPDATE	= 2 << 0,
 
     /* True on an GIMPLE_OMP_RETURN statement if the return does not require
        a thread synchronization via some sort of barrier.  The exact barrier
@@ -618,7 +622,7 @@ struct GTY(()) gimple_statement_omp_continue {
   tree control_use;
 };
 
-/* GIMPLE_OMP_SINGLE */
+/* GIMPLE_OMP_SINGLE, GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS */
 
 struct GTY(()) gimple_statement_omp_single {
   /* [ WORD 1-7 ]  */
@@ -805,7 +809,7 @@ gimple gimple_build_switch_nlabels (unsigned, tree, tree);
 gimple gimple_build_switch (tree, tree, vec<tree> );
 gimple gimple_build_omp_parallel (gimple_seq, tree, tree, tree);
 gimple gimple_build_omp_task (gimple_seq, tree, tree, tree, tree, tree, tree);
-gimple gimple_build_omp_for (gimple_seq, tree, size_t, gimple_seq);
+gimple gimple_build_omp_for (gimple_seq, int, tree, size_t, gimple_seq);
 gimple gimple_build_omp_critical (gimple_seq, tree);
 gimple gimple_build_omp_section (gimple_seq);
 gimple gimple_build_omp_continue (tree, tree);
@@ -815,6 +819,8 @@ gimple gimple_build_omp_ordered (gimple_seq);
 gimple gimple_build_omp_sections (gimple_seq, tree);
 gimple gimple_build_omp_sections_switch (void);
 gimple gimple_build_omp_single (gimple_seq, tree);
+gimple gimple_build_omp_target (gimple_seq, int, tree);
+gimple gimple_build_omp_teams (gimple_seq, tree);
 gimple gimple_build_cdt (tree, tree);
 gimple gimple_build_omp_atomic_load (tree, tree);
 gimple gimple_build_omp_atomic_store (tree);
@@ -1264,6 +1270,8 @@ gimple_has_substatements (gimple g)
     case GIMPLE_OMP_TASK:
     case GIMPLE_OMP_SECTIONS:
     case GIMPLE_OMP_SINGLE:
+    case GIMPLE_OMP_TARGET:
+    case GIMPLE_OMP_TEAMS:
     case GIMPLE_OMP_CRITICAL:
     case GIMPLE_WITH_CLEANUP_EXPR:
     case GIMPLE_TRANSACTION:
@@ -1691,7 +1699,7 @@ static inline unsigned
 gimple_omp_subcode (const_gimple s)
 {
   gcc_gimple_checking_assert (gimple_code (s) >= GIMPLE_OMP_ATOMIC_LOAD
-	      && gimple_code (s) <= GIMPLE_OMP_SINGLE);
+	      && gimple_code (s) <= GIMPLE_OMP_TEAMS);
   return s->gsbase.subcode;
 }
 
@@ -4604,6 +4612,87 @@ gimple_omp_single_set_clauses (gimple gs, tree clauses)
 }
 
 
+/* Return the clauses associated with OMP_TARGET GS.  */
+
+static inline tree
+gimple_omp_target_clauses (const_gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  return gs->gimple_omp_single.clauses;
+}
+
+
+/* Return a pointer to the clauses associated with OMP_TARGET GS.  */
+
+static inline tree *
+gimple_omp_target_clauses_ptr (gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  return &gs->gimple_omp_single.clauses;
+}
+
+
+/* Set CLAUSES to be the clauses associated with OMP_TARGET GS.  */
+
+static inline void
+gimple_omp_target_set_clauses (gimple gs, tree clauses)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TARGET);
+  gs->gimple_omp_single.clauses = clauses;
+}
+
+
+/* Return the kind of OMP target statemement.  */
+
+static inline int
+gimple_omp_target_kind (const_gimple g)
+{
+  GIMPLE_CHECK (g, GIMPLE_OMP_TARGET);
+  return (gimple_omp_subcode (g) & GF_OMP_TARGET_KIND_MASK);
+}
+
+
+/* Set the OMP target kind.  */
+
+static inline void
+gimple_omp_target_set_kind (gimple g, int kind)
+{
+  GIMPLE_CHECK (g, GIMPLE_OMP_TARGET);
+  g->gsbase.subcode = (g->gsbase.subcode & ~GF_OMP_TARGET_KIND_MASK)
+		      | (kind & GF_OMP_TARGET_KIND_MASK);
+}
+
+
+/* Return the clauses associated with OMP_TEAMS GS.  */
+
+static inline tree
+gimple_omp_teams_clauses (const_gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS);
+  return gs->gimple_omp_single.clauses;
+}
+
+
+/* Return a pointer to the clauses associated with OMP_TEAMS GS.  */
+
+static inline tree *
+gimple_omp_teams_clauses_ptr (gimple gs)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS);
+  return &gs->gimple_omp_single.clauses;
+}
+
+
+/* Set CLAUSES to be the clauses associated with OMP_TEAMS GS.  */
+
+static inline void
+gimple_omp_teams_set_clauses (gimple gs, tree clauses)
+{
+  GIMPLE_CHECK (gs, GIMPLE_OMP_TEAMS);
+  gs->gimple_omp_single.clauses = clauses;
+}
+
+
 /* Return the clauses associated with OMP_SECTIONS GS.  */
 
 static inline tree
@@ -4946,6 +5035,8 @@ gimple_return_set_retval (gimple gs, tree retval)
     case GIMPLE_OMP_SECTIONS:			\
     case GIMPLE_OMP_SECTIONS_SWITCH:		\
     case GIMPLE_OMP_SINGLE:			\
+    case GIMPLE_OMP_TARGET:			\
+    case GIMPLE_OMP_TEAMS:			\
     case GIMPLE_OMP_SECTION:			\
     case GIMPLE_OMP_MASTER:			\
     case GIMPLE_OMP_ORDERED:			\
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index ec377f8fc84..4fea48ab37a 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -57,10 +57,11 @@ enum gimplify_omp_var_data
   GOVD_LASTPRIVATE = 32,
   GOVD_REDUCTION = 64,
   GOVD_LOCAL = 128,
-  GOVD_DEBUG_PRIVATE = 256,
-  GOVD_PRIVATE_OUTER_REF = 512,
-  GOVD_LINEAR = 1024,
-  GOVD_ALIGNED = 2048,
+  GOVD_MAP = 256,
+  GOVD_DEBUG_PRIVATE = 512,
+  GOVD_PRIVATE_OUTER_REF = 1024,
+  GOVD_LINEAR = 2048,
+  GOVD_ALIGNED = 4096,
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -74,7 +75,10 @@ enum omp_region_type
   ORT_PARALLEL = 2,
   ORT_COMBINED_PARALLEL = 3,
   ORT_TASK = 4,
-  ORT_UNTIED_TASK = 5
+  ORT_UNTIED_TASK = 5,
+  ORT_TEAMS = 8,
+  ORT_TARGET_DATA = 16,
+  ORT_TARGET = 32
 };
 
 struct gimplify_omp_ctx
@@ -5829,6 +5833,9 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
      the parameters of the type.  */
   if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
     {
+      /* To be handled later.  */
+      gcc_assert ((flags & GOVD_MAP) == 0);
+
       /* Add the pointer replacement variable as PRIVATE if the variable
 	 replacement is private, else FIRSTPRIVATE since we'll need the
 	 address of the original variable either for SHARED, or for the
@@ -5870,6 +5877,9 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
     }
   else if (lang_hooks.decls.omp_privatize_by_reference (decl))
     {
+      /* To be handled later.  */
+      gcc_assert ((flags & GOVD_MAP) == 0);
+
       gcc_assert ((flags & GOVD_LOCAL) == 0);
       omp_firstprivatize_type_sizes (ctx, TREE_TYPE (decl));
 
@@ -5896,6 +5906,22 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
 				   tree decl2)
 {
   splay_tree_node n;
+  struct gimplify_omp_ctx *octx;
+
+  for (octx = ctx; octx; octx = octx->outer_context)
+    if (octx->region_type == ORT_TARGET)
+      {
+	n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+	if (n == NULL)
+	  {
+	    error ("threadprivate variable %qE used in target region",
+		   DECL_NAME (decl));
+	    error_at (octx->location, "enclosing target region");
+	    splay_tree_insert (octx->variables, (splay_tree_key)decl, 0);
+	  }
+	if (decl2)
+	  splay_tree_insert (octx->variables, (splay_tree_key)decl2, 0);
+      }
 
   if (ctx->region_type != ORT_UNTIED_TASK)
     return false;
@@ -5944,13 +5970,24 @@ 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)
+    {
+      if (n == NULL)
+	omp_add_variable (ctx, decl, GOVD_MAP | flags);
+      else
+	n->value |= flags;
+      ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
+      goto do_outer;
+    }
+
   if (n == NULL)
     {
       enum omp_clause_default_kind default_kind, kind;
       struct gimplify_omp_ctx *octx;
 
       if (ctx->region_type == ORT_WORKSHARE
-	  || ctx->region_type == ORT_SIMD)
+	  || ctx->region_type == ORT_SIMD
+	  || ctx->region_type == ORT_TARGET_DATA)
 	goto do_outer;
 
       /* ??? Some compiler-generated variables (like SAVE_EXPRs) could be
@@ -5964,12 +6001,24 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
       switch (default_kind)
 	{
 	case OMP_CLAUSE_DEFAULT_NONE:
-	  error ("%qE not specified in enclosing parallel",
-		 DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
 	  if ((ctx->region_type & ORT_TASK) != 0)
-	    error_at (ctx->location, "enclosing task");
+	    {
+	      error ("%qE not specified in enclosing task",
+		     DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
+	      error_at (ctx->location, "enclosing task");
+	    }
+	  else if (ctx->region_type == ORT_TEAMS)
+	    {
+	      error ("%qE not specified in enclosing teams construct",
+		     DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
+	      error_at (ctx->location, "enclosing teams construct");
+	    }
 	  else
-	    error_at (ctx->location, "enclosing parallel");
+	    {
+	      error ("%qE not specified in enclosing parallel",
+		     DECL_NAME (lang_hooks.decls.omp_report_decl (decl)));
+	      error_at (ctx->location, "enclosing parallel");
+	    }
 	  /* FALLTHRU */
 	case OMP_CLAUSE_DEFAULT_SHARED:
 	  flags |= GOVD_SHARED;
@@ -5989,13 +6038,15 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
 	    {
 	      splay_tree_node n2;
 
+	      if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0)
+		continue;
 	      n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
 	      if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED)
 		{
 		  flags |= GOVD_FIRSTPRIVATE;
 		  break;
 		}
-	      if ((octx->region_type & ORT_PARALLEL) != 0)
+	      if ((octx->region_type & (ORT_PARALLEL | ORT_TEAMS)) != 0)
 		break;
 	    }
 	  if (flags & GOVD_FIRSTPRIVATE)
@@ -6137,6 +6188,9 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl)
 		 /* References might be private, but might be shared too.  */
 		 || lang_hooks.decls.omp_privatize_by_reference (decl));
 
+      if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
+	continue;
+
       n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
       if (n != NULL)
 	return (n->value & GOVD_SHARED) == 0;
@@ -6204,6 +6258,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  flags = GOVD_LINEAR | GOVD_EXPLICIT;
 	  goto do_add;
+	case OMP_CLAUSE_MAP:
+	  flags = GOVD_MAP | GOVD_EXPLICIT;
+	  notice_outer = false;
+	  goto do_add;
+
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
+	  decl = OMP_CLAUSE_DECL (c);
+	  if (error_operand_p (decl))
+	    {
+	      remove = true;
+	      break;
+	    }
+	  goto do_notice;
 
 	do_add:
 	  decl = OMP_CLAUSE_DECL (c);
@@ -6292,6 +6360,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	case OMP_CLAUSE_SCHEDULE:
 	case OMP_CLAUSE_NUM_THREADS:
+	case OMP_CLAUSE_NUM_TEAMS:
+	case OMP_CLAUSE_DIST_SCHEDULE:
+	case OMP_CLAUSE_DEVICE:
 	  if (gimplify_expr (&OMP_CLAUSE_OPERAND (c, 0), pre_p, NULL,
 			     is_gimple_val, fb_rvalue) == GS_ERROR)
 	    remove = true;
@@ -6357,12 +6428,40 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
       gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_PRIVATE);
       private_debug = true;
     }
+  else if (flags & GOVD_MAP)
+    private_debug = false;
   else
     private_debug
       = lang_hooks.decls.omp_private_debug_clause (decl,
 						   !!(flags & GOVD_SHARED));
   if (private_debug)
     code = OMP_CLAUSE_PRIVATE;
+  else if (flags & GOVD_MAP)
+    {
+      /* If decl is already in the enclosing device data environment,
+	 the spec says that it should just be used and no init/assignment
+	 should be done.  If there was any privatization in between though,
+	 it means that original decl might be in the enclosing device data
+	 environment, but the privatized might not.  */
+      struct gimplify_omp_ctx *ctx;
+      for (ctx = gimplify_omp_ctxp->outer_context;
+	   ctx; ctx = ctx->outer_context)
+	{
+	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+	  if (n == NULL)
+	    continue;
+	  if (ctx->region_type == ORT_TARGET_DATA)
+	    {
+	      if ((n->value & GOVD_MAP) != 0)
+		return 0;
+	    }
+	  else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
+				| GOVD_PRIVATE | GOVD_REDUCTION
+				| GOVD_LINEAR)) != 0)
+	    break;
+	}
+      code = OMP_CLAUSE_MAP;
+    }
   else if (flags & GOVD_SHARED)
     {
       if (is_global_var (decl))
@@ -6373,7 +6472,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 	      splay_tree_node on
 		= splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
 	      if (on && (on->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
-				      | GOVD_PRIVATE | GOVD_REDUCTION)) != 0)
+				      | GOVD_PRIVATE | GOVD_REDUCTION
+				      | GOVD_LINEAR)) != 0)
 		break;
 	      ctx = ctx->outer_context;
 	    }
@@ -6400,6 +6500,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
     OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
   else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
     OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
+  else if (code == OMP_CLAUSE_MAP)
+    OMP_CLAUSE_MAP_KIND (clause) = OMP_CLAUSE_MAP_TOFROM;
   *list_p = clause;
   lang_hooks.decls.omp_finish_clause (clause);
 
@@ -6517,11 +6619,47 @@ gimplify_adjust_omp_clauses (tree *list_p)
 	    }
 	  break;
 
+	case OMP_CLAUSE_MAP:
+	  decl = OMP_CLAUSE_DECL (c);
+	  n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+	  remove = false;
+	  if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
+	    remove = true;
+	  else
+	    {
+	      /* If decl is already in the enclosing device data environment,
+		 the spec says that it should just be used and no init/assignment
+		 should be done.  If there was any privatization in between though,
+		 it means that original decl might be in the enclosing device data
+		 environment, but the privatized might not.  */
+	      struct gimplify_omp_ctx *octx;
+	      for (octx = ctx->outer_context; octx; octx = octx->outer_context)
+		{
+		  n = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) decl);
+		  if (n == NULL)
+		    continue;
+		  if (octx->region_type == ORT_TARGET_DATA)
+		    {
+		      if ((n->value & GOVD_MAP) != 0)
+			remove = true;
+		    }
+		  else if ((n->value & (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE
+					| GOVD_PRIVATE | GOVD_REDUCTION
+					| GOVD_LINEAR)) != 0)
+		    break;
+		}
+	    }
+	  break;
+
 	case OMP_CLAUSE_REDUCTION:
 	case OMP_CLAUSE_COPYIN:
 	case OMP_CLAUSE_COPYPRIVATE:
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
+	case OMP_CLAUSE_NUM_TEAMS:
+	case OMP_CLAUSE_DIST_SCHEDULE:
+	case OMP_CLAUSE_DEVICE:
 	case OMP_CLAUSE_SCHEDULE:
 	case OMP_CLAUSE_NOWAIT:
 	case OMP_CLAUSE_ORDERED:
@@ -6532,6 +6670,8 @@ gimplify_adjust_omp_clauses (tree *list_p)
 	case OMP_CLAUSE_MERGEABLE:
 	case OMP_CLAUSE_PROC_BIND:
 	case OMP_CLAUSE_SAFELEN:
+	case OMP_CLAUSE_TO:
+	case OMP_CLAUSE_FROM:
 	  break;
 
 	default:
@@ -6847,25 +6987,19 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 
   gimplify_adjust_omp_clauses (&OMP_FOR_CLAUSES (for_stmt));
 
-  gfor = gimple_build_omp_for (for_body, OMP_FOR_CLAUSES (for_stmt),
-			       TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)),
-			       for_pre_body);
+  int kind;
   switch (TREE_CODE (for_stmt))
     {
-    case OMP_FOR:
-      break;
-    case OMP_SIMD:
-      gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_SIMD);
-      break;
-    case OMP_FOR_SIMD:
-      gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_FOR_SIMD);
-      break;
-    case OMP_DISTRIBUTE:
-      gimple_omp_for_set_kind (gfor, GF_OMP_FOR_KIND_DISTRIBUTE);
-      break;
+    case OMP_FOR: kind = GF_OMP_FOR_KIND_FOR; break;
+    case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break;
+    case OMP_FOR_SIMD: kind = GF_OMP_FOR_KIND_FOR_SIMD; break;
+    case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
     default:
       gcc_unreachable ();
     }
+  gfor = gimple_build_omp_for (for_body, kind, OMP_FOR_CLAUSES (for_stmt),
+			       TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)),
+			       for_pre_body);
 
   for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++)
     {
@@ -6880,11 +7014,15 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
     }
 
   gimplify_seq_add_stmt (pre_p, gfor);
-  return ret == GS_ALL_DONE ? GS_ALL_DONE : GS_ERROR;
+  if (ret != GS_ALL_DONE)
+    return GS_ERROR;
+  *expr_p = NULL_TREE;
+  return GS_ALL_DONE;
 }
 
-/* Gimplify the gross structure of other OpenMP worksharing constructs.
-   In particular, OMP_SECTIONS and OMP_SINGLE.  */
+/* Gimplify the gross structure of other OpenMP constructs.
+   In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA
+   and OMP_TEAMS.  */
 
 static void
 gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
@@ -6892,19 +7030,72 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
   tree expr = *expr_p;
   gimple stmt;
   gimple_seq body = NULL;
+  enum omp_region_type ort = ORT_WORKSHARE;
 
-  gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ORT_WORKSHARE);
+  switch (TREE_CODE (expr))
+    {
+    case OMP_SECTIONS:
+    case OMP_SINGLE:
+      break;
+    case OMP_TARGET:
+      ort = ORT_TARGET;
+      break;
+    case OMP_TARGET_DATA:
+      ort = ORT_TARGET_DATA;
+      break;
+    case OMP_TEAMS:
+      ort = ORT_TEAMS;
+      break;
+    default:
+      gcc_unreachable ();
+    }
+  gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
   gimplify_and_add (OMP_BODY (expr), &body);
   gimplify_adjust_omp_clauses (&OMP_CLAUSES (expr));
 
-  if (TREE_CODE (expr) == OMP_SECTIONS)
-    stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
-  else if (TREE_CODE (expr) == OMP_SINGLE)
-    stmt = gimple_build_omp_single (body, OMP_CLAUSES (expr));
-  else
-    gcc_unreachable ();
+  switch (TREE_CODE (expr))
+    {
+    case OMP_SECTIONS:
+      stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
+      break;
+    case OMP_SINGLE:
+      stmt = gimple_build_omp_single (body, OMP_CLAUSES (expr));
+      break;
+    case OMP_TARGET:
+      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_REGION,
+				      OMP_CLAUSES (expr));
+      break;
+    case OMP_TARGET_DATA:
+      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA,
+				      OMP_CLAUSES (expr));
+      break;
+    case OMP_TEAMS:
+      stmt = gimple_build_omp_teams (body, OMP_CLAUSES (expr));
+      break;
+    default:
+      gcc_unreachable ();
+    }
+
+  gimplify_seq_add_stmt (pre_p, stmt);
+  *expr_p = NULL_TREE;
+}
+
+/* Gimplify the gross structure of OpenMP target update construct.  */
+
+static void
+gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
+{
+  tree expr = *expr_p;
+  gimple stmt;
+
+  gimplify_scan_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr), pre_p,
+			     ORT_WORKSHARE);
+  gimplify_adjust_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr));
+  stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_UPDATE,
+				  OMP_TARGET_UPDATE_CLAUSES (expr));
 
   gimplify_seq_add_stmt (pre_p, stmt);
+  *expr_p = NULL_TREE;
 }
 
 /* A subroutine of gimplify_omp_atomic.  The front end is supposed to have
@@ -7811,10 +8002,18 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 
 	case OMP_SECTIONS:
 	case OMP_SINGLE:
+	case OMP_TARGET:
+	case OMP_TARGET_DATA:
+	case OMP_TEAMS:
 	  gimplify_omp_workshare (expr_p, pre_p);
 	  ret = GS_ALL_DONE;
 	  break;
 
+	case OMP_TARGET_UPDATE:
+	  gimplify_omp_target_update (expr_p, pre_p);
+	  ret = GS_ALL_DONE;
+	  break;
+
 	case OMP_SECTION:
 	case OMP_MASTER:
 	case OMP_ORDERED:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 6d964840839..7c795624a7a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1483,6 +1483,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_SCHEDULE:
+	case OMP_CLAUSE_DIST_SCHEDULE:
 	  if (ctx->outer)
 	    scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
 	  break;
@@ -1548,6 +1549,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_IF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_SCHEDULE:
+	case OMP_CLAUSE_DIST_SCHEDULE:
 	case OMP_CLAUSE_NOWAIT:
 	case OMP_CLAUSE_ORDERED:
 	case OMP_CLAUSE_COLLAPSE:
diff --git a/gcc/tree-cfg.c b/gcc/tree-cfg.c
index 721c4f77d4a..c9a218bda5d 100644
--- a/gcc/tree-cfg.c
+++ b/gcc/tree-cfg.c
@@ -592,6 +592,8 @@ make_edges (void)
 	    case GIMPLE_OMP_TASK:
 	    case GIMPLE_OMP_FOR:
 	    case GIMPLE_OMP_SINGLE:
+	    case GIMPLE_OMP_TARGET:
+	    case GIMPLE_OMP_TEAMS:
 	    case GIMPLE_OMP_MASTER:
 	    case GIMPLE_OMP_ORDERED:
 	    case GIMPLE_OMP_CRITICAL:
diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c
index 57af7de2c09..b7b2b364dcc 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -1298,7 +1298,8 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
 	case GIMPLE_OMP_FOR:
 	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
 	  s2 = remap_gimple_seq (gimple_omp_for_pre_body (stmt), id);
-	  copy = gimple_build_omp_for (s1, gimple_omp_for_clauses (stmt),
+	  copy = gimple_build_omp_for (s1, gimple_omp_for_kind (stmt),
+				       gimple_omp_for_clauses (stmt),
 				       gimple_omp_for_collapse (stmt), s2);
 	  {
 	    size_t i;
@@ -1345,6 +1346,19 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
 	           (s1, gimple_omp_single_clauses (stmt));
 	  break;
 
+	case GIMPLE_OMP_TARGET:
+	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
+	  copy = gimple_build_omp_target
+		   (s1, gimple_omp_target_kind (stmt),
+		    gimple_omp_target_clauses (stmt));
+	  break;
+
+	case GIMPLE_OMP_TEAMS:
+	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
+	  copy = gimple_build_omp_teams
+		   (s1, gimple_omp_teams_clauses (stmt));
+	  break;
+
 	case GIMPLE_OMP_CRITICAL:
 	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
 	  copy
@@ -3716,6 +3730,8 @@ estimate_num_insns (gimple stmt, eni_weights *weights)
     case GIMPLE_OMP_SECTION:
     case GIMPLE_OMP_SECTIONS:
     case GIMPLE_OMP_SINGLE:
+    case GIMPLE_OMP_TARGET:
+    case GIMPLE_OMP_TEAMS:
       return (weights->omp_cost
               + estimate_num_insns_seq (gimple_omp_body (stmt), weights));
 
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index fe44679a013..41548821acf 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1291,6 +1291,22 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       info->suppress_expansion = save_suppress;
       break;
 
+    case GIMPLE_OMP_TARGET:
+      save_suppress = info->suppress_expansion;
+      convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
+      walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op,
+		 info, gimple_omp_body_ptr (stmt));
+      info->suppress_expansion = save_suppress;
+      break;
+
+    case GIMPLE_OMP_TEAMS:
+      save_suppress = info->suppress_expansion;
+      convert_nonlocal_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi);
+      walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op,
+		 info, gimple_omp_body_ptr (stmt));
+      info->suppress_expansion = save_suppress;
+      break;
+
     case GIMPLE_OMP_SECTION:
     case GIMPLE_OMP_MASTER:
     case GIMPLE_OMP_ORDERED:
@@ -1714,6 +1730,22 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       info->suppress_expansion = save_suppress;
       break;
 
+    case GIMPLE_OMP_TARGET:
+      save_suppress = info->suppress_expansion;
+      convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
+      walk_body (convert_local_reference_stmt, convert_local_reference_op,
+		 info, gimple_omp_body_ptr (stmt));
+      info->suppress_expansion = save_suppress;
+      break;
+
+    case GIMPLE_OMP_TEAMS:
+      save_suppress = info->suppress_expansion;
+      convert_local_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi);
+      walk_body (convert_local_reference_stmt, convert_local_reference_op,
+		 info, gimple_omp_body_ptr (stmt));
+      info->suppress_expansion = save_suppress;
+      break;
+
     case GIMPLE_OMP_SECTION:
     case GIMPLE_OMP_MASTER:
     case GIMPLE_OMP_ORDERED:
@@ -2071,6 +2103,8 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
     case GIMPLE_OMP_SECTIONS:
     case GIMPLE_OMP_SECTION:
     case GIMPLE_OMP_SINGLE:
+    case GIMPLE_OMP_TARGET:
+    case GIMPLE_OMP_TEAMS:
     case GIMPLE_OMP_MASTER:
     case GIMPLE_OMP_ORDERED:
     case GIMPLE_OMP_CRITICAL:
diff --git a/gcc/tree-parloops.c b/gcc/tree-parloops.c
index cea6f030c0a..1e6bb07b398 100644
--- a/gcc/tree-parloops.c
+++ b/gcc/tree-parloops.c
@@ -1686,7 +1686,7 @@ create_parallel_loop (struct loop *loop, tree loop_fn, tree data,
   t = build_omp_clause (loc, OMP_CLAUSE_SCHEDULE);
   OMP_CLAUSE_SCHEDULE_KIND (t) = OMP_CLAUSE_SCHEDULE_STATIC;
 
-  for_stmt = gimple_build_omp_for (NULL, t, 1, NULL);
+  for_stmt = gimple_build_omp_for (NULL, GF_OMP_FOR_KIND_FOR, t, 1, NULL);
   gimple_set_location (for_stmt, loc);
   gimple_omp_for_set_index (for_stmt, 0, initvar);
   gimple_omp_for_set_initial (for_stmt, 0, cvar_init);
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 2bdf9dcd232..ace9e22ebd1 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -2347,6 +2347,27 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
       pp_string (buffer, "#pragma omp distribute");
       goto dump_omp_loop;
 
+    case OMP_TEAMS:
+      pp_string (buffer, "#pragma omp teams");
+      dump_omp_clauses (buffer, OMP_TEAMS_CLAUSES (node), spc, flags);
+      goto dump_omp_body;
+
+    case OMP_TARGET_DATA:
+      pp_string (buffer, "#pragma omp target data");
+      dump_omp_clauses (buffer, OMP_TARGET_DATA_CLAUSES (node), spc, flags);
+      goto dump_omp_body;
+
+    case OMP_TARGET:
+      pp_string (buffer, "#pragma omp target");
+      dump_omp_clauses (buffer, OMP_TARGET_CLAUSES (node), spc, flags);
+      goto dump_omp_body;
+
+    case OMP_TARGET_UPDATE:
+      pp_string (buffer, "#pragma omp target update");
+      dump_omp_clauses (buffer, OMP_TARGET_UPDATE_CLAUSES (node), spc, flags);
+      is_expr = false;
+      break;
+
     dump_omp_loop:
       dump_omp_clauses (buffer, OMP_FOR_CLAUSES (node), spc, flags);
 
diff --git a/gcc/tree.def b/gcc/tree.def
index 147b6805b0a..a2e53e65f72 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1042,6 +1042,21 @@ DEFTREECODE (OMP_FOR_SIMD, "omp_for_simd", tcc_statement, 6)
    Operands like for OMP_FOR.  */
 DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 6)
 
+/* OpenMP - #pragma omp teams [clause1 ... clauseN]
+   Operand 0: OMP_TEAMS_BODY: Teams body.
+   Operand 1: OMP_TEAMS_CLAUSES: List of clauses.  */
+DEFTREECODE (OMP_TEAMS, "omp_teams", tcc_statement, 2)
+
+/* OpenMP - #pragma omp target data [clause1 ... clauseN]
+   Operand 0: OMP_TARGET_DATA_BODY: Target data construct body.
+   Operand 1: OMP_TARGET_DATA_CLAUSES: List of clauses.  */
+DEFTREECODE (OMP_TARGET_DATA, "omp_target_data", tcc_statement, 2)
+
+/* OpenMP - #pragma omp target [clause1 ... clauseN]
+   Operand 0: OMP_TARGET_BODY: Target construct body.
+   Operand 1: OMP_TARGET_CLAUSES: List of clauses.  */
+DEFTREECODE (OMP_TARGET, "omp_target", tcc_statement, 2)
+
 /* OpenMP - #pragma omp sections [clause1 ... clauseN]
    Operand 0: OMP_SECTIONS_BODY: Sections body.
    Operand 1: OMP_SECTIONS_CLAUSES: List of clauses.  */
@@ -1069,6 +1084,10 @@ DEFTREECODE (OMP_ORDERED, "omp_ordered", tcc_statement, 1)
    Operand 1: OMP_CRITICAL_NAME: Identifier for critical section.  */
 DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2)
 
+/* OpenMP - #pragma omp target update [clause1 ... clauseN]
+   Operand 0: OMP_TARGET_UPDATE_CLAUSES: List of clauses.  */
+DEFTREECODE (OMP_TARGET_UPDATE, "omp_target_update", tcc_statement, 1)
+
 /* OMP_ATOMIC through OMP_ATOMIC_CAPTURE_NEW must be consecutive,
    or OMP_ATOMIC_SEQ_CST needs adjusting.  */
 
diff --git a/gcc/tree.h b/gcc/tree.h
index 49da30a3186..f2a58a09e0d 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1863,6 +1863,20 @@ extern void protected_set_expr_location (tree, location_t);
 #define OMP_CRITICAL_BODY(NODE)    TREE_OPERAND (OMP_CRITICAL_CHECK (NODE), 0)
 #define OMP_CRITICAL_NAME(NODE)    TREE_OPERAND (OMP_CRITICAL_CHECK (NODE), 1)
 
+#define OMP_TEAMS_BODY(NODE)	   TREE_OPERAND (OMP_TEAMS_CHECK (NODE), 0)
+#define OMP_TEAMS_CLAUSES(NODE)	   TREE_OPERAND (OMP_TEAMS_CHECK (NODE), 1)
+
+#define OMP_TARGET_DATA_BODY(NODE) \
+  TREE_OPERAND (OMP_TARGET_DATA_CHECK (NODE), 0)
+#define OMP_TARGET_DATA_CLAUSES(NODE)\
+  TREE_OPERAND (OMP_TARGET_DATA_CHECK (NODE), 1)
+
+#define OMP_TARGET_BODY(NODE)	   TREE_OPERAND (OMP_TARGET_CHECK (NODE), 0)
+#define OMP_TARGET_CLAUSES(NODE)   TREE_OPERAND (OMP_TARGET_CHECK (NODE), 1)
+
+#define OMP_TARGET_UPDATE_CLAUSES(NODE)\
+  TREE_OPERAND (OMP_TARGET_UPDATE_CHECK (NODE), 0)
+
 #define OMP_CLAUSE_CHAIN(NODE)     TREE_CHAIN (OMP_CLAUSE_CHECK (NODE))
 #define OMP_CLAUSE_DECL(NODE)      					\
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE),	\