--- gcc/omp-expand.cc.jj	2022-09-06 09:19:14.735561333 +0200
+++ gcc/omp-expand.cc	2022-09-07 21:08:18.649782291 +0200
@@ -3287,7 +3287,8 @@ expand_omp_ordered_source (gimple_stmt_i
 
 static void
 expand_omp_ordered_sink (gimple_stmt_iterator *gsi, struct omp_for_data *fd,
-			 tree *counts, tree c, location_t loc)
+			 tree *counts, tree c, location_t loc,
+			 basic_block cont_bb)
 {
   auto_vec<tree, 10> args;
   enum built_in_function sink_ix
@@ -3300,7 +3301,93 @@ expand_omp_ordered_sink (gimple_stmt_ite
 
   if (deps == NULL)
     {
-      sorry_at (loc, "%<doacross(sink:omp_cur_iteration-1)%> not supported yet");
+      /* Handle doacross(sink: omp_cur_iteration - 1).  */
+      gsi_prev (&gsi2);
+      edge e1 = split_block (gsi_bb (gsi2), gsi_stmt (gsi2));
+      edge e2 = split_block_after_labels (e1->dest);
+      gsi2 = gsi_after_labels (e1->dest);
+      *gsi = gsi_last_bb (e1->src);
+      gimple_stmt_iterator gsi3 = *gsi;
+
+      if (counts[fd->collapse - 1])
+	{
+	  gcc_assert (fd->collapse == 1);
+	  t = counts[fd->collapse - 1];
+	}
+      else if (fd->collapse > 1)
+	t = fd->loop.v;
+      else
+	{
+	  t = fold_build2 (MINUS_EXPR, TREE_TYPE (fd->loops[0].v),
+			   fd->loops[0].v, fd->loops[0].n1);
+	  t = fold_convert (fd->iter_type, t);
+	}
+
+      t = force_gimple_operand_gsi (gsi, t, true, NULL_TREE,
+				    false, GSI_CONTINUE_LINKING);
+      gsi_insert_after (gsi, gimple_build_cond (NE_EXPR, t,
+						build_zero_cst (TREE_TYPE (t)),
+						NULL_TREE, NULL_TREE),
+			GSI_NEW_STMT);
+
+      t = fold_build2 (PLUS_EXPR, TREE_TYPE (t), t,
+		       build_minus_one_cst (TREE_TYPE (t)));
+      t = force_gimple_operand_gsi (&gsi2, t, true, NULL_TREE,
+				    true, GSI_SAME_STMT);
+      args.safe_push (t);
+      for (i = fd->collapse; i < fd->ordered; i++)
+	{
+	  t = counts[fd->ordered + 2 + (i - fd->collapse)];
+	  t = fold_build2 (PLUS_EXPR, TREE_TYPE (t), t,
+			   build_minus_one_cst (TREE_TYPE (t)));
+	  t = fold_convert (fd->iter_type, t);
+	  t = force_gimple_operand_gsi (&gsi2, t, true, NULL_TREE,
+					true, GSI_SAME_STMT);
+	  args.safe_push (t);
+	}
+
+      gimple *g = gimple_build_call_vec (builtin_decl_explicit (sink_ix),
+					 args);
+      gimple_set_location (g, loc);
+      gsi_insert_before (&gsi2, g, GSI_SAME_STMT);
+
+      edge e3 = make_edge (e1->src, e2->dest, EDGE_FALSE_VALUE);
+      e3->probability = profile_probability::guessed_always () / 8;
+      e1->probability = e3->probability.invert ();
+      e1->flags = EDGE_TRUE_VALUE;
+      set_immediate_dominator (CDI_DOMINATORS, e2->dest, e1->src);
+
+      if (fd->ordered > fd->collapse && cont_bb)
+	{
+	  if (counts[fd->ordered + 1] == NULL_TREE)
+	    counts[fd->ordered + 1]
+	      = create_tmp_var (boolean_type_node, ".first");
+
+	  edge e4;
+	  if (gsi_end_p (gsi3))
+	    e4 = split_block_after_labels (e1->src);
+	  else
+	    {
+	      gsi_prev (&gsi3);
+	      e4 = split_block (gsi_bb (gsi3), gsi_stmt (gsi3));
+	    }
+	  gsi3 = gsi_last_bb (e4->src);
+
+	  gsi_insert_after (&gsi3,
+			    gimple_build_cond (NE_EXPR,
+					       counts[fd->ordered + 1],
+					       boolean_false_node,
+					       NULL_TREE, NULL_TREE),
+			    GSI_NEW_STMT);
+
+	  edge e5 = make_edge (e4->src, e2->dest, EDGE_FALSE_VALUE);
+	  e4->probability = profile_probability::guessed_always () / 8;
+	  e5->probability = e4->probability.invert ();
+	  e4->flags = EDGE_TRUE_VALUE;
+	  set_immediate_dominator (CDI_DOMINATORS, e2->dest, e4->src);
+	}
+
+      *gsi = gsi_after_labels (e2->dest);
       return;
     }
   for (i = 0; i < fd->ordered; i++)
@@ -3558,6 +3645,7 @@ expand_omp_ordered_source_sink (struct o
     = build_array_type_nelts (fd->iter_type, fd->ordered - fd->collapse + 1);
   counts[fd->ordered] = create_tmp_var (atype, ".orditera");
   TREE_ADDRESSABLE (counts[fd->ordered]) = 1;
+  counts[fd->ordered + 1] = NULL_TREE;
 
   for (inner = region->inner; inner; inner = inner->next)
     if (inner->type == GIMPLE_OMP_ORDERED)
@@ -3575,7 +3663,7 @@ expand_omp_ordered_source_sink (struct o
 	for (c = gimple_omp_ordered_clauses (ord_stmt);
 	     c; c = OMP_CLAUSE_CHAIN (c))
 	  if (OMP_CLAUSE_DOACROSS_KIND (c) == OMP_CLAUSE_DOACROSS_SINK)
-	    expand_omp_ordered_sink (&gsi, fd, counts, c, loc);
+	    expand_omp_ordered_sink (&gsi, fd, counts, c, loc, cont_bb);
 	gsi_remove (&gsi, true);
       }
 }
@@ -3611,6 +3699,9 @@ expand_omp_for_ordered_loops (struct omp
     {
       tree t, type = TREE_TYPE (fd->loops[i].v);
       gimple_stmt_iterator gsi = gsi_after_labels (body_bb);
+      if (counts[fd->ordered + 1] && i == fd->collapse)
+	expand_omp_build_assign (&gsi, counts[fd->ordered + 1],
+				 boolean_true_node);
       expand_omp_build_assign (&gsi, fd->loops[i].v,
 			       fold_convert (type, fd->loops[i].n1));
       if (counts[i])
@@ -3658,6 +3749,9 @@ expand_omp_for_ordered_loops (struct omp
 			 size_int (i - fd->collapse + 1),
 			 NULL_TREE, NULL_TREE);
 	  expand_omp_build_assign (&gsi, aref, t);
+	  if (counts[fd->ordered + 1] && i == fd->ordered - 1)
+	    expand_omp_build_assign (&gsi, counts[fd->ordered + 1],
+				     boolean_false_node);
 	  gsi_prev (&gsi);
 	  e2 = split_block (cont_bb, gsi_stmt (gsi));
 	  new_header = e2->dest;
@@ -3915,7 +4009,10 @@ expand_omp_for_generic (struct omp_regio
       int first_zero_iter1 = -1, first_zero_iter2 = -1;
       basic_block zero_iter1_bb = NULL, zero_iter2_bb = NULL, l2_dom_bb = NULL;
 
-      counts = XALLOCAVEC (tree, fd->ordered ? fd->ordered + 1 : fd->collapse);
+      counts = XALLOCAVEC (tree, fd->ordered
+				 ? fd->ordered + 2
+				   + (fd->ordered - fd->collapse)
+				 : fd->collapse);
       expand_omp_for_init_counts (fd, &gsi, entry_bb, counts,
 				  zero_iter1_bb, first_zero_iter1,
 				  zero_iter2_bb, first_zero_iter2, l2_dom_bb);
@@ -4352,13 +4449,21 @@ expand_omp_for_generic (struct omp_regio
   if (fd->ordered)
     {
       /* Until now, counts array contained number of iterations or
-	 variable containing it for ith loop.  From now on, we need
+	 variable containing it for ith loop.  From now on, we usually need
 	 those counts only for collapsed loops, and only for the 2nd
 	 till the last collapsed one.  Move those one element earlier,
 	 we'll use counts[fd->collapse - 1] for the first source/sink
 	 iteration counter and so on and counts[fd->ordered]
 	 as the array holding the current counter values for
-	 depend(source).  */
+	 depend(source).  For doacross(sink:omp_cur_iteration - 1) we need
+	 the counts from fd->collapse to fd->ordered - 1; make a copy of
+	 those to counts[fd->ordered + 2] and onwards.
+	 counts[fd->ordered + 1] can be a flag whether it is the first
+	 iteration with a new collapsed counter (used only if
+	 fd->ordered > fd->collapse).  */
+      if (fd->ordered > fd->collapse)
+	memcpy (counts + fd->ordered + 2, counts + fd->collapse,
+		(fd->ordered - fd->collapse) * sizeof (counts[0]));
       if (fd->collapse > 1)
 	memmove (counts, counts + 1, (fd->collapse - 1) * sizeof (counts[0]));
       if (broken_loop)
--- gcc/testsuite/c-c++-common/gomp/doacross-7.c.jj	2022-09-08 11:01:50.295306390 +0200
+++ gcc/testsuite/c-c++-common/gomp/doacross-7.c	2022-09-08 10:33:37.589881715 +0200
@@ -0,0 +1,78 @@
+void
+foo (int l)
+{
+  int i, j, k;
+  #pragma omp parallel
+  {
+    #pragma omp for schedule(static) ordered (3)
+    for (i = 2; i < 256 / 16 - 1; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	    #pragma omp ordered doacross(source:)
+	  }
+    #pragma omp for schedule(static) ordered (3) collapse(2)
+    for (i = 2; i < 256 / 16 - 1; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	    #pragma omp ordered doacross(source:)
+	  }
+    #pragma omp for schedule(static) ordered (3) collapse(3)
+    for (i = 2; i < 256 / 16 - 1; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	    #pragma omp ordered doacross(source: omp_cur_iteration)
+	  }
+    #pragma omp for schedule(static) ordered (1) nowait
+    for (i = 2; i < 256 / 16 - 1; i += l)
+      {
+	#pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	#pragma omp ordered doacross(source:)
+      }
+  }
+}
+
+void
+bar (int l, int m, int n, int o)
+{
+  int i, j, k;
+  #pragma omp for schedule(static) ordered (3)
+  for (i = 2; i < 256 / 16 - 1; i++)
+    for (j = 0; j < m; j += n)
+      for (k = o; k <= 3; k++)
+	{
+	  foo (l);
+	  #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	  #pragma omp ordered doacross(source:omp_cur_iteration)
+	}
+  #pragma omp for schedule(static) ordered (3) collapse(2)
+  for (i = 2; i < 256 / 16 - m; i += n)
+    for (j = 0; j < 8; j += o)
+      for (k = 1; k <= 3; k++)
+	{
+	  foo (l);
+	  #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	  #pragma omp ordered doacross(source : omp_cur_iteration)
+	}
+  #pragma omp for schedule(static) ordered (3) collapse(3)
+  for (i = m; i < 256 / 16 - 1; i++)
+    for (j = 0; j < n; j += 2)
+      for (k = 1; k <= o; k++)
+	{
+	  foo (l);
+	  #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	  #pragma omp ordered doacross(source :)
+	}
+  #pragma omp for schedule(static) ordered
+  for (i = m; i < n / 16 - 1; i += l)
+    {
+      foo (l);
+      #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+      #pragma omp ordered doacross(source: omp_cur_iteration)
+    }
+}
--- libgomp/libgomp.texi.jj	2022-09-05 23:25:28.671018459 +0200
+++ libgomp/libgomp.texi	2022-09-08 12:49:17.575393448 +0200
@@ -397,8 +397,7 @@ to address of matching mapped list item
       @code{source}/@code{sink} modifier @tab Y @tab
 @item Deprecation of @code{depend} with @code{source}/@code{sink} modifier
       @tab N @tab
-@item @code{omp_cur_iteration} keyword @tab P
-      @tab @code{sink: omp_cur_iteration - 1} unsupported
+@item @code{omp_cur_iteration} keyword @tab Y @tab
 @end multitable
 
 @unnumberedsubsec Other new OpenMP 5.2 features
--- libgomp/testsuite/libgomp.c/doacross-4.c.jj	2022-09-08 11:13:20.274112539 +0200
+++ libgomp/testsuite/libgomp.c/doacross-4.c	2022-09-08 11:51:27.187644498 +0200
@@ -0,0 +1,228 @@
+extern void abort (void);
+
+#define N 256
+int a[N], b[N / 16][8][4], c[N / 32][8][8], g[N / 16][8][6];
+volatile int d, e;
+volatile unsigned long long f;
+
+int
+main ()
+{
+  unsigned long long i;
+  int j, k, l, m;
+  #pragma omp parallel private (l)
+  {
+    #pragma omp for schedule(static, 1) ordered nowait
+    for (i = 1; i < N + f; i++)
+      {
+	#pragma omp atomic write
+	a[i] = 1;
+	#pragma omp ordered doacross(sink: i - 1)
+	if (i > 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i - 1];
+	    if (l < 2)
+	      abort ();
+	  }
+	#pragma omp atomic write
+	a[i] = 2;
+	if (i < N - 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i + 1];
+	    if (l == 3)
+	      abort ();
+	  }
+	#pragma omp ordered doacross(source : omp_cur_iteration)
+	#pragma omp atomic write
+	a[i] = 3;
+      }
+    #pragma omp for schedule(static) ordered (3) nowait
+    for (i = 3; i < N / 16 - 1 + f; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp atomic write
+	    b[i][j][k] = 1;
+	    #pragma omp ordered doacross(sink: i, j - 2, k - 1) \
+				doacross(sink: i - 2, j - 2, k + 1)
+	    #pragma omp ordered doacross(sink: i - 3, j + 2, k - 2)
+	    if (j >= 2 && k > 1)
+	      {
+		#pragma omp atomic read
+		l = b[i][j - 2][k - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    b[i][j][k] = 2;
+	    if (i >= 5 && j >= 2 && k < 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 2][j - 2][k + 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (i >= 6 && j < N / 16 - 3 && k == 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 3][j + 2][k - 2];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered doacross(source : )
+	    #pragma omp atomic write
+	    b[i][j][k] = 3;
+	  }
+#define A(n) int n;
+#define B(n) A(n##0) A(n##1) A(n##2) A(n##3)
+#define C(n) B(n##0) B(n##1) B(n##2) B(n##3)
+#define D(n) C(n##0) C(n##1) C(n##2) C(n##3)
+    D(m)
+#undef A
+    #pragma omp for collapse (2) ordered(61) schedule(dynamic, 15)
+    for (i = 2; i < N / 32 + f; i++)
+      for (j = 7; j > 1; j--)
+	for (k = 6; k >= 0; k -= 2)
+#define A(n) for (n = 4; n < 5; n++)
+	  D(m)
+#undef A
+	    {
+	      #pragma omp atomic write
+	      c[i][j][k] = 1;
+#define A(n) ,n
+#define E(n) C(n##0) C(n##1) C(n##2) B(n##30) B(n##31) A(n##320) A(n##321)
+	      #pragma omp ordered doacross (sink: i, j, k + 2 E(m)) \
+				  doacross (sink:i - 2, j + 1, k - 4 E(m)) \
+				  doacross(sink: i - 1, j - 2, k - 2 E(m))
+	      if (k <= 4)
+		{
+		  #pragma omp atomic read
+		  l = c[i][j][k + 2];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp atomic write
+	      c[i][j][k] = 2;
+	      if (i >= 4 && j < 7 && k >= 4)
+		{
+		  #pragma omp atomic read
+		  l = c[i - 2][j + 1][k - 4];
+		  if (l < 2)
+		    abort ();
+		}
+	      if (i >= 3 && j >= 4 && k >= 2)
+		{
+		  #pragma omp atomic read
+		  l = c[i - 1][j - 2][k - 2];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp atomic write
+	      c[i][j][k] = 3;
+	    }
+    #pragma omp for schedule(static) ordered (3) nowait
+    for (j = 0; j < N / 16 - 1; j++)
+      for (k = 0; k < 8; k += 2)
+	for (i = 3; i <= 5 + f; i++)
+	  {
+	    #pragma omp atomic write
+	    g[j][k][i] = 1;
+	    #pragma omp ordered doacross(sink: j, k - 2, i - 1) \
+				doacross(sink: j - 2, k - 2, i + 1)
+	    #pragma omp ordered doacross(sink: j - 3, k + 2, i - 2)
+	    if (k >= 2 && i > 3)
+	      {
+		#pragma omp atomic read
+		l = g[j][k - 2][i - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    g[j][k][i] = 2;
+	    if (j >= 2 && k >= 2 && i < 5)
+	      {
+		#pragma omp atomic read
+		l = g[j - 2][k - 2][i + 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (j >= 3 && k < N / 16 - 3 && i == 5)
+	      {
+		#pragma omp atomic read
+		l = g[j - 3][k + 2][i - 2];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered doacross(source :)
+	    #pragma omp atomic write
+	    g[j][k][i] = 3;
+	  }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k)
+    for (i = 2; i < f + 3; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d; k++)
+	  for (l = 0; l < d + 2; l++)
+	    {
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp ordered doacross (sink:i - 2, j + 2, k - 2, l)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp single
+    {
+      if (i != 3 || j != -1 || k != 0)
+	abort ();
+      i = 8; j = 9; k = 10;
+    }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k, m)
+    for (i = 2; i < f + 3; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (m = 0; m < d; m++)
+	    {
+	      #pragma omp ordered doacross (source:)
+	      #pragma omp ordered doacross (sink:i - 2, j + 2, k - 2, m)
+	      abort ();
+	    }
+    #pragma omp single
+    if (i != 3 || j != -1 || k != 2 || m != 0)
+      abort ();
+    #pragma omp for collapse(2) ordered(4) nowait
+    for (i = 2; i < f + 3; i++)
+      for (j = d; j > 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (l = 0; l < d + 4; l++)
+	    {
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp ordered doacross (sink:i - 2, j + 2, k - 2, l)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp for nowait
+    for (i = 0; i < N; i++)
+      if (a[i] != 3 * (i >= 1))
+	abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 4; k++)
+	  if (b[i][j][k] != 3 * (i >= 3 && i < N / 16 - 1 && (j & 1) == 0 && k >= 1))
+	    abort ();
+    #pragma omp for collapse(3) nowait
+    for (i = 0; i < N / 32; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 8; k++)
+	  if (c[i][j][k] != 3 * (i >= 2 && j >= 2 && (k & 1) == 0))
+	    abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 6; k++)
+	  if (g[i][j][k] != 3 * (i < N / 16 - 1 && (j & 1) == 0 && k >= 3))
+	    abort ();
+  }
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/doacross-5.c.jj	2022-09-08 11:24:26.787237389 +0200
+++ libgomp/testsuite/libgomp.c/doacross-5.c	2022-09-08 11:47:42.248642381 +0200
@@ -0,0 +1,198 @@
+extern void abort (void);
+
+#define N 256
+int a[N], b[N / 16][8][4], c[N / 32][8][8];
+volatile int d, e;
+
+int
+main ()
+{
+  int i, j, k, l, m;
+  #pragma omp parallel private (l)
+  {
+    #pragma omp for schedule(static, 1) ordered (1) nowait
+    for (i = 0; i < N; i++)
+      {
+	#pragma omp atomic write
+	a[i] = 1;
+	#pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	if (i)
+	  {
+	    #pragma omp atomic read
+	    l = a[i - 1];
+	    if (l < 2)
+	      abort ();
+	  }
+	#pragma omp atomic write
+	a[i] = 2;
+	if (i < N - 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i + 1];
+	    if (l == 3)
+	      abort ();
+	  }
+	#pragma omp ordered doacross(source :)
+	#pragma omp atomic write
+	a[i] = 3;
+      }
+    #pragma omp for schedule(static) ordered (3) nowait
+    for (i = 2; i < N / 16 - 1; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp atomic write
+	    b[i][j][k] = 1;
+	    #pragma omp ordered doacross(sink: omp_cur_iteration - 1) \
+				doacross(sink: i - 2, j - 2, k + 1)
+	    #pragma omp ordered doacross(sink: i - 3, j + 2, k - 2)
+	    if (i != 2 || j || k != 1)
+	      {
+		if (k != 1)
+		  #pragma omp atomic read
+		  l = b[i][j][k - 1];
+		else if (j)
+		  #pragma omp atomic read
+		  l = b[i][j - 2][3];
+		else
+		  #pragma omp atomic read
+		  l = b[i - 1][6][3];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    b[i][j][k] = 2;
+	    if (i >= 4 && j >= 2 && k < 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 2][j - 2][k + 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (i >= 5 && j < N / 16 - 3 && k == 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 3][j + 2][k - 2];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered doacross(source : omp_cur_iteration)
+	    #pragma omp atomic write
+	    b[i][j][k] = 3;
+	  }
+#define A(n) int n;
+#define B(n) A(n##0) A(n##1) A(n##2) A(n##3)
+#define C(n) B(n##0) B(n##1) B(n##2) B(n##3)
+#define D(n) C(n##0) C(n##1) C(n##2) C(n##3)
+    D(m)
+#undef A
+    #pragma omp for collapse (2) ordered(61) schedule(dynamic, 15)
+    for (i = 0; i < N / 32; i++)
+      for (j = 7; j > 1; j--)
+	for (k = 6; k >= 0; k -= 2)
+#define A(n) for (n = 4; n < 5; n++)
+	  D(m)
+#undef A
+	    {
+	      #pragma omp atomic write
+	      c[i][j][k] = 1;
+#define A(n) ,n
+#define E(n) C(n##0) C(n##1) C(n##2) B(n##30) B(n##31) A(n##320) A(n##321)
+	      #pragma omp ordered doacross (sink: i, j, k + 2 E(m)) \
+				  doacross (sink:omp_cur_iteration - 1) \
+				  doacross(sink: i - 1, j - 2, k - 2 E(m))
+	      if (k <= 4)
+		{
+		  #pragma omp atomic read
+		  l = c[i][j][k + 2];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp atomic write
+	      c[i][j][k] = 2;
+	      if (i || j != 7 && k != 6)
+		{
+		  if (k != 6)
+		    #pragma omp atomic read
+		    l = c[i][j][k + 2];
+		  else if (j != 7)
+		    #pragma omp atomic read
+		    l = c[i][j + 1][0];
+		  else
+		    #pragma omp atomic read
+		    l = c[i - 1][2][0];
+		  if (l < 2)
+		    abort ();
+		}
+	      if (i >= 1 && j >= 4 && k >= 2)
+		{
+		  #pragma omp atomic read
+		  l = c[i - 1][j - 2][k - 2];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp ordered doacross (source: )
+	      #pragma omp atomic write
+	      c[i][j][k] = 3;
+	    }
+
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k)
+    for (i = 0; i < d + 1; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d; k++)
+	  for (l = 0; l < d + 2; l++)
+	    {
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp ordered doacross (sink: omp_cur_iteration - 1)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp single
+    {
+      if (i != 1 || j != -1 || k != 0)
+	abort ();
+      i = 8; j = 9; k = 10;
+    }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k, m)
+    for (i = 0; i < d + 1; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (m = 0; m < d; m++)
+	    {
+	      #pragma omp ordered doacross (source : )
+	      #pragma omp ordered doacross (sink:omp_cur_iteration - 1)
+	      abort ();
+	    }
+    #pragma omp single
+    if (i != 1 || j != -1 || k != 2 || m != 0)
+      abort ();
+    #pragma omp for collapse(2) ordered(4) nowait
+    for (i = 0; i < d + 1; i++)
+      for (j = d; j > 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (l = 0; l < d + 4; l++)
+	    {
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp ordered doacross (sink:omp_cur_iteration - 1)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp for nowait
+    for (i = 0; i < N; i++)
+      if (a[i] != 3)
+	abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 4; k++)
+	  if (b[i][j][k] != 3 * (i >= 2 && i < N / 16 - 1 && (j & 1) == 0 && k >= 1))
+	    abort ();
+    #pragma omp for collapse(3) nowait
+    for (i = 0; i < N / 32; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 8; k++)
+	  if (c[i][j][k] != 3 * (j >= 2 && (k & 1) == 0))
+	    abort ();
+  }
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/doacross-6.c.jj	2022-09-08 12:12:09.894080951 +0200
+++ libgomp/testsuite/libgomp.c/doacross-6.c	2022-09-08 12:36:18.902768322 +0200
@@ -0,0 +1,231 @@
+extern void abort (void);
+
+#define N 256
+int a[N], b[N / 16][8][4], c[N / 32][8][8], g[N / 16][8][6];
+volatile int d, e;
+volatile unsigned long long f;
+
+int
+main ()
+{
+  unsigned long long i;
+  int j, k, l, m;
+  #pragma omp parallel private (l)
+  {
+    #pragma omp for schedule(static, 1) ordered (1) nowait
+    for (i = 1; i < N + f; i++)
+      {
+	#pragma omp atomic write
+	a[i] = 1;
+	#pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	if (i > 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i - 1];
+	    if (l < 2)
+	      abort ();
+	  }
+	#pragma omp atomic write
+	a[i] = 2;
+	if (i < N - 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i + 1];
+	    if (l == 3)
+	      abort ();
+	  }
+	#pragma omp ordered doacross(source : omp_cur_iteration)
+	#pragma omp atomic write
+	a[i] = 3;
+      }
+    #pragma omp for schedule(static) ordered (3) nowait
+    for (i = 3; i < N / 16 - 1 + f; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp atomic write
+	    b[i][j][k] = 1;
+	    #pragma omp ordered doacross(sink: i, j - 2, k - 1) \
+				doacross(sink: i - 2, j - 2, k + 1)
+	    #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	    if (j >= 2 && k > 1)
+	      {
+		#pragma omp atomic read
+		l = b[i][j - 2][k - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    b[i][j][k] = 2;
+	    if (i >= 5 && j >= 2 && k < 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 2][j - 2][k + 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (i != 3 || j || k != 1)
+	      {
+		if (k != 1)
+		  #pragma omp atomic read
+		  l = b[i][j][k - 1];
+		else if (j)
+		  #pragma omp atomic read
+		  l = b[i][j - 2][3];
+		else
+		  #pragma omp atomic read
+		  l = b[i - 1][6][3];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered doacross(source:)
+	    #pragma omp atomic write
+	    b[i][j][k] = 3;
+	  }
+#define A(n) int n;
+#define B(n) A(n##0) A(n##1) A(n##2) A(n##3)
+#define C(n) B(n##0) B(n##1) B(n##2) B(n##3)
+#define D(n) C(n##0) C(n##1) C(n##2) C(n##3)
+    D(m)
+#undef A
+    #pragma omp for collapse (2) ordered(61) schedule(dynamic, 15)
+    for (i = 2; i < N / 32 + f; i++)
+      for (j = 7; j > 1; j--)
+	for (k = 6; k >= 0; k -= 2)
+#define A(n) for (n = 4; n < 5; n++)
+	  D(m)
+#undef A
+	    {
+	      #pragma omp atomic write
+	      c[i][j][k] = 1;
+	      #pragma omp ordered doacross (sink: omp_cur_iteration - 1)
+	      if (i != 2 || j != 7 || k != 6)
+		{
+		  if (k != 6)
+		    #pragma omp atomic read
+		    l = c[i][j][k + 2];
+		  else if (j != 7)
+		    #pragma omp atomic read
+		    l = c[i][j + 1][0];
+		  else
+		    #pragma omp atomic read
+		    l = c[i - 1][2][0];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp atomic write
+	      c[i][j][k] = 2;
+	      #pragma omp ordered doacross (source:)
+	      #pragma omp atomic write
+	      c[i][j][k] = 3;
+	    }
+    #pragma omp for schedule(static) ordered (3) nowait
+    for (j = 0; j < N / 16 - 1; j++)
+      for (k = 0; k < 8; k += 2)
+	for (i = 3; i <= 5 + f; i++)
+	  {
+	    #pragma omp atomic write
+	    g[j][k][i] = 1;
+	    #pragma omp ordered doacross(sink: j, k - 2, i - 1) \
+				doacross(sink: omp_cur_iteration - 1)
+	    #pragma omp ordered doacross(sink: j - 3, k + 2, i - 2)
+	    if (k >= 2 && i > 3)
+	      {
+		#pragma omp atomic read
+		l = g[j][k - 2][i - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    g[j][k][i] = 2;
+	    if (j || k || i != 3)
+	      {
+		if (i != 3)
+		  #pragma omp atomic read
+		  l = g[j][k][i - 1];
+		else if (k)
+		  #pragma omp atomic read
+		  l = g[j][k - 2][5 + f];
+		else
+		  #pragma omp atomic read
+		  l = g[j - 1][6][5 + f];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (j >= 3 && k < N / 16 - 3 && i == 5)
+	      {
+		#pragma omp atomic read
+		l = g[j - 3][k + 2][i - 2];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered doacross(source : omp_cur_iteration)
+	    #pragma omp atomic write
+	    g[j][k][i] = 3;
+	  }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k)
+    for (i = 2; i < f + 3; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d; k++)
+	  for (l = 0; l < d + 2; l++)
+	    {
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp ordered doacross (sink:omp_cur_iteration - 1)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp single
+    {
+      if (i != 3 || j != -1 || k != 0)
+	abort ();
+      i = 8; j = 9; k = 10;
+    }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k, m)
+    for (i = 2; i < f + 3; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (m = 0; m < d; m++)
+	    {
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp ordered doacross (sink:omp_cur_iteration - 1)
+	      abort ();
+	    }
+    #pragma omp single
+    if (i != 3 || j != -1 || k != 2 || m != 0)
+      abort ();
+    #pragma omp for collapse(2) ordered(4) nowait
+    for (i = 2; i < f + 3; i++)
+      for (j = d; j > 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (l = 0; l < d + 4; l++)
+	    {
+	      #pragma omp ordered doacross (source:)
+	      #pragma omp ordered doacross (sink:omp_cur_iteration-1)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp for nowait
+    for (i = 0; i < N; i++)
+      if (a[i] != 3 * (i >= 1))
+	abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 4; k++)
+	  if (b[i][j][k] != 3 * (i >= 3 && i < N / 16 - 1 && (j & 1) == 0 && k >= 1))
+	    abort ();
+    #pragma omp for collapse(3) nowait
+    for (i = 0; i < N / 32; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 8; k++)
+	  if (c[i][j][k] != 3 * (i >= 2 && j >= 2 && (k & 1) == 0))
+	    abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 6; k++)
+	  if (g[i][j][k] != 3 * (i < N / 16 - 1 && (j & 1) == 0 && k >= 3))
+	    abort ();
+  }
+  return 0;
+}
--- libgomp/testsuite/libgomp.c/doacross-7.c.jj	2022-09-08 12:37:05.964141298 +0200
+++ libgomp/testsuite/libgomp.c/doacross-7.c	2022-09-08 12:37:52.923515613 +0200
@@ -0,0 +1,231 @@
+extern void abort (void);
+
+#define N 256
+int a[N], b[N / 16][8][4], c[N / 32][8][8], g[N / 16][8][6];
+volatile int d, e;
+volatile unsigned long long f;
+
+int
+main ()
+{
+  unsigned long long i;
+  int j, k, l, m;
+  #pragma omp parallel private (l)
+  {
+    #pragma omp for schedule(guided, 3) ordered (1) nowait
+    for (i = 1; i < N + f; i++)
+      {
+	#pragma omp atomic write
+	a[i] = 1;
+	#pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	if (i > 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i - 1];
+	    if (l < 2)
+	      abort ();
+	  }
+	#pragma omp atomic write
+	a[i] = 2;
+	if (i < N - 1)
+	  {
+	    #pragma omp atomic read
+	    l = a[i + 1];
+	    if (l == 3)
+	      abort ();
+	  }
+	#pragma omp ordered doacross(source : omp_cur_iteration)
+	#pragma omp atomic write
+	a[i] = 3;
+      }
+    #pragma omp for schedule(guided) ordered (3) nowait
+    for (i = 3; i < N / 16 - 1 + f; i++)
+      for (j = 0; j < 8; j += 2)
+	for (k = 1; k <= 3; k++)
+	  {
+	    #pragma omp atomic write
+	    b[i][j][k] = 1;
+	    #pragma omp ordered doacross(sink: i, j - 2, k - 1) \
+				doacross(sink: i - 2, j - 2, k + 1)
+	    #pragma omp ordered doacross(sink: omp_cur_iteration - 1)
+	    if (j >= 2 && k > 1)
+	      {
+		#pragma omp atomic read
+		l = b[i][j - 2][k - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    b[i][j][k] = 2;
+	    if (i >= 5 && j >= 2 && k < 3)
+	      {
+		#pragma omp atomic read
+		l = b[i - 2][j - 2][k + 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (i != 3 || j || k != 1)
+	      {
+		if (k != 1)
+		  #pragma omp atomic read
+		  l = b[i][j][k - 1];
+		else if (j)
+		  #pragma omp atomic read
+		  l = b[i][j - 2][3];
+		else
+		  #pragma omp atomic read
+		  l = b[i - 1][6][3];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered doacross(source:)
+	    #pragma omp atomic write
+	    b[i][j][k] = 3;
+	  }
+#define A(n) int n;
+#define B(n) A(n##0) A(n##1) A(n##2) A(n##3)
+#define C(n) B(n##0) B(n##1) B(n##2) B(n##3)
+#define D(n) C(n##0) C(n##1) C(n##2) C(n##3)
+    D(m)
+#undef A
+    #pragma omp for collapse (2) ordered(61) schedule(guided, 15)
+    for (i = 2; i < N / 32 + f; i++)
+      for (j = 7; j > 1; j--)
+	for (k = 6; k >= 0; k -= 2)
+#define A(n) for (n = 4; n < 5; n++)
+	  D(m)
+#undef A
+	    {
+	      #pragma omp atomic write
+	      c[i][j][k] = 1;
+	      #pragma omp ordered doacross (sink: omp_cur_iteration - 1)
+	      if (i != 2 || j != 7 || k != 6)
+		{
+		  if (k != 6)
+		    #pragma omp atomic read
+		    l = c[i][j][k + 2];
+		  else if (j != 7)
+		    #pragma omp atomic read
+		    l = c[i][j + 1][0];
+		  else
+		    #pragma omp atomic read
+		    l = c[i - 1][2][0];
+		  if (l < 2)
+		    abort ();
+		}
+	      #pragma omp atomic write
+	      c[i][j][k] = 2;
+	      #pragma omp ordered doacross (source:)
+	      #pragma omp atomic write
+	      c[i][j][k] = 3;
+	    }
+    #pragma omp for schedule(guided, 5) ordered (3) nowait
+    for (j = 0; j < N / 16 - 1; j++)
+      for (k = 0; k < 8; k += 2)
+	for (i = 3; i <= 5 + f; i++)
+	  {
+	    #pragma omp atomic write
+	    g[j][k][i] = 1;
+	    #pragma omp ordered doacross(sink: j, k - 2, i - 1) \
+				doacross(sink: omp_cur_iteration - 1)
+	    #pragma omp ordered doacross(sink: j - 3, k + 2, i - 2)
+	    if (k >= 2 && i > 3)
+	      {
+		#pragma omp atomic read
+		l = g[j][k - 2][i - 1];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp atomic write
+	    g[j][k][i] = 2;
+	    if (j || k || i != 3)
+	      {
+		if (i != 3)
+		  #pragma omp atomic read
+		  l = g[j][k][i - 1];
+		else if (k)
+		  #pragma omp atomic read
+		  l = g[j][k - 2][5 + f];
+		else
+		  #pragma omp atomic read
+		  l = g[j - 1][6][5 + f];
+		if (l < 2)
+		  abort ();
+	      }
+	    if (j >= 3 && k < N / 16 - 3 && i == 5)
+	      {
+		#pragma omp atomic read
+		l = g[j - 3][k + 2][i - 2];
+		if (l < 2)
+		  abort ();
+	      }
+	    #pragma omp ordered doacross(source : omp_cur_iteration)
+	    #pragma omp atomic write
+	    g[j][k][i] = 3;
+	  }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k)
+    for (i = 2; i < f + 3; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d; k++)
+	  for (l = 0; l < d + 2; l++)
+	    {
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp ordered doacross (sink:omp_cur_iteration - 1)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp single
+    {
+      if (i != 3 || j != -1 || k != 0)
+	abort ();
+      i = 8; j = 9; k = 10;
+    }
+    #pragma omp for collapse(2) ordered(4) lastprivate (i, j, k, m)
+    for (i = 2; i < f + 3; i++)
+      for (j = d + 1; j >= 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (m = 0; m < d; m++)
+	    {
+	      #pragma omp ordered doacross (source : omp_cur_iteration)
+	      #pragma omp ordered doacross (sink:omp_cur_iteration - 1)
+	      abort ();
+	    }
+    #pragma omp single
+    if (i != 3 || j != -1 || k != 2 || m != 0)
+      abort ();
+    #pragma omp for collapse(2) ordered(4) nowait
+    for (i = 2; i < f + 3; i++)
+      for (j = d; j > 0; j--)
+	for (k = 0; k < d + 2; k++)
+	  for (l = 0; l < d + 4; l++)
+	    {
+	      #pragma omp ordered doacross (source:)
+	      #pragma omp ordered doacross (sink:omp_cur_iteration-1)
+	      if (!e)
+		abort ();
+	    }
+    #pragma omp for nowait
+    for (i = 0; i < N; i++)
+      if (a[i] != 3 * (i >= 1))
+	abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 4; k++)
+	  if (b[i][j][k] != 3 * (i >= 3 && i < N / 16 - 1 && (j & 1) == 0 && k >= 1))
+	    abort ();
+    #pragma omp for collapse(3) nowait
+    for (i = 0; i < N / 32; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 8; k++)
+	  if (c[i][j][k] != 3 * (i >= 2 && j >= 2 && (k & 1) == 0))
+	    abort ();
+    #pragma omp for collapse(2) private(k) nowait
+    for (i = 0; i < N / 16; i++)
+      for (j = 0; j < 8; j++)
+	for (k = 0; k < 6; k++)
+	  if (g[i][j][k] != 3 * (i < N / 16 - 1 && (j & 1) == 0 && k >= 3))
+	    abort ();
+  }
+  return 0;
+}
