diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index cd3bd5a..211857a 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11701,6 +11701,95 @@ c_parser_omp_clause_simdlen (c_parser *parser, tree list)
   return c;
 }
 
+/* OpenMP 4.1:
+   vec:
+     identifier [+/- integer]
+     vec , identifier [+/- integer]
+*/
+
+static tree
+c_parser_omp_clause_depend_sink (c_parser *parser, location_t clause_loc,
+				 tree list)
+{
+  tree vec = NULL;
+  if (c_parser_next_token_is_not (parser, CPP_NAME)
+      || c_parser_peek_token (parser)->id_kind != C_ID_ID)
+    {
+      c_parser_error (parser, "expected identifier");
+      return list;
+    }
+
+  while (c_parser_next_token_is (parser, CPP_NAME)
+	 && c_parser_peek_token (parser)->id_kind == C_ID_ID)
+    {
+      tree t = lookup_name (c_parser_peek_token (parser)->value);
+      tree addend = NULL;
+
+      if (t == NULL_TREE)
+	{
+	  undeclared_variable (c_parser_peek_token (parser)->location,
+			       c_parser_peek_token (parser)->value);
+	  t = error_mark_node;
+	}
+
+      c_parser_consume_token (parser);
+
+      if (t != error_mark_node)
+	{
+	  bool neg;
+
+	  if (c_parser_next_token_is (parser, CPP_MINUS))
+	    neg = true;
+	  else if (c_parser_next_token_is (parser, CPP_PLUS))
+	    neg = false;
+	  else
+	    {
+	      addend = integer_zero_node;
+	      goto add_to_vector;
+	    }
+	  c_parser_consume_token (parser);
+
+	  if (c_parser_next_token_is_not (parser, CPP_NUMBER))
+	    {
+	      c_parser_error (parser, "expected integer");
+	      return list;
+	    }
+
+	  addend = c_parser_peek_token (parser)->value;
+	  if (TREE_CODE (addend) != INTEGER_CST)
+	    {
+	      c_parser_error (parser, "expected integer");
+	      return list;
+	    }
+	  if (neg)
+	    {
+	      bool overflow;
+	      wide_int offset = wi::neg (addend, &overflow);
+	      addend = wide_int_to_tree (TREE_TYPE (addend), offset);
+	      if (overflow)
+		warning_at (c_parser_peek_token (parser)->location,
+			    OPT_Woverflow,
+			    "overflow in implicit constant conversion");
+	    }
+	  c_parser_consume_token (parser);
+
+	add_to_vector:
+	  vec = tree_cons (addend, t, vec);
+
+	  if (c_parser_next_token_is_not (parser, CPP_COMMA))
+	    break;
+
+	  c_parser_consume_token (parser);
+	}
+    }
+
+  tree u = build_omp_clause (clause_loc, OMP_CLAUSE_DEPEND);
+  OMP_CLAUSE_DEPEND_KIND (u) = OMP_CLAUSE_DEPEND_SINK;
+  OMP_CLAUSE_DECL (u) = nreverse (vec);
+  OMP_CLAUSE_CHAIN (u) = list;
+  return u;
+}
+
 /* OpenMP 4.0:
    depend ( depend-kind: variable-list )
 
@@ -11708,10 +11797,9 @@ c_parser_omp_clause_simdlen (c_parser *parser, tree list)
      in | out | inout
 
    OpenMP 4.1:
-   depend ( depend-loop-kind [ : vec ] )
+   depend ( source )
 
-   depend-loop-kind:
-     source | sink  */
+   depend ( sink  : vec )  */
 
 static tree
 c_parser_omp_clause_depend (c_parser *parser, tree list)
@@ -11754,16 +11842,19 @@ c_parser_omp_clause_depend (c_parser *parser, tree list)
       return c;
     }
 
-  /* FIXME: Handle OMP_CLAUSE_DEPEND_SINK.  */
-
   if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
     goto resync_fail;
 
-  nl = c_parser_omp_variable_list (parser, clause_loc,
-				   OMP_CLAUSE_DEPEND, list);
+  if (kind == OMP_CLAUSE_DEPEND_SINK)
+    nl = c_parser_omp_clause_depend_sink (parser, clause_loc, list);
+  else
+    {
+      nl = c_parser_omp_variable_list (parser, clause_loc,
+				       OMP_CLAUSE_DEPEND, list);
 
-  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_DEPEND_KIND (c) = kind;
+      for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+	OMP_CLAUSE_DEPEND_KIND (c) = kind;
+    }
 
   c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
   return nl;
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 469cd88..0b332e8 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12489,6 +12489,11 @@ c_finish_omp_clauses (tree clauses, bool declare_simd)
 			  == OMP_CLAUSE_DEPEND_SOURCE);
 	      break;
 	    }
+	  if (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK)
+	    {
+	      gcc_assert (TREE_CODE (t) == TREE_LIST);
+	      break;
+	    }
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
 	      if (handle_omp_array_sections (c))
diff --git a/gcc/coretypes.h b/gcc/coretypes.h
index e3cec01..a708896 100644
--- a/gcc/coretypes.h
+++ b/gcc/coretypes.h
@@ -113,6 +113,7 @@ struct gomp_atomic_load;
 struct gomp_atomic_store;
 struct gomp_continue;
 struct gomp_critical;
+struct gomp_ordered;
 struct gomp_for;
 struct gomp_parallel;
 struct gomp_task;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index b18f54f..64a8aec 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -29314,6 +29314,99 @@ cp_parser_omp_clause_simdlen (cp_parser *parser, tree list,
   return c;
 }
 
+/* OpenMP 4.1:
+   vec:
+     identifier [+/- integer]
+     vec , identifier [+/- integer]
+*/
+
+static tree
+cp_parser_omp_clause_depend_sink (cp_parser *parser, location_t clause_loc,
+				  tree list)
+{
+  tree vec = NULL;
+
+  if (cp_lexer_next_token_is_not (parser->lexer, CPP_NAME))
+    {
+      cp_parser_error (parser, "expected identifier");
+      return list;
+    }
+
+  while (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+    {
+      location_t id_loc = cp_lexer_peek_token (parser->lexer)->location;
+      tree t, identifier = cp_parser_identifier (parser);
+      tree addend = NULL;
+
+      if (identifier == error_mark_node)
+	t = error_mark_node;
+      else
+	{
+	  t = cp_parser_lookup_name_simple
+	        (parser, identifier,
+		 cp_lexer_peek_token (parser->lexer)->location);
+	  if (t == error_mark_node)
+	    cp_parser_name_lookup_error (parser, identifier, t, NLE_NULL,
+					 id_loc);
+	}
+
+      if (t != error_mark_node)
+	{
+	  bool neg;
+
+	  if (cp_lexer_next_token_is (parser->lexer, CPP_MINUS))
+	    neg = true;
+	  else if (cp_lexer_next_token_is (parser->lexer, CPP_PLUS))
+	    neg = false;
+	  else
+	    {
+	      addend = integer_zero_node;
+	      goto add_to_vector;
+	    }
+	  cp_lexer_consume_token (parser->lexer);
+
+	  if (cp_lexer_next_token_is_not (parser->lexer, CPP_NUMBER))
+	    {
+	      cp_parser_error (parser, "expected integer");
+	      return list;
+	    }
+
+	  addend = cp_lexer_peek_token (parser->lexer)->u.value;
+	  if (TREE_CODE (addend) != INTEGER_CST)
+	    {
+	      cp_parser_error (parser, "expected integer");
+	      return list;
+	    }
+	  if (neg)
+	    {
+	      bool overflow;
+	      wide_int offset = wi::neg (addend, &overflow);
+	      addend = wide_int_to_tree (TREE_TYPE (addend), offset);
+	      if (overflow)
+		warning_at (cp_lexer_peek_token (parser->lexer)->location,
+			    OPT_Woverflow,
+			    "overflow in implicit constant conversion");
+	    }
+	  cp_lexer_consume_token (parser->lexer);
+
+	add_to_vector:
+	  vec = tree_cons (addend, t, vec);
+
+	  if (cp_lexer_next_token_is_not (parser->lexer, CPP_COMMA))
+	    break;
+
+	  cp_lexer_consume_token (parser->lexer);
+	}
+    }
+  cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN);
+
+  tree u = build_omp_clause (clause_loc, OMP_CLAUSE_DEPEND);
+  OMP_CLAUSE_DEPEND_KIND (u) = OMP_CLAUSE_DEPEND_SINK;
+  OMP_CLAUSE_DECL (u) = nreverse (vec);
+  OMP_CLAUSE_CHAIN (u) = list;
+  return u;
+}
+
 /* OpenMP 4.0:
    depend ( depend-kind : variable-list )
 
@@ -29321,10 +29414,9 @@ cp_parser_omp_clause_simdlen (cp_parser *parser, tree list,
      in | out | inout
 
    OpenMP 4.1:
-   depend ( depend-loop-kind [ : vec ] )
+   depend ( source )
 
-   depend-loop-kind:
-     source | sink  */
+   depend ( sink : vec ) */
 
 static tree
 cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc)
@@ -29371,17 +29463,19 @@ cp_parser_omp_clause_depend (cp_parser *parser, tree list, location_t loc)
       return c;
     }
 
-  /* FIXME: Handle OMP_CLAUSE_DEPEND_SINK.  */
-
   if (!cp_parser_require (parser, CPP_COLON, RT_COLON))
     goto resync_fail;
 
-  nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_DEPEND, list,
-					  NULL);
-
-  for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_DEPEND_KIND (c) = kind;
+  if (kind == OMP_CLAUSE_DEPEND_SINK)
+    nlist = cp_parser_omp_clause_depend_sink (parser, loc, list);
+  else
+    {
+      nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_DEPEND,
+					      list, NULL);
 
+      for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
+	OMP_CLAUSE_DEPEND_KIND (c) = kind;
+    }
   return nlist;
 
  invalid_kind:
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 6e3aabd..a7b19ba 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -13532,13 +13532,21 @@ tsubst_copy (tree t, tree args, tsubst_flags_t complain, tree in_decl)
 }
 
 /* Helper function for tsubst_omp_clauses, used for instantiation of
-   OMP_CLAUSE_DECL of clauses that handles also OpenMP array sections
-   represented with TREE_LIST.  */
+   OMP_CLAUSE_DECL of clauses.  */
 
 static tree
 tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain,
 			tree in_decl)
 {
+  if (decl == NULL_TREE)
+    return NULL_TREE;
+
+  /* Handle an OpenMP array section represented as a TREE_LIST (or
+     OMP_CLAUSE_DEPEND_KIND).  An OMP_CLAUSE_DEPEND (with a depend
+     kind of OMP_CLAUSE_DEPEND_SINK) can also be represented as a
+     TREE_LIST.  We can handle it exactly the same as an array section
+     (purpose, value, and a chain), even though the nomenclature
+     (low_bound, length, etc) is different.  */
   if (TREE_CODE (decl) == TREE_LIST)
     {
       tree low_bound
@@ -14542,10 +14550,22 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       add_stmt (t);
       break;
 
+    case OMP_ORDERED:
+      tmp = tsubst_omp_clauses (OMP_ORDERED_CLAUSES (t), false, true,
+				args, complain, in_decl);
+      stmt = push_stmt_list ();
+      RECUR (OMP_BODY (t));
+      stmt = pop_stmt_list (stmt);
+
+      t = copy_node (t);
+      OMP_BODY (t) = stmt;
+      OMP_ORDERED_CLAUSES (t) = tmp;
+      add_stmt (t);
+      break;
+
     case OMP_SECTION:
     case OMP_MASTER:
     case OMP_TASKGROUP:
-    case OMP_ORDERED:
       stmt = push_stmt_list ();
       RECUR (OMP_BODY (t));
       stmt = pop_stmt_list (stmt);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index cca6ea1..3ad4c36 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6148,6 +6148,11 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
 			  == OMP_CLAUSE_DEPEND_SOURCE);
 	      break;
 	    }
+	  if (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK)
+	    {
+	      gcc_assert (TREE_CODE (t) == TREE_LIST);
+	      break;
+	    }
 	  if (TREE_CODE (t) == TREE_LIST)
 	    {
 	      if (handle_omp_array_sections (c))
diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index cae8883..83ce3ec 100644
--- a/gcc/gimple-pretty-print.c
+++ b/gcc/gimple-pretty-print.c
@@ -1488,9 +1488,6 @@ dump_gimple_omp_block (pretty_printer *buffer, gimple gs, int spc, int flags)
 	case GIMPLE_OMP_TASKGROUP:
 	  pp_string (buffer, "#pragma omp taskgroup");
 	  break;
-	case GIMPLE_OMP_ORDERED:
-	  pp_string (buffer, "#pragma omp ordered");
-	  break;
 	case GIMPLE_OMP_SECTION:
 	  pp_string (buffer, "#pragma omp section");
 	  break;
@@ -1541,6 +1538,31 @@ dump_gimple_omp_critical (pretty_printer *buffer, gomp_critical *gs,
     }
 }
 
+/* Dump a GIMPLE_OMP_ORDERED tuple on the pretty_printer BUFFER.  */
+
+static void
+dump_gimple_omp_ordered (pretty_printer *buffer, gomp_ordered *gs,
+			 int spc, int flags)
+{
+  if (flags & TDF_RAW)
+    dump_gimple_fmt (buffer, spc, flags, "%G <%+BODY <%S> >", gs,
+		     gimple_omp_body (gs));
+  else
+    {
+      pp_string (buffer, "#pragma omp ordered");
+      dump_omp_clauses (buffer, gimple_omp_ordered_clauses (gs), spc, flags);
+      if (!gimple_seq_empty_p (gimple_omp_body (gs)))
+	{
+	  newline_and_indent (buffer, spc + 2);
+	  pp_left_brace (buffer);
+	  pp_newline (buffer);
+	  dump_gimple_seq (buffer, gimple_omp_body (gs), spc + 4, flags);
+	  newline_and_indent (buffer, spc + 2);
+	  pp_right_brace (buffer);
+	}
+    }
+}
+
 /* Dump a GIMPLE_OMP_RETURN tuple on the pretty_printer BUFFER.  */
 
 static void
@@ -2250,11 +2272,15 @@ pp_gimple_stmt_1 (pretty_printer *buffer, gimple gs, int spc, int flags)
 
     case GIMPLE_OMP_MASTER:
     case GIMPLE_OMP_TASKGROUP:
-    case GIMPLE_OMP_ORDERED:
     case GIMPLE_OMP_SECTION:
       dump_gimple_omp_block (buffer, gs, spc, flags);
       break;
 
+    case GIMPLE_OMP_ORDERED:
+      dump_gimple_omp_ordered (buffer, as_a <gomp_ordered *> (gs), spc,
+			       flags);
+      break;
+
     case GIMPLE_OMP_CRITICAL:
       dump_gimple_omp_critical (buffer, as_a <gomp_critical *> (gs), spc,
 				flags);
diff --git a/gcc/gimple-walk.c b/gcc/gimple-walk.c
index f0e2c67..1fe4365 100644
--- a/gcc/gimple-walk.c
+++ b/gcc/gimple-walk.c
@@ -327,6 +327,16 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op,
       }
       break;
 
+    case GIMPLE_OMP_ORDERED:
+      {
+	gomp_ordered *omp_stmt = as_a <gomp_ordered *> (stmt);
+	ret = walk_tree (gimple_omp_ordered_clauses_ptr (omp_stmt),
+			 callback_op, wi, pset);
+	if (ret)
+	  return ret;
+      }
+      break;
+
     case GIMPLE_OMP_FOR:
       ret = walk_tree (gimple_omp_for_clauses_ptr (stmt), callback_op, wi,
 		       pset);
diff --git a/gcc/gimple.c b/gcc/gimple.c
index d8eb77a..b87d44d 100644
--- a/gcc/gimple.c
+++ b/gcc/gimple.c
@@ -991,12 +991,15 @@ gimple_build_omp_continue (tree control_def, tree control_use)
 /* Build a GIMPLE_OMP_ORDERED statement.
 
    BODY is the sequence of statements inside a loop that will executed in
-   sequence.  */
+   sequence.
+   CLAUSES are clauses for this statement.  */
 
-gimple
-gimple_build_omp_ordered (gimple_seq body)
+gomp_ordered *
+gimple_build_omp_ordered (gimple_seq body, tree clauses)
 {
-  gimple p = gimple_alloc (GIMPLE_OMP_ORDERED, 0);
+  gomp_ordered *p
+    = as_a <gomp_ordered *> (gimple_alloc (GIMPLE_OMP_ORDERED, 0));
+  gimple_omp_ordered_set_clauses (p, clauses);
   if (body)
     gimple_omp_set_body (p, body);
 
diff --git a/gcc/gimple.def b/gcc/gimple.def
index 96602df..d3ca402 100644
--- a/gcc/gimple.def
+++ b/gcc/gimple.def
@@ -283,10 +283,6 @@ DEFGSCODE(GIMPLE_OMP_MASTER, "gimple_omp_master", GSS_OMP)
    BODY is the sequence of statements to execute in the taskgroup section.  */
 DEFGSCODE(GIMPLE_OMP_TASKGROUP, "gimple_omp_taskgroup", GSS_OMP)
 
-/* GIMPLE_OMP_ORDERED <BODY> represents #pragma omp ordered.
-   BODY is the sequence of statements to execute in the ordered section.  */
-DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_omp_ordered", GSS_OMP)
-
 /* GIMPLE_OMP_PARALLEL <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
 
    #pragma omp parallel [CLAUSES]
@@ -375,6 +371,11 @@ DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_PARALLEL_LAYOUT)
    CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
 DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_SINGLE_LAYOUT)
 
+/* GIMPLE_OMP_ORDERED <BODY, CLAUSES> represents #pragma omp ordered.
+   BODY is the sequence of statements to execute in the ordered section.
+   CLAUSES is an OMP_CLAUSE chain holding the associated clauses.  */
+DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_omp_ordered", GSS_OMP_SINGLE_LAYOUT)
+
 /* 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 6057ea0..65f662f 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -668,7 +668,7 @@ struct GTY((tag("GSS_OMP_CONTINUE")))
   tree control_use;
 };
 
-/* GIMPLE_OMP_SINGLE, GIMPLE_OMP_TEAMS */
+/* GIMPLE_OMP_SINGLE, GIMPLE_OMP_TEAMS, GIMPLE_OMP_ORDERED */
 
 struct GTY((tag("GSS_OMP_SINGLE_LAYOUT")))
   gimple_statement_omp_single_layout : public gimple_statement_omp
@@ -693,6 +693,13 @@ struct GTY((tag("GSS_OMP_SINGLE_LAYOUT")))
          stmt->code == GIMPLE_OMP_TEAMS.  */
 };
 
+struct GTY((tag("GSS_OMP_SINGLE_LAYOUT")))
+  gomp_ordered : public gimple_statement_omp_single_layout
+{
+    /* No extra fields; adds invariant:
+         stmt->code == GIMPLE_OMP_ORDERED.  */
+};
+
 
 /* GIMPLE_OMP_ATOMIC_LOAD.
    Note: This is based on gimple_statement_base, not g_s_omp, because g_s_omp
@@ -1007,6 +1014,14 @@ is_a_helper <gomp_critical *>::test (gimple gs)
 template <>
 template <>
 inline bool
+is_a_helper <gomp_ordered *>::test (gimple gs)
+{
+  return gs->code == GIMPLE_OMP_ORDERED;
+}
+
+template <>
+template <>
+inline bool
 is_a_helper <gomp_for *>::test (gimple gs)
 {
   return gs->code == GIMPLE_OMP_FOR;
@@ -1215,6 +1230,14 @@ is_a_helper <const gomp_critical *>::test (const_gimple gs)
 template <>
 template <>
 inline bool
+is_a_helper <const gomp_ordered *>::test (const_gimple gs)
+{
+  return gs->code == GIMPLE_OMP_ORDERED;
+}
+
+template <>
+template <>
+inline bool
 is_a_helper <const gomp_for *>::test (const_gimple gs)
 {
   return gs->code == GIMPLE_OMP_FOR;
@@ -1355,7 +1378,7 @@ gimple gimple_build_omp_section (gimple_seq);
 gimple gimple_build_omp_master (gimple_seq);
 gimple gimple_build_omp_taskgroup (gimple_seq);
 gomp_continue *gimple_build_omp_continue (tree, tree);
-gimple gimple_build_omp_ordered (gimple_seq);
+gomp_ordered *gimple_build_omp_ordered (gimple_seq, tree);
 gimple gimple_build_omp_return (bool);
 gomp_sections *gimple_build_omp_sections (gimple_seq, tree);
 gimple gimple_build_omp_sections_switch (void);
@@ -4460,6 +4483,35 @@ gimple_omp_critical_set_clauses (gomp_critical *crit_stmt, tree clauses)
 }
 
 
+/* Return the clauses associated with OMP_ORDERED statement ORD_STMT.  */
+
+static inline tree
+gimple_omp_ordered_clauses (const gomp_ordered *ord_stmt)
+{
+  return ord_stmt->clauses;
+}
+
+
+/* Return a pointer to the clauses associated with OMP ordered statement
+   ORD_STMT.  */
+
+static inline tree *
+gimple_omp_ordered_clauses_ptr (gomp_ordered *ord_stmt)
+{
+  return &ord_stmt->clauses;
+}
+
+
+/* Set CLAUSES to be the clauses associated with OMP ordered statement
+   ORD_STMT.  */
+
+static inline void
+gimple_omp_ordered_set_clauses (gomp_ordered *ord_stmt, tree clauses)
+{
+  ord_stmt->clauses = clauses;
+}
+
+
 /* Return the kind of the OMP_FOR statemement G.  */
 
 static inline int
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 21f8223..2273ac1 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -149,6 +149,8 @@ struct gimplify_omp_ctx
   struct gimplify_omp_ctx *outer_context;
   splay_tree variables;
   hash_set<tree> *privatized_types;
+  /* Iteration variables in an OMP_FOR.  */
+  vec<tree> iter_vars;
   location_t location;
   enum omp_clause_default_kind default_kind;
   enum omp_region_type region_type;
@@ -384,6 +386,7 @@ delete_omp_context (struct gimplify_omp_ctx *c)
 {
   splay_tree_delete (c->variables);
   delete c->privatized_types;
+  c->iter_vars.release ();
   XDELETE (c);
 }
 
@@ -6343,6 +6346,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  goto do_add;
 
 	case OMP_CLAUSE_DEPEND:
+	  if (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK
+	      || OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SOURCE)
+	    {
+	      /* Nothing to do.  OMP_CLAUSE_DECL will be lowered in
+		 omp-low.c.  */
+	      break;
+	    }
 	  if (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPOUND_EXPR)
 	    {
 	      gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (c), 0), pre_p,
@@ -7282,6 +7292,8 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	      == TREE_VEC_LENGTH (OMP_FOR_COND (for_stmt)));
   gcc_assert (TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt))
 	      == TREE_VEC_LENGTH (OMP_FOR_INCR (for_stmt)));
+  gimplify_omp_ctxp->iter_vars.create (TREE_VEC_LENGTH
+				       (OMP_FOR_INIT (for_stmt)));
   for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++)
     {
       t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
@@ -7290,6 +7302,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       gcc_assert (DECL_P (decl));
       gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
 		  || POINTER_TYPE_P (TREE_TYPE (decl)));
+      gimplify_omp_ctxp->iter_vars.quick_push (decl);
 
       /* Make sure the iteration variable is private.  */
       tree c = NULL_TREE;
@@ -8982,7 +8995,38 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 		}
 		break;
 	      case OMP_ORDERED:
-		g = gimple_build_omp_ordered (body);
+		if (gimplify_omp_ctxp)
+		  for (tree c = OMP_ORDERED_CLAUSES (*expr_p);
+		       c; c = OMP_CLAUSE_CHAIN (c))
+		    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
+			&& OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK)
+		      {
+			unsigned int n = 0;
+			bool fail = false;
+			for (tree decls = OMP_CLAUSE_DECL (c);
+			     decls && TREE_CODE (decls) == TREE_LIST;
+			     decls = TREE_CHAIN (decls), ++n)
+			  if (n < gimplify_omp_ctxp->iter_vars.length ()
+			      && TREE_VALUE (decls)
+			      != gimplify_omp_ctxp->iter_vars[n])
+			    {
+			      error_at (OMP_CLAUSE_LOCATION (c),
+					"variable %qE is not an iteration "
+					"of outermost loop %d, expected %qE",
+					TREE_VALUE (decls), n + 1,
+					gimplify_omp_ctxp->iter_vars[n]);
+			      fail = true;
+			    }
+			/* Avoid being too redundant.  */
+			if (!fail
+			    && n != gimplify_omp_ctxp->iter_vars.length ())
+			  error_at (OMP_CLAUSE_LOCATION (c),
+			     "number of variables in depend(sink) clause "
+			     "does not match number of iteration variables");
+		      }
+
+		g = gimple_build_omp_ordered (body,
+					      OMP_ORDERED_CLAUSES (*expr_p));
 		break;
 	      case OMP_CRITICAL:
 		gimplify_scan_omp_clauses (&OMP_CRITICAL_CLAUSES (*expr_p),
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 83677ea..e2f00ef 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2996,6 +2996,8 @@ scan_omp_teams (gomp_teams *stmt, omp_context *outer_ctx)
 static bool
 check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 {
+  tree c;
+
   /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin)
      inside an OpenACC CTX.  */
   if (!(is_gimple_omp (stmt)
@@ -3216,7 +3218,57 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
 	    break;
 	  }
       break;
+    case GIMPLE_OMP_TASK:
+      for (c = gimple_omp_task_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
+	    && (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SOURCE
+		|| OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK))
+	  {
+	    enum omp_clause_depend_kind kind = OMP_CLAUSE_DEPEND_KIND (c);
+	    gcc_assert (kind == OMP_CLAUSE_DEPEND_SOURCE
+			|| kind == OMP_CLAUSE_DEPEND_SINK);
+	    error_at (OMP_CLAUSE_LOCATION (c),
+		      "%<depend(%s)%> is only allowed in %<omp ordered%>",
+		      kind == OMP_CLAUSE_DEPEND_SOURCE ? "source" : "sink");
+	    return false;
+	  }
+      break;
     case GIMPLE_OMP_ORDERED:
+      for (c = gimple_omp_ordered_clauses (as_a <gomp_ordered *> (stmt));
+	   c; c = OMP_CLAUSE_CHAIN (c))
+	{
+	  enum omp_clause_depend_kind kind = OMP_CLAUSE_DEPEND_KIND (c);
+	  if (kind == OMP_CLAUSE_DEPEND_SOURCE
+	      || kind == OMP_CLAUSE_DEPEND_SINK)
+	    {
+	      bool have_ordered = false;
+	      tree oclause;
+	      /* Look for containing ordered(N) loop.  */
+	      for (omp_context *octx = ctx; octx; octx = octx->outer)
+		if (gimple_code (octx->stmt) == GIMPLE_OMP_FOR
+		    && (oclause = find_omp_clause
+			            (gimple_omp_for_clauses (octx->stmt),
+				     OMP_CLAUSE_ORDERED))
+		    && OMP_CLAUSE_ORDERED_EXPR (oclause) != NULL_TREE)
+		  {
+		    have_ordered = true;
+		    break;
+		  }
+	      if (!have_ordered)
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "%<depend%> clause must be closely nested "
+			    "inside an ordered loop");
+		  return false;
+		}
+	    }
+	  else
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"invalid depend kind in omp ordered depend");
+	      return false;
+	    }
+	}
       for (; ctx != NULL; ctx = ctx->outer)
 	switch (gimple_code (ctx->stmt))
 	  {
@@ -12562,6 +12614,10 @@ lower_depend_clauses (gimple stmt, gimple_seq *iseq, gimple_seq *oseq)
 	case OMP_CLAUSE_DEPEND_INOUT:
 	  n_out++;
 	  break;
+	case OMP_CLAUSE_DEPEND_SOURCE:
+	case OMP_CLAUSE_DEPEND_SINK:
+	  /* FIXME:  */
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/testsuite/c-c++-common/gomp/sink-1.c b/gcc/testsuite/c-c++-common/gomp/sink-1.c
new file mode 100644
index 0000000..1a5e3e9
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/sink-1.c
@@ -0,0 +1,96 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -Wunknown-pragmas -Werror" } */
+
+extern void bark (void);
+int i,j,k;
+int array[555];
+
+int
+main()
+{
+#pragma omp parallel for ordered(2)
+  for (i=0; i < 100; ++i)
+    for (j=0; j < 100; ++j)
+      {
+/* OUT variant does not apply to ORDERED construct.  */
+#pragma omp ordered depend(out:i) /* { dg-error "invalid depend kind" } */
+
+/* depend(sink...) is allowed without an offset.  */
+#pragma omp ordered depend(sink:i,j-1)
+
+#pragma omp ordered depend(sink:i-1,j+2)
+      bark ();
+      }
+
+/* depend(sink...) does not apply to `omp task'.  */
+#pragma omp task depend(sink:i+3) /* { dg-error "only allowed in 'omp ordered'" } */
+  bark();
+
+#pragma omp ordered depend(source) /* { dg-error "'depend' clause must be closely nested" } */
+
+#pragma omp parallel for ordered(2)
+  for (i=0; i < 100; ++i)
+    for (j=0; j < 100; ++j)
+      {
+/* Multiple depend(source) allowed.  */
+#pragma omp ordered depend(source)
+#pragma omp ordered depend(source)
+      }
+
+#pragma omp parallel for ordered(2)
+  for (i=0; i < 100; ++i)
+    for (j=0; j < 100; ++j)
+      {
+#pragma omp ordered depend(sink:i-2,j-2,k+2) /* { dg-error "does not match number of iteration var" } */
+	bark();
+      }
+
+#pragma omp parallel for ordered(2)
+  for (i=0; i < 100; ++i)
+    for (j=0; j < 100; ++j)
+      {
+#pragma omp ordered depend(sink:i-2) /* { dg-error "does not match number of iteration variables" } */
+	bark();
+      }
+
+#pragma omp parallel for ordered(2)
+  for (i=0; i < 100; ++i)
+    for (j=0; j < 100; ++j)
+      {
+#pragma omp ordered depend(sink:k,i) /* { dg-error "is not an iteration" } */
+	bark();
+      }
+}
+
+void bar (int, int, int);
+
+void
+foo (int n, int m, int o)
+{
+  int i, j, k;
+  #pragma omp for collapse(2) ordered(2)
+  for (i = 0; i < m; i++)
+    {
+      for (j = 0; j < n; j++)
+        for (k = 0; k < o; k++)
+          {
+#pragma omp ordered depend(sink: i-1,j,k) depend(sink: i,j-1,k-1) depend(sink: i-1,j-1,k+1)
+	    bar (i, j, k);
+#pragma omp ordered depend(source)
+	  }
+    }
+}
+
+int
+baz ()
+{
+  int i, j;
+#pragma omp parallel for ordered(2)
+  for (i=0; i < 100; ++i)
+    for (j=0; j < 100; ++j)
+      {
+#pragma omp ordered depend(sink:i-1,j-3)
+        bar (i, j, 0);
+#pragma omp ordered depend(source)
+      }
+}
diff --git a/gcc/testsuite/g++.dg/gomp/sink-1.C b/gcc/testsuite/g++.dg/gomp/sink-1.C
new file mode 100644
index 0000000..d56d940
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/sink-1.C
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+
+void bar (int, int, int);
+
+template<typename T>
+void baz ()
+{
+  T i, j;
+#pragma omp parallel for ordered(2)
+  for (i=0; i < 100; ++i)
+    for (j=0; j < 100; ++j)
+      {
+#pragma omp ordered depend(sink:i-3,j)
+        bar (i, j, 0);
+      }
+}
+
+int main()
+{
+  baz<int>();
+}
diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c
index 06b4077..f4e8cf2 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -1479,7 +1479,9 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
 
 	case GIMPLE_OMP_ORDERED:
 	  s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
-	  copy = gimple_build_omp_ordered (s1);
+	  copy = gimple_build_omp_ordered
+	           (s1,
+		    gimple_omp_ordered_clauses (as_a <gomp_ordered *> (stmt)));
 	  break;
 
 	case GIMPLE_OMP_SECTION:
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 7f8e0fe..3c5d0f5 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -533,6 +533,22 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
 	case OMP_CLAUSE_DEPEND_SOURCE:
 	  pp_string (pp, "source)");
 	  return;
+	case OMP_CLAUSE_DEPEND_SINK:
+	  pp_string (pp, "sink:");
+	  for (tree t = OMP_CLAUSE_DECL (clause); t; t = TREE_CHAIN (t))
+	    if (TREE_CODE (t) == TREE_LIST)
+	      {
+		dump_generic_node (pp, TREE_VALUE (t), spc, flags, false);
+		if (TREE_PURPOSE (t) != integer_zero_node)
+		  dump_generic_node (pp, TREE_PURPOSE (t), spc, flags,
+				     false);
+		if (TREE_CHAIN (t))
+		  pp_comma (pp);
+	      }
+	    else
+	      gcc_unreachable ();
+	  pp_right_paren (pp);
+	  return;
 	default:
 	  gcc_unreachable ();
 	}
