[gomp4.1] depend(sink) and depend(source) parsing for C
diff mbox

Message ID 559EBC6C.70109@redhat.com
State New
Headers show

Commit Message

Aldy Hernandez July 9, 2015, 6:24 p.m. UTC
The following patch goes along with Jakub's parsing of ordered(n) loops. 
  With it, we can now parse his testcase, along with a variety of other 
tests with appropriate diagnostics.

The lowering to gimple is still not done, as we should agree on what 
needs to be emitted first.

I'll follow up with the C++ version once we agree on C.

How does this look?

Patch
diff mbox

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index cd3bd5a..98a6b02 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11701,6 +11701,92 @@  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");
+
+  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,
+			    "possible overflow in %<depend(sink)%> offset");
+	    }
+	  c_parser_consume_token (parser);
+
+	add_to_vector:
+	  vec = tree_cons (t, addend, 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 +11794,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 +11839,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/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..ba79977 100644
--- a/gcc/gimple-walk.c
+++ b/gcc/gimple-walk.c
@@ -327,6 +327,10 @@  walk_gimple_op (gimple stmt, walk_tree_fn callback_op,
       }
       break;
 
+    case GIMPLE_OMP_ORDERED:
+      /* Ignore clauses.  */
+      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..827f6a3 100644
--- a/gcc/gimple.def
+++ b/gcc/gimple.def
@@ -285,7 +285,7 @@  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)
+DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_omp_ordered", GSS_OMP_ORDERED)
 
 /* GIMPLE_OMP_PARALLEL <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
 
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 6057ea0..e33fe1e 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -527,6 +527,17 @@  struct GTY((tag("GSS_OMP_CRITICAL")))
   tree name;
 };
 
+/* GIMPLE_OMP_ORDERED */
+
+struct GTY((tag("GSS_OMP_ORDERED")))
+  gomp_ordered : public gimple_statement_omp
+{
+  /* [ WORD 1-7 ] : base class */
+
+  /* [ WORD 8 ]  */
+  tree clauses;
+};
+
 
 struct GTY(()) gimple_omp_for_iter {
   /* Condition code.  */
@@ -1007,6 +1018,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 +1234,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 +1382,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 +4487,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..254c37e 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -149,6 +149,9 @@  struct gimplify_omp_ctx
   struct gimplify_omp_ctx *outer_context;
   splay_tree variables;
   hash_set<tree> *privatized_types;
+  /* Iteration variables in an OMP_FOR.  */
+  tree *iter_vars;
+  int niter_vars;
   location_t location;
   enum omp_clause_default_kind default_kind;
   enum omp_region_type region_type;
@@ -384,6 +387,8 @@  delete_omp_context (struct gimplify_omp_ctx *c)
 {
   splay_tree_delete (c->variables);
   delete c->privatized_types;
+  if (c->iter_vars)
+    delete c->iter_vars;
   XDELETE (c);
 }
 
@@ -6343,6 +6348,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 +7294,9 @@  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->niter_vars = TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt));
+  gimplify_omp_ctxp->iter_vars
+    = new tree[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 +7305,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[i] = decl;
 
       /* Make sure the iteration variable is private.  */
       tree c = NULL_TREE;
@@ -8169,6 +8185,19 @@  gimplify_transaction (tree *expr_p, gimple_seq *pre_p)
   return GS_ALL_DONE;
 }
 
+/* Verify the validity of the depend(sink:...) variable VAR.
+   Return TRUE if everything is OK, otherwise return FALSE.  */
+
+static bool
+verify_sink_var (location_t loc, tree var)
+{
+  for (int i = 0; i < gimplify_omp_ctxp->niter_vars; ++i)
+    if (var == gimplify_omp_ctxp->iter_vars[i])
+      return true;
+  error_at (loc, "variable %qE is not an iteration variable", var);
+  return false;
+}
+
 /* Convert the GENERIC expression tree *EXPR_P to GIMPLE.  If the
    expression produces a value to be used as an operand inside a GIMPLE
    statement, the value will be stored back in *EXPR_P.  This value will
@@ -8982,7 +9011,32 @@  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)
+		      {
+			int n = 0;
+			bool fail = false;
+			for (tree decls = OMP_CLAUSE_DECL (c);
+			     decls && TREE_CODE (decls) == TREE_LIST;
+			     decls = TREE_CHAIN (decls))
+			  {
+			    if (!verify_sink_var (OMP_CLAUSE_LOCATION (c),
+						  TREE_PURPOSE (decls)))
+			      fail = true;
+			    ++n;
+			  }
+			/* Avoid being too redundant.  */
+			if (!fail && n != gimplify_omp_ctxp->niter_vars)
+			  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/gsstruct.def b/gcc/gsstruct.def
index 18cf403..00c07f2 100644
--- a/gcc/gsstruct.def
+++ b/gcc/gsstruct.def
@@ -42,6 +42,7 @@  DEFGSSTRUCT(GSS_EH_ELSE, geh_else, false)
 DEFGSSTRUCT(GSS_WCE, gimple_statement_wce, false)
 DEFGSSTRUCT(GSS_OMP, gimple_statement_omp, false)
 DEFGSSTRUCT(GSS_OMP_CRITICAL, gomp_critical, false)
+DEFGSSTRUCT(GSS_OMP_ORDERED, gomp_ordered, false)
 DEFGSSTRUCT(GSS_OMP_FOR, gomp_for, false)
 DEFGSSTRUCT(GSS_OMP_PARALLEL_LAYOUT, gimple_statement_omp_parallel_layout, false)
 DEFGSSTRUCT(GSS_OMP_TASK, gomp_task, false)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 83677ea..1f22b1b 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,51 @@  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))
+	  {
+	    error_at (OMP_CLAUSE_LOCATION (c),
+		      "depend(%s) is only available in 'omp ordered'",
+		      OMP_CLAUSE_DEPEND_KIND (c) == 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;
+	      /* Look for containing ordered(N) loop.  */
+	      for (omp_context *ctx_ = ctx; ctx_; ctx_ = ctx_->outer)
+		if (gimple_code (ctx_->stmt) == GIMPLE_OMP_FOR
+		    && find_omp_clause (gimple_omp_for_clauses (ctx_->stmt),
+					OMP_CLAUSE_ORDERED))
+		  {
+		    have_ordered = true;
+		    break;
+		  }
+	      if (!have_ordered)
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "depend clause is not within 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 +12608,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..353339f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/sink-1.c
@@ -0,0 +1,88 @@ 
+/* { 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 available in 'omp ordered'" } */
+  bark();
+
+#pragma omp ordered depend(source) /* { dg-error "depend clause is not within an ordered loop" } */
+
+#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 "is not an 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();
+      }
+}
+
+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/tree-inline.c b/gcc/tree-inline.c
index 06b4077..2ac4221 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -1479,7 +1479,7 @@  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, NULL);
 	  break;
 
 	case GIMPLE_OMP_SECTION:
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 7f8e0fe..27c3981 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -533,6 +533,9 @@  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");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}