@@ -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;
@@ -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))
@@ -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;
@@ -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);
@@ -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);
@@ -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);
@@ -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
@@ -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
@@ -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),
@@ -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)
@@ -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 ();
}
new file mode 100644
@@ -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)
+ }
+}
@@ -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:
@@ -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 ();
}