diff mbox series

[OG13,4/6] OpenMP: New c/c++ testcases for imperfectly-nested loops

Message ID 20230614220804.917436-5-sandra@codesourcery.com
State New
Headers show
Series OpenMP: Support for imperfectly-nested loops | expand

Commit Message

Sandra Loosemore June 14, 2023, 10:08 p.m. UTC
gcc/testsuite/ChangeLog
	* c-c++-common/gomp/imperfect1.c: New.
	* c-c++-common/gomp/imperfect2.c: New.
	* c-c++-common/gomp/imperfect3.c: New.
	* c-c++-common/gomp/imperfect4.c: New.
	* c-c++-common/gomp/imperfect5.c: New.

libgomp/ChangeLog
	* testsuite/libgomp.c-c++-common/imperfect-transform-1.c: New.
	* testsuite/libgomp.c-c++-common/imperfect-transform-2.c: New.
	* testsuite/libgomp.c-c++-common/imperfect1.c: New.
	* testsuite/libgomp.c-c++-common/imperfect2.c: New.
	* testsuite/libgomp.c-c++-common/imperfect3.c: New.
	* testsuite/libgomp.c-c++-common/imperfect4.c: New.
	* testsuite/libgomp.c-c++-common/imperfect5.c: New.
	* testsuite/libgomp.c-c++-common/imperfect6.c: New.
	* testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c: New.
	* testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c: New.
	* testsuite/libgomp.c-c++-common/target-imperfect1.c: New.
	* testsuite/libgomp.c-c++-common/target-imperfect2.c: New.
	* testsuite/libgomp.c-c++-common/target-imperfect3.c: New.
	* testsuite/libgomp.c-c++-common/target-imperfect4.c: New.
---
 gcc/testsuite/ChangeLog.omp                   |   8 ++
 gcc/testsuite/c-c++-common/gomp/imperfect1.c  |  38 ++++++
 gcc/testsuite/c-c++-common/gomp/imperfect2.c  |  34 +++++
 gcc/testsuite/c-c++-common/gomp/imperfect3.c  |  33 +++++
 gcc/testsuite/c-c++-common/gomp/imperfect4.c  |  33 +++++
 gcc/testsuite/c-c++-common/gomp/imperfect5.c  |  57 ++++++++
 libgomp/ChangeLog.omp                         |  17 +++
 .../imperfect-transform-1.c                   |  79 +++++++++++
 .../imperfect-transform-2.c                   |  79 +++++++++++
 .../libgomp.c-c++-common/imperfect1.c         |  76 +++++++++++
 .../libgomp.c-c++-common/imperfect2.c         | 114 ++++++++++++++++
 .../libgomp.c-c++-common/imperfect3.c         | 119 +++++++++++++++++
 .../libgomp.c-c++-common/imperfect4.c         | 117 ++++++++++++++++
 .../libgomp.c-c++-common/imperfect5.c         |  49 +++++++
 .../libgomp.c-c++-common/imperfect6.c         | 115 ++++++++++++++++
 .../target-imperfect-transform-1.c            |  82 ++++++++++++
 .../target-imperfect-transform-2.c            |  82 ++++++++++++
 .../libgomp.c-c++-common/target-imperfect1.c  |  81 ++++++++++++
 .../libgomp.c-c++-common/target-imperfect2.c  | 122 +++++++++++++++++
 .../libgomp.c-c++-common/target-imperfect3.c  | 125 ++++++++++++++++++
 .../libgomp.c-c++-common/target-imperfect4.c  | 122 +++++++++++++++++
 21 files changed, 1582 insertions(+)
 create mode 100644 gcc/testsuite/c-c++-common/gomp/imperfect1.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/imperfect2.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/imperfect3.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/imperfect4.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/imperfect5.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/imperfect-transform-1.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/imperfect-transform-2.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/imperfect1.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/imperfect2.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/imperfect3.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/imperfect4.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/imperfect5.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/imperfect6.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-imperfect1.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-imperfect2.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-imperfect3.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-imperfect4.c
diff mbox series

Patch

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index d42813684e2..72d7b52256a 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,11 @@ 
+2023-06-13  Sandra Loosemore  <sandra@codesourcery.com>
+
+	* c-c++-common/gomp/imperfect1.c: New.
+	* c-c++-common/gomp/imperfect2.c: New.
+	* c-c++-common/gomp/imperfect3.c: New.
+	* c-c++-common/gomp/imperfect4.c: New.
+	* c-c++-common/gomp/imperfect5.c: New.
+
 2023-06-13  Sandra Loosemore  <sandra@codesourcery.com>
 
 	* c-c++-common/goacc/tile-2.c: Adjust expected error patterns.
diff --git a/gcc/testsuite/c-c++-common/gomp/imperfect1.c b/gcc/testsuite/c-c++-common/gomp/imperfect1.c
new file mode 100644
index 00000000000..705626ad169
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/imperfect1.c
@@ -0,0 +1,38 @@ 
+/* { dg-do compile } */
+
+/* This test case is expected to fail due to errors.  */
+
+int f1 (int depth, int iter);
+int f2 (int depth, int iter);
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for collapse(3)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+#pragma omp barrier	/* { dg-error "intervening code must not contain OpenMP directives" } */
+	  f1 (1, j);
+	  if (i == 2)
+	    continue;	/* { dg-error "invalid exit" } */
+	  else
+	    break;	/* { dg-error "invalid exit" } */
+	  for (k = 0; k < a3; k++)
+	    {
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      for (k = 0; k < a3; k++)	/* { dg-error "loop not permitted in intervening code " } */
+	{
+	  f1 (2, k);
+	  f2 (2, k);
+	}
+      f2 (0, i);
+    }
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/imperfect2.c b/gcc/testsuite/c-c++-common/gomp/imperfect2.c
new file mode 100644
index 00000000000..dff17dd3ca5
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/imperfect2.c
@@ -0,0 +1,34 @@ 
+/* { dg-do compile } */
+
+/* This test case is expected to fail due to errors.  */
+
+/* These functions that are part of the OpenMP runtime API would ordinarily
+   be declared in omp.h, but we don't have that here.  */
+extern int omp_get_num_threads(void);
+extern int omp_get_max_threads(void);
+
+int f1 (int depth, int iter);
+int f2 (int depth, int iter);
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+#pragma omp for collapse(3)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      for (j = 0; j < omp_get_num_threads (); j++)  /* This is OK */
+	{
+	  f1 (1, omp_get_num_threads ());  /* { dg-error "not permitted in intervening code" } */
+	  for (k = omp_get_num_threads (); k < a3; k++)  /* This is OK */
+	    {
+	      f1 (2, omp_get_num_threads ());
+	      f2 (2, omp_get_max_threads ());
+	    }
+	  f2 (1, omp_get_max_threads ());  /* { dg-error "not permitted in intervening code" } */
+	}
+      f2 (0, i);
+    }
+}
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/imperfect3.c b/gcc/testsuite/c-c++-common/gomp/imperfect3.c
new file mode 100644
index 00000000000..ad727ed3170
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/imperfect3.c
@@ -0,0 +1,33 @@ 
+/* { dg-do compile } */
+
+/* This test case is expected to fail due to errors.  */
+
+/* Test that the imperfectly-nested loops with the ordered clause gives
+   an error, and that there is only one error (and not one on every
+   intervening statement).  */
+
+int f1 (int depth, int iter);
+int f2 (int depth, int iter);
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for ordered(3)
+  for (i = 0; i < a1; i++)  /* { dg-error "inner loops must be perfectly nested" } */
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+	  for (k = 0; k < a3; k++)
+	    {
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+}
+
diff --git a/gcc/testsuite/c-c++-common/gomp/imperfect4.c b/gcc/testsuite/c-c++-common/gomp/imperfect4.c
new file mode 100644
index 00000000000..1a0c07cd48e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/imperfect4.c
@@ -0,0 +1,33 @@ 
+/* { dg-do compile } */
+
+/* This test case is expected to fail due to errors.  */
+
+int f1 (int depth, int iter);
+int f2 (int depth, int iter);
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for collapse(4)
+  for (i = 0; i < a1; i++)	/* { dg-error "not enough nested loops" } */
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+	  for (k = 0; k < a3; k++)
+	    {
+	      /* According to the grammar, this is intervening code; we
+		 don't know that we are also missing a nested for loop
+		 until we have parsed this whole compound expression.  */
+#pragma omp barrier	/* { dg-error "intervening code must not contain OpenMP directives" } */
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+}
+
diff --git a/gcc/testsuite/c-c++-common/gomp/imperfect5.c b/gcc/testsuite/c-c++-common/gomp/imperfect5.c
new file mode 100644
index 00000000000..585d89ff789
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/imperfect5.c
@@ -0,0 +1,57 @@ 
+/* { dg-do compile } */
+
+/* This test case is expected to fail due to errors.  */
+
+int f1 (int depth, int iter);
+int f2 (int depth, int iter);
+int ijk (int x, int y, int z);
+void f3 (int sum);
+
+/* This function isn't particularly meaningful, but it should compile without
+   error.  */
+int s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+  int r = 0;
+
+#pragma omp simd collapse(3) reduction (inscan, +:r)
+  for (i = 0; i < a1; i++)
+    {
+      for (j = 0; j < a2; j++)
+	{
+	  for (k = 0; k < a3; k++)
+	    {
+	      r = r + ijk (i, j, k);
+#pragma omp scan exclusive (r)
+	      f3 (r);
+	    }
+	}
+    }
+  return r;
+}
+
+/* Adding intervening code should trigger an error.  */
+int s2 (int a1, int a2, int a3)
+{
+  int i, j, k;
+  int r = 0;
+
+#pragma omp simd collapse(3) reduction (inscan, +:r)
+  for (i = 0; i < a1; i++)  /* { dg-error "inner loops must be perfectly nested" } */
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+	  for (k = 0; k < a3; k++)
+	    {
+	      r = r + ijk (i, j, k);
+#pragma omp scan exclusive (r)
+	      f3 (r);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+  return r;
+}
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 29ef68843bc..cd6dfa6af2d 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,20 @@ 
+2023-06-13  Sandra Loosemore  <sandra@codesourcery.com>
+
+	* testsuite/libgomp.c-c++-common/imperfect-transform-1.c: New.
+	* testsuite/libgomp.c-c++-common/imperfect-transform-2.c: New.
+	* testsuite/libgomp.c-c++-common/imperfect1.c: New.
+	* testsuite/libgomp.c-c++-common/imperfect2.c: New.
+	* testsuite/libgomp.c-c++-common/imperfect3.c: New.
+	* testsuite/libgomp.c-c++-common/imperfect4.c: New.
+	* testsuite/libgomp.c-c++-common/imperfect5.c: New.
+	* testsuite/libgomp.c-c++-common/imperfect6.c: New.
+	* testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c: New.
+	* testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c: New.
+	* testsuite/libgomp.c-c++-common/target-imperfect1.c: New.
+	* testsuite/libgomp.c-c++-common/target-imperfect2.c: New.
+	* testsuite/libgomp.c-c++-common/target-imperfect3.c: New.
+	* testsuite/libgomp.c-c++-common/target-imperfect4.c: New.
+
 2023-06-13  Sandra Loosemore  <sandra@codesourcery.com>
 
 	* testsuite/libgomp.c++/imperfect-class-1.C: New.
diff --git a/libgomp/testsuite/libgomp.c-c++-common/imperfect-transform-1.c b/libgomp/testsuite/libgomp.c-c++-common/imperfect-transform-1.c
new file mode 100644
index 00000000000..6743594b2eb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/imperfect-transform-1.c
@@ -0,0 +1,79 @@ 
+/* { dg-do run } */
+
+/* Like imperfect1.c, but also includes loop transforms.  */
+
+static int f1count[3], f2count[3];
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  f2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for collapse(2)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+#pragma omp unroll partial
+	  for (k = 0; k < a3; k++)
+	    {
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/imperfect-transform-2.c b/libgomp/testsuite/libgomp.c-c++-common/imperfect-transform-2.c
new file mode 100644
index 00000000000..e7d6a9941b4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/imperfect-transform-2.c
@@ -0,0 +1,79 @@ 
+/* { dg-do run } */
+
+/* Like imperfect1.c, but also includes loop transforms.  */
+
+static int f1count[3], f2count[3];
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  f2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for collapse(2)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+#pragma omp tile sizes(5)
+	  for (k = 0; k < a3; k++)
+	    {
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/imperfect1.c b/libgomp/testsuite/libgomp.c-c++-common/imperfect1.c
new file mode 100644
index 00000000000..cafdcaf25b0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/imperfect1.c
@@ -0,0 +1,76 @@ 
+/* { dg-do run } */
+
+static int f1count[3], f2count[3];
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  f2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for collapse(3)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+	  for (k = 0; k < a3; k++)
+	    {
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/imperfect2.c b/libgomp/testsuite/libgomp.c-c++-common/imperfect2.c
new file mode 100644
index 00000000000..e2098006eab
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/imperfect2.c
@@ -0,0 +1,114 @@ 
+/* { dg-do run } */
+
+static int f1count[3], f2count[3];
+static int g1count[3], g2count[3];
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  f2count[depth]++;
+  return iter;
+}
+
+int g1 (int depth, int iter)
+{
+  g1count[depth]++;
+  return iter;
+}
+
+int g2 (int depth, int iter)
+{
+  g2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for collapse(3)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      {
+	g1 (0, i);
+	for (j = 0; j < a2; j++)
+	  {
+	    f1 (1, j);
+	    {
+	      g1 (1, j);
+	      for (k = 0; k < a3; k++)
+		{
+		  f1 (2, k);
+		  {
+		    g1 (2, k);
+		    g2 (2, k);
+		  }
+		  f2 (2, k);
+		}
+	      g2 (1, j);
+	    }
+	  f2 (1, j);
+	  }
+	g2 (0, i);
+      }
+      f2 (0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  g1count[0] = 0;
+  g1count[1] = 0;
+  g1count[2] = 0;
+  g2count[0] = 0;
+  g2count[1] = 0;
+  g2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+  if (g1count[0] != f1count[0]) abort ();
+  if (g2count[0] != f1count[0]) abort ();
+  if (g1count[1] != f1count[1]) abort ();
+  if (g2count[1] != f1count[1]) abort ();
+  if (g1count[2] != f1count[2]) abort ();
+  if (g2count[2] != f1count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/imperfect3.c b/libgomp/testsuite/libgomp.c-c++-common/imperfect3.c
new file mode 100644
index 00000000000..feb5e32d1d6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/imperfect3.c
@@ -0,0 +1,119 @@ 
+/* { dg-do run } */
+
+/* Like imperfect2.c, but includes bindings in the blocks.  */
+
+static int f1count[3], f2count[3];
+static int g1count[3], g2count[3];
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  f2count[depth]++;
+  return iter;
+}
+
+int g1 (int depth, int iter)
+{
+  g1count[depth]++;
+  return iter;
+}
+
+int g2 (int depth, int iter)
+{
+  g2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for collapse(3)
+  for (i = 0; i < a1; i++)
+    {
+      int local0 = 0;
+      f1 (local0, i);
+      {
+	g1 (local0, i);
+	for (j = 0; j < a2; j++)
+	  {
+	    int local1 = 1;
+	    f1 (local1, j);
+	    {
+	      g1 (local1, j);
+	      for (k = 0; k < a3; k++)
+		{
+		  int local2 = 2;
+		  f1 (local2, k);
+		  {
+		    g1 (local2, k);
+		    g2 (local2, k);
+		  }
+		  f2 (local2, k);
+		}
+	      g2 (local1, j);
+	    }
+	  f2 (local1, j);
+	  }
+	g2 (local0, i);
+      }
+      f2 (local0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  g1count[0] = 0;
+  g1count[1] = 0;
+  g1count[2] = 0;
+  g2count[0] = 0;
+  g2count[1] = 0;
+  g2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+  if (g1count[0] != f1count[0]) abort ();
+  if (g2count[0] != f1count[0]) abort ();
+  if (g1count[1] != f1count[1]) abort ();
+  if (g2count[1] != f1count[1]) abort ();
+  if (g1count[2] != f1count[2]) abort ();
+  if (g2count[2] != f1count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/imperfect4.c b/libgomp/testsuite/libgomp.c-c++-common/imperfect4.c
new file mode 100644
index 00000000000..e29301bfbad
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/imperfect4.c
@@ -0,0 +1,117 @@ 
+/* { dg-do run } */
+
+/* Like imperfect2.c, but includes blocks that are themselves intervening
+   code.  */
+
+static int f1count[3], f2count[3];
+static int g1count[3], g2count[3];
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  f2count[depth]++;
+  return iter;
+}
+
+int g1 (int depth, int iter)
+{
+  g1count[depth]++;
+  return iter;
+}
+
+int g2 (int depth, int iter)
+{
+  g2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp for collapse(3)
+  for (i = 0; i < a1; i++)
+    {
+      { f1 (0, i); }
+      {
+	g1 (0, i);
+	for (j = 0; j < a2; j++)
+	  {
+	    { f1 (1, j); }
+	    {
+	      { g1 (1, j); }
+	      for (k = 0; k < a3; k++)
+		{
+		  f1 (2, k);
+		  {
+		    g1 (2, k);
+		    g2 (2, k);
+		  }
+		  f2 (2, k);
+		}
+	      { g2 (1, j); }
+	    }
+	    { f2 (1, j); }
+	  }
+	{ g2 (0, i); }
+      }
+      { f2 (0, i); }
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  g1count[0] = 0;
+  g1count[1] = 0;
+  g1count[2] = 0;
+  g2count[0] = 0;
+  g2count[1] = 0;
+  g2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+  if (g1count[0] != f1count[0]) abort ();
+  if (g2count[0] != f1count[0]) abort ();
+  if (g1count[1] != f1count[1]) abort ();
+  if (g2count[1] != f1count[1]) abort ();
+  if (g1count[2] != f1count[2]) abort ();
+  if (g2count[2] != f1count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/imperfect5.c b/libgomp/testsuite/libgomp.c-c++-common/imperfect5.c
new file mode 100644
index 00000000000..7bd4f12d472
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/imperfect5.c
@@ -0,0 +1,49 @@ 
+/* { dg-do run } */
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+static int inner_loop_count = 0;
+static int intervening_code_count = 0;
+
+void
+g (int x, int y)
+{
+  inner_loop_count++;
+}
+
+int
+foo (int imax, int jmax)
+{
+  int j = 0;
+
+#pragma omp for collapse(2)
+  for (int i = 0; i < imax; ++i)
+    {
+      /* All the intervening code at the same level must be executed
+	 the same number of times.  */
+      ++intervening_code_count;
+      for (int j = 0; j < jmax; ++j)
+	{
+	  g (i, j);
+	}
+      /* This is the outer j, not the one from the inner collapsed loop.  */
+      ++j;
+    }
+  return j;
+}
+
+int
+main (void)
+{
+  int j = foo (5, 3);
+  if (j != intervening_code_count)
+    abort ();
+  if (inner_loop_count != 5 * 3)
+    abort ();
+  if (intervening_code_count < 5 || intervening_code_count > 5 * 3)
+    abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/imperfect6.c b/libgomp/testsuite/libgomp.c-c++-common/imperfect6.c
new file mode 100644
index 00000000000..808c6540890
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/imperfect6.c
@@ -0,0 +1,115 @@ 
+/* { dg-do run } */
+
+/* Like imperfect4.c, but bind the iteration variables in the loops.  */
+
+static int f1count[3], f2count[3];
+static int g1count[3], g2count[3];
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  f2count[depth]++;
+  return iter;
+}
+
+int g1 (int depth, int iter)
+{
+  g1count[depth]++;
+  return iter;
+}
+
+int g2 (int depth, int iter)
+{
+  g2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+
+#pragma omp for collapse(3)
+  for (int i = 0; i < a1; i++)
+    {
+      { f1 (0, i); }
+      {
+	g1 (0, i);
+	for (int j = 0; j < a2; j++)
+	  {
+	    { f1 (1, j); }
+	    {
+	      { g1 (1, j); }
+	      for (int k = 0; k < a3; k++)
+		{
+		  f1 (2, k);
+		  {
+		    g1 (2, k);
+		    g2 (2, k);
+		  }
+		  f2 (2, k);
+		}
+	      { g2 (1, j); }
+	    }
+	    { f2 (1, j); }
+	  }
+	{ g2 (0, i); }
+      }
+      { f2 (0, i); }
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  g1count[0] = 0;
+  g1count[1] = 0;
+  g1count[2] = 0;
+  g2count[0] = 0;
+  g2count[1] = 0;
+  g2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+  if (g1count[0] != f1count[0]) abort ();
+  if (g2count[0] != f1count[0]) abort ();
+  if (g1count[1] != f1count[1]) abort ();
+  if (g2count[1] != f1count[1]) abort ();
+  if (g1count[2] != f1count[2]) abort ();
+  if (g2count[2] != f1count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c
new file mode 100644
index 00000000000..0e33e028ac2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c
@@ -0,0 +1,82 @@ 
+/* { dg-do run } */
+
+/* Like imperfect-transform.c, but enables offloading.  */
+
+static int f1count[3], f2count[3];
+#pragma omp declare target enter (f1count, f2count)
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  #pragma omp atomic
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  #pragma omp atomic
+  f2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp target parallel for collapse(2) map(always, tofrom:f1count, f2count)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+#pragma omp unroll partial
+	  for (k = 0; k < a3; k++)
+	    {
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c
new file mode 100644
index 00000000000..78986e8d3ae
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c
@@ -0,0 +1,82 @@ 
+/* { dg-do run } */
+
+/* Like imperfect-transform.c, but enables offloading.  */
+
+static int f1count[3], f2count[3];
+#pragma omp declare target enter (f1count, f2count)
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  #pragma omp atomic
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  #pragma omp atomic
+  f2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp target parallel for collapse(2) map(always, tofrom:f1count, f2count)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+#pragma omp tile sizes(5)
+	  for (k = 0; k < a3; k++)
+	    {
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-imperfect1.c b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect1.c
new file mode 100644
index 00000000000..53bc611ace3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect1.c
@@ -0,0 +1,81 @@ 
+/* { dg-do run } */
+
+/* Like imperfect1.c, but enables offloading.  */
+
+static int f1count[3], f2count[3];
+#pragma omp declare target enter (f1count, f2count)
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  #pragma omp atomic
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  #pragma omp atomic
+  f2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp target parallel for collapse(3) map(always, tofrom:f1count, f2count)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      for (j = 0; j < a2; j++)
+	{
+	  f1 (1, j);
+	  for (k = 0; k < a3; k++)
+	    {
+	      f1 (2, k);
+	      f2 (2, k);
+	    }
+	  f2 (1, j);
+	}
+      f2 (0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-imperfect2.c b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect2.c
new file mode 100644
index 00000000000..bc2901a517e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect2.c
@@ -0,0 +1,122 @@ 
+/* { dg-do run } */
+
+/* Like imperfect2.c, but enables offloading.  */
+
+static int f1count[3], f2count[3];
+static int g1count[3], g2count[3];
+#pragma omp declare target enter (f1count, f2count)
+#pragma omp declare target enter (g1count, g2count)
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  #pragma omp atomic
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  #pragma omp atomic
+  f2count[depth]++;
+  return iter;
+}
+
+int g1 (int depth, int iter)
+{
+  #pragma omp atomic
+  g1count[depth]++;
+  return iter;
+}
+
+int g2 (int depth, int iter)
+{
+  #pragma omp atomic
+  g2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp target parallel for collapse(3) map(always, tofrom:f1count, f2count, g1count, g2count)
+  for (i = 0; i < a1; i++)
+    {
+      f1 (0, i);
+      {
+	g1 (0, i);
+	for (j = 0; j < a2; j++)
+	  {
+	    f1 (1, j);
+	    {
+	      g1 (1, j);
+	      for (k = 0; k < a3; k++)
+		{
+		  f1 (2, k);
+		  {
+		    g1 (2, k);
+		    g2 (2, k);
+		  }
+		  f2 (2, k);
+		}
+	      g2 (1, j);
+	    }
+	  f2 (1, j);
+	  }
+	g2 (0, i);
+      }
+      f2 (0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  g1count[0] = 0;
+  g1count[1] = 0;
+  g1count[2] = 0;
+  g2count[0] = 0;
+  g2count[1] = 0;
+  g2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+  if (g1count[0] != f1count[0]) abort ();
+  if (g2count[0] != f1count[0]) abort ();
+  if (g1count[1] != f1count[1]) abort ();
+  if (g2count[1] != f1count[1]) abort ();
+  if (g1count[2] != f1count[2]) abort ();
+  if (g2count[2] != f1count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-imperfect3.c b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect3.c
new file mode 100644
index 00000000000..ddcfcf4b7eb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect3.c
@@ -0,0 +1,125 @@ 
+/* { dg-do run } */
+
+/* Like imperfect3.c, but enables offloading.  */
+
+static int f1count[3], f2count[3];
+static int g1count[3], g2count[3];
+#pragma omp declare target enter (f1count, f2count)
+#pragma omp declare target enter (g1count, g2count)
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  #pragma omp atomic
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  #pragma omp atomic
+  f2count[depth]++;
+  return iter;
+}
+
+int g1 (int depth, int iter)
+{
+  #pragma omp atomic
+  g1count[depth]++;
+  return iter;
+}
+
+int g2 (int depth, int iter)
+{
+  #pragma omp atomic
+  g2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp target parallel for collapse(3) map(always, tofrom:f1count, f2count, g1count, g2count)
+  for (i = 0; i < a1; i++)
+    {
+      int local0 = 0;
+      f1 (local0, i);
+      {
+	g1 (local0, i);
+	for (j = 0; j < a2; j++)
+	  {
+	    int local1 = 1;
+	    f1 (local1, j);
+	    {
+	      g1 (local1, j);
+	      for (k = 0; k < a3; k++)
+		{
+		  int local2 = 2;
+		  f1 (local2, k);
+		  {
+		    g1 (local2, k);
+		    g2 (local2, k);
+		  }
+		  f2 (local2, k);
+		}
+	      g2 (local1, j);
+	    }
+	  f2 (local1, j);
+	  }
+	g2 (local0, i);
+      }
+      f2 (local0, i);
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  g1count[0] = 0;
+  g1count[1] = 0;
+  g1count[2] = 0;
+  g2count[0] = 0;
+  g2count[1] = 0;
+  g2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+  if (g1count[0] != f1count[0]) abort ();
+  if (g2count[0] != f1count[0]) abort ();
+  if (g1count[1] != f1count[1]) abort ();
+  if (g2count[1] != f1count[1]) abort ();
+  if (g1count[2] != f1count[2]) abort ();
+  if (g2count[2] != f1count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-imperfect4.c b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect4.c
new file mode 100644
index 00000000000..ede488977b8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-imperfect4.c
@@ -0,0 +1,122 @@ 
+/* { dg-do run } */
+
+/* Like imperfect4.c, but enables offloading.  */
+
+static int f1count[3], f2count[3];
+static int g1count[3], g2count[3];
+#pragma omp declare target enter (f1count, f2count)
+#pragma omp declare target enter (g1count, g2count)
+
+#ifndef __cplusplus
+extern void abort (void);
+#else
+extern "C" void abort (void);
+#endif
+
+int f1 (int depth, int iter)
+{
+  #pragma omp atomic
+  f1count[depth]++;
+  return iter;
+}
+
+int f2 (int depth, int iter)
+{
+  #pragma omp atomic
+  f2count[depth]++;
+  return iter;
+}
+
+int g1 (int depth, int iter)
+{
+  #pragma omp atomic
+  g1count[depth]++;
+  return iter;
+}
+
+int g2 (int depth, int iter)
+{
+  #pragma omp atomic
+  g2count[depth]++;
+  return iter;
+}
+
+void s1 (int a1, int a2, int a3)
+{
+  int i, j, k;
+
+#pragma omp target parallel for collapse(3) map(always, tofrom:f1count, f2count, g1count, g2count)
+  for (i = 0; i < a1; i++)
+    {
+      { f1 (0, i); }
+      {
+	g1 (0, i);
+	for (j = 0; j < a2; j++)
+	  {
+	    { f1 (1, j); }
+	    {
+	      { g1 (1, j); }
+	      for (k = 0; k < a3; k++)
+		{
+		  f1 (2, k);
+		  {
+		    g1 (2, k);
+		    g2 (2, k);
+		  }
+		  f2 (2, k);
+		}
+	      { g2 (1, j); }
+	    }
+	    { f2 (1, j); }
+	  }
+	{ g2 (0, i); }
+      }
+      { f2 (0, i); }
+    }
+}
+
+int
+main (void)
+{
+  f1count[0] = 0;
+  f1count[1] = 0;
+  f1count[2] = 0;
+  f2count[0] = 0;
+  f2count[1] = 0;
+  f2count[2] = 0;
+
+  g1count[0] = 0;
+  g1count[1] = 0;
+  g1count[2] = 0;
+  g2count[0] = 0;
+  g2count[1] = 0;
+  g2count[2] = 0;
+
+  s1 (3, 4, 5);
+
+  /* All intervening code at the same depth must be executed the same
+     number of times. */
+  if (f1count[0] != f2count[0]) abort ();
+  if (f1count[1] != f2count[1]) abort ();
+  if (f1count[2] != f2count[2]) abort ();
+  if (g1count[0] != f1count[0]) abort ();
+  if (g2count[0] != f1count[0]) abort ();
+  if (g1count[1] != f1count[1]) abort ();
+  if (g2count[1] != f1count[1]) abort ();
+  if (g1count[2] != f1count[2]) abort ();
+  if (g2count[2] != f1count[2]) abort ();
+
+  /* Intervening code must be executed at least as many times as the loop
+     that encloses it. */
+  if (f1count[0] < 3) abort ();
+  if (f1count[1] < 3 * 4) abort ();
+
+  /* Intervening code must not be executed more times than the number
+     of logical iterations. */
+  if (f1count[0] > 3 * 4 * 5) abort ();
+  if (f1count[1] > 3 * 4 * 5) abort ();
+
+  /* Check that the innermost loop body is executed exactly the number
+     of logical iterations expected. */
+  if (f1count[2] != 3 * 4 * 5) abort ();
+}