Merge branch 'master' into devel/sphinx

10 files changed, 1151 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-5.c b/libgomp/testsuite/libgomp.c-c++-common/icv-5.c
new file mode 100644
index 0000000..431cfc7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-5.c
@@ -0,0 +1,25 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_0 "42" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1 "43" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_2 "44" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "45" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "46" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS "47" } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  if (omp_get_max_teams () != 47)
+    abort ();
+
+  int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
+  for (int i=0; i < num_devices; i++)
+    #pragma omp target device (i)
+      if (omp_get_max_teams () != 42 + i)
+	abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-6.c b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
new file mode 100644
index 0000000..7151bd1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
@@ -0,0 +1,45 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "42" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "43" } */
+/* { dg-set-target-env-var OMP_SCHEDULE_ALL "guided,4" } */
+/* { dg-set-target-env-var OMP_DYNAMIC_ALL "true" } */
+/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "44" } */
+/* { dg-set-target-env-var OMP_THREAD_LIMIT_ALL "45" } */
+/* { dg-set-target-env-var OMP_NUM_THREADS_ALL "46,3,2" } */
+/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_ALL "47" } */
+/* { dg-set-target-env-var OMP_PROC_BIND_ALL "spread" } */
+/* { dg-set-target-env-var OMP_WAIT_POLICY_ALL "active" } */
+
+/* This tests the hierarchical usage of ICVs on the device, i.e. if
+   OMP_NUM_TEAMS_DEV_<device_num> is not configured, then the value of
+   OMP_NUM_TEAMS_DEV should be used.  And if there is no environment variable
+   without suffix, then the corresponding _ALL variant should be used.  */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  enum omp_sched_t kind;
+  int chunk_size;
+  omp_get_schedule(&kind, &chunk_size);
+
+  if (omp_get_max_teams () != 42
+      || !omp_get_dynamic ()
+      || kind != 3 || chunk_size != 4
+      || omp_get_teams_thread_limit () != 44
+      || omp_get_thread_limit () != 45
+      || omp_get_max_threads () != 46
+      || omp_get_proc_bind () != omp_proc_bind_spread
+      || omp_get_max_active_levels () != 47)
+    abort ();
+
+  int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
+  for (int i=0; i < num_devices; i++)
+    #pragma omp target device (i)
+      if (omp_get_max_teams () != 43)
+	abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-7.c b/libgomp/testsuite/libgomp.c-c++-common/icv-7.c
new file mode 100644
index 0000000..70a716d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-7.c
@@ -0,0 +1,26 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "42" } */
+
+/* This tests the hierarchical usage of ICVs on the host and on devices, i.e. if
+   OMP_NUM_TEAMS_DEV_<device_num>, OMP_NUM_TEAMS_DEV, and
+   OMP_NUM_TEAMS are not configured, then the value of
+   OMP_NUM_TEAMS_ALL should be used for the host as well as for the
+   devices.  */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  if (omp_get_max_teams () != 42)
+    abort ();
+
+  int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
+  for (int i=0; i < num_devices; i++)
+    #pragma omp target device (i)
+      if (omp_get_max_teams () != 42)
+	abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-8.c b/libgomp/testsuite/libgomp.c-c++-common/icv-8.c
new file mode 100644
index 0000000..f25ce45
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-8.c
@@ -0,0 +1,26 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1234567890 "42" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_ "43" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_01 "44" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_a "45" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_12345678901 "46" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_-1 "47" } */
+/* { dg-set-target-env-var "OMP_NUM_TEAMS_DEV_ 1" "48" } */
+/* { dg-set-target-env-var "OMP_NUM_TEAMS_DEV_00" "49" } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  return 0;
+}
+
+/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_=43.*" { target native } } */
+/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_01=44.*" { target native } } */
+/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_a=45.*" { target native } } */
+/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_12345678901=46.*" { target native } } */
+/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_-1=47.*" { target native } } */
+/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_ 1=48.*" { target native } } */
+/* { dg-output ".*Invalid device number in OMP_NUM_TEAMS_DEV_00=49.*" { target native } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-1.c b/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-1.c
new file mode 100644
index 0000000..9ea7ade
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-1.c
@@ -0,0 +1,119 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_THREAD_LIMIT_DEV_24 "42" } */
+/* { dg-set-target-env-var OMP_THREAD_LIMIT_ALL "43" } */
+/* { dg-set-target-env-var OMP_THREAD_LIMIT_DEV "44" } */
+/* { dg-set-target-env-var OMP_THREAD_LIMIT "45" } */
+/* { dg-set-target-env-var OMP_DEFAULT_DEVICE "42" } */
+/* { dg-set-target-env-var OMP_SCHEDULE_DEV_24 "guided,4" } */
+/* { dg-set-target-env-var OMP_SCHEDULE_ALL "dynamic" } */
+/* { dg-set-target-env-var OMP_SCHEDULE_DEV "guided,1" } */
+/* { dg-set-target-env-var OMP_SCHEDULE "guided,2" } */
+/* { dg-set-target-env-var OMP_DYNAMIC_DEV_24 "true" } */
+
+/* { dg-set-target-env-var OMP_DYNAMIC_ALL "true" } */
+/* { dg-set-target-env-var OMP_DYNAMIC_DEV "true" } */
+/* { dg-set-target-env-var OMP_DYNAMIC "true" } */
+/* { dg-set-target-env-var OMP_NUM_THREADS "4,3,2" } */
+/* { dg-set-target-env-var OMP_NUM_THREADS_ALL "45,46,47" } */
+/* { dg-set-target-env-var OMP_NUM_THREADS_DEV "42,43,44" } */
+/* { dg-set-target-env-var OMP_NUM_THREADS_DEV_24 "14,13,12" } */
+/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS "42" } */
+/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_ALL "43" } */
+/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_DEV "44" } */
+
+/* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_DEV_24 "45" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS "42" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "43" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "44" } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_24 "45" } */
+/* { dg-set-target-env-var OMP_PROC_BIND "spread" } */
+/* { dg-set-target-env-var OMP_PROC_BIND_ALL "close" } */
+/* { dg-set-target-env-var OMP_PROC_BIND_DEV "spread,spread" } */
+/* { dg-set-target-env-var OMP_PROC_BIND_DEV_24 "spread,close" } */
+/* { dg-set-target-env-var OMP_STACKSIZE "42" } */
+
+/* { dg-set-target-env-var OMP_STACKSIZE_ALL "42 M" } */
+/* { dg-set-target-env-var OMP_STACKSIZE_DEV "43 k" } */
+/* { dg-set-target-env-var OMP_STACKSIZE_DEV_24 "44" } */
+/* { dg-set-target-env-var OMP_WAIT_POLICY "active" } */
+/* { dg-set-target-env-var OMP_WAIT_POLICY_ALL "ACTIVE" } */
+/* { dg-set-target-env-var OMP_WAIT_POLICY_DEV "passive" } */
+/* { dg-set-target-env-var OMP_WAIT_POLICY_DEV_24 "PASSIVE" } */
+/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT "42" } */
+/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "43" } */
+/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV "44" } */
+
+/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_24 "45" } */
+/* { dg-set-target-env-var OMP_CANCELLATION "true" } */
+/* { dg-set-target-env-var OMP_DISPLAY_AFFINITY "true" } */
+/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */
+/* { dg-set-target-env-var OMP_MAX_TASK_PRIORITY "20" } */
+/* { dg-set-target-env-var OMP_ALLOCATOR "omp_const_mem_alloc" } */
+/* { dg-set-target-env-var OMP_NESTED "false" } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  omp_display_env (1);
+  return 0;
+}
+
+/* { dg-output ".*\\\[host] OMP_DYNAMIC = 'TRUE'.*" { target native } } */
+/* { dg-output ".*\\\[all] OMP_DYNAMIC = 'TRUE'.*" { target native } } */
+/* { dg-output ".*\\\[device] OMP_DYNAMIC = 'TRUE'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_DYNAMIC = 'TRUE'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_NUM_THREADS = '4,3,2'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_NUM_THREADS = '45,46,47'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_NUM_THREADS = '42,43,44'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_NUM_THREADS = '14,13,12'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_SCHEDULE = 'GUIDED,2'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_SCHEDULE = 'DYNAMIC'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_SCHEDULE = 'GUIDED'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_SCHEDULE = 'GUIDED,4'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_PROC_BIND = 'SPREAD'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_PROC_BIND = 'CLOSE'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_PROC_BIND = 'SPREAD,SPREAD'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_PROC_BIND = 'SPREAD,CLOSE'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_STACKSIZE = '43008'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_STACKSIZE = '44040192'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_STACKSIZE = '44032'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_STACKSIZE = '45056'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_WAIT_POLICY = 'ACTIVE'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_WAIT_POLICY = 'ACTIVE'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_WAIT_POLICY = 'PASSIVE'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_WAIT_POLICY = 'PASSIVE'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_THREAD_LIMIT = '45'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_THREAD_LIMIT = '43'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_THREAD_LIMIT = '44'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_THREAD_LIMIT = '42'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_MAX_ACTIVE_LEVELS = '42'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_MAX_ACTIVE_LEVELS = '43'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_MAX_ACTIVE_LEVELS = '44'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_MAX_ACTIVE_LEVELS = '45'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_NUM_TEAMS = '42'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_NUM_TEAMS = '43'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_NUM_TEAMS = '44'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_NUM_TEAMS = '45'.*" { target native } } */
+
+/* { dg-output ".*\\\[host] OMP_TEAMS_THREAD_LIMIT = '42'.*" { target native } } */
+/* { dg-output ".*\\\[all\] OMP_TEAMS_THREAD_LIMIT = '43'.*" { target native } } */
+/* { dg-output ".*\\\[device\] OMP_TEAMS_THREAD_LIMIT = '44'.*" { target native } } */
+/* { dg-output ".*\\\[24\] OMP_TEAMS_THREAD_LIMIT = '45'.*" { target native } } */
+
+/* { dg-output ".*\\\[all] OMP_CANCELLATION = 'TRUE'.*" { target native } } */
+/* { dg-output ".*\\\[all] OMP_DEFAULT_DEVICE = '42'.*" { target native } } */
+/* { dg-output ".*\\\[all] OMP_MAX_TASK_PRIORITY = '20'.*" { target native } } */
+/* { dg-output ".*\\\[all] OMP_DISPLAY_AFFINITY = 'TRUE'.*" { target native } } */
+/* { dg-output ".*\\\[host] OMP_ALLOCATOR = 'omp_const_mem_alloc'.*" { target native } } */
+/* { dg-output ".*\\\[all] OMP_TARGET_OFFLOAD = 'MANDATORY'.*" { target native } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-2.c b/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-2.c
new file mode 100644
index 0000000..e1beef4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/omp-display-env-2.c
@@ -0,0 +1,22 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var OMP_NUM_TEAMS "42" } */
+
+/* This test checks if omp_display_env outputs the initial ICV values although
+   the value was updated.  */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  omp_display_env (1);
+  omp_set_num_teams (24);
+  if (omp_get_max_teams () != 24)
+    abort ();
+  omp_display_env (1);
+
+  return 0;
+}
+
+/* { dg-output ".*\\\[host] OMP_NUM_TEAMS = '42'.*\\\[host] OMP_NUM_TEAMS = '42'" { target native } } */
diff --git a/libgomp/testsuite/libgomp.c/doacross-4.c b/libgomp/testsuite/libgomp.c/doacross-4.c
new file mode 100644
index 0000000..4bc451d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/doacross-4.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.c/doacross-5.c b/libgomp/testsuite/libgomp.c/doacross-5.c
new file mode 100644
index 0000000..7bb1993
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/doacross-5.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.c/doacross-6.c b/libgomp/testsuite/libgomp.c/doacross-6.c
new file mode 100644
index 0000000..de5dee6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/doacross-6.c
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.c/doacross-7.c b/libgomp/testsuite/libgomp.c/doacross-7.c
new file mode 100644
index 0000000..f0ad388
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/doacross-7.c
@@ -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;
+}