Add libgomp.oacc-c-c++-common/f-asyncwait-{1,2,3}.c

2017-11-15  Tom de Vries  <tom@codesourcery.com>

	* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: New test, copied
	from asyncwait-1.f90.  Rewrite into C.  Rewrite from float to int.
	* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-2.c: New test, copied
	from asyncwait-2.f90.  Rewrite into C.  Rewrite from float to int.
	* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-3.c: New test, copied
	from asyncwait-3.f90.  Rewrite into C.  Rewrite from float to int.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@254769 138bc75d-0d04-0410-961f-82ee72b054a4

4 files changed, 430 insertions, 0 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index eaa50ddc66b..7331d412a82 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,12 @@
+2017-11-15  Tom de Vries  <tom@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: New test, copied
+	from asyncwait-1.f90.  Rewrite into C.  Rewrite from float to int.
+	* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-2.c: New test, copied
+	from asyncwait-2.f90.  Rewrite into C.  Rewrite from float to int.
+	* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-3.c: New test, copied
+	from asyncwait-3.f90.  Rewrite into C.  Rewrite from float to int.
+
 2017-11-14  Tom de Vries  <tom@codesourcery.com>
 
 	* testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Allow to run for
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c
new file mode 100644
index 00000000000..cf851707dc7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c
@@ -0,0 +1,297 @@
+/* { dg-do run } */
+
+/* Based on asyncwait-1.f90.  */
+
+#include <stdlib.h>
+
+#define N 64
+
+int
+main (void)
+{
+  int *a, *b, *c, *d, *e;
+
+  a = (int*)malloc (N * sizeof (*a));
+  b = (int*)malloc (N * sizeof (*b));
+  c = (int*)malloc (N * sizeof (*c));
+  d = (int*)malloc (N * sizeof (*d));
+  e = (int*)malloc (N * sizeof (*e));
+
+  for (int i = 0; i < N; ++i)
+    {
+      a[i] = 3;
+      b[i] = 0;
+    }
+
+#pragma acc data copy (a[0:N]) copy (b[0:N])
+  {
+
+#pragma acc parallel async
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      b[i] = a[i];
+
+#pragma acc wait
+  }
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (a[i] != 3)
+	abort ();
+      if (b[i] != 3)
+	abort ();
+    }
+
+  for (int i = 0; i < N; ++i)
+    {
+      a[i] = 2;
+      b[i] = 0;
+    }
+
+#pragma acc data copy (a[0:N]) copy (b[0:N])
+  {
+#pragma acc parallel async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      b[i] = a[i];
+
+#pragma acc wait (1)
+  }
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (a[i] != 2) abort ();
+      if (b[i] != 2) abort ();
+    }
+
+  for (int i = 0; i < N; ++i)
+    {
+      a[i] = 3;
+      b[i] = 0;
+      c[i] = 0;
+      d[i] = 0;
+    }
+
+#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N])
+  {
+
+#pragma acc parallel async (1)
+    for (int i = 0; i < N; ++i)
+      b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc parallel async (1)
+    for (int i = 0; i < N; ++i)
+      c[i] = (a[i] * 4) / a[i];
+
+
+#pragma acc parallel async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+#pragma acc wait (1)
+  }
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (a[i] != 3)
+	abort ();
+      if (b[i] != 9)
+	abort ();
+      if (c[i] != 4)
+	abort ();
+      if (d[i] != 1)
+	abort ();
+    }
+
+  for (int i = 0; i < N; ++i)
+    {
+      a[i] = 2;
+      b[i] = 0;
+      c[i] = 0;
+      d[i] = 0;
+      e[i] = 0;
+    }
+
+#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N])
+  {
+
+#pragma acc parallel async (1)
+    for (int i = 0; i < N; ++i)
+      b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc parallel async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      c[i] = (a[i] * 4) / a[i];
+
+#pragma acc parallel async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+
+#pragma acc parallel wait (1) async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      e[i] = a[i] + b[i] + c[i] + d[i];
+
+#pragma acc wait (1)
+  }
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (a[i] != 2)
+	abort ();
+      if (b[i] != 4)
+	abort ();
+      if (c[i] != 4)
+	abort ();
+      if (d[i] != 1)
+	abort ();
+      if (e[i] != 11)
+	abort ();
+    }
+
+  for (int i = 0; i < N; ++i)
+    {
+      a[i] = 3;
+      b[i] = 0;
+    }
+
+#pragma acc data copy (a[0:N]) copy (b[0:N])
+  {
+
+#pragma acc kernels async
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      b[i] = a[i];
+
+#pragma acc wait
+  }
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (a[i] != 3)
+	abort ();
+      if (b[i] != 3)
+	abort ();
+    }
+
+  for (int i = 0; i < N; ++i)
+    {
+      a[i] = 2;
+      b[i] = 0;
+    }
+
+#pragma acc data copy (a[0:N]) copy (b[0:N])
+  {
+#pragma acc kernels async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      b[i] = a[i];
+
+#pragma acc wait (1)
+  }
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (a[i] != 2)
+	abort ();
+      if (b[i] != 2)
+	abort ();
+    }
+
+  for (int i = 0; i < N; ++i)
+    {
+      a[i] = 3;
+      b[i] = 0;
+      c[i] = 0;
+      d[i] = 0;
+    }
+
+#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N])
+  {
+#pragma acc kernels async (1)
+    for (int i = 0; i < N; ++i)
+      b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc kernels async (1)
+    for (int i = 0; i < N; ++i)
+      c[i] = (a[i] * 4) / a[i];
+
+#pragma acc kernels async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+#pragma acc wait (1)
+  }
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (a[i] != 3)
+	abort ();
+      if (b[i] != 9)
+	abort ();
+      if (c[i] != 4)
+	abort ();
+      if (d[i] != 1)
+	abort ();
+    }
+
+  for (int i = 0; i < N; ++i)
+    {
+      a[i] = 2;
+      b[i] = 0;
+      c[i] = 0;
+      d[i] = 0;
+      e[i] = 0;
+    }
+
+#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N])
+  {
+#pragma acc kernels async (1)
+    for (int i = 0; i < N; ++i)
+      b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc kernels async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      c[i] = (a[i] * 4) / a[i];
+
+#pragma acc kernels async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+#pragma acc kernels wait (1) async (1)
+#pragma acc loop
+    for (int i = 0; i < N; ++i)
+      e[i] = a[i] + b[i] + c[i] + d[i];
+
+#pragma acc wait (1)
+  }
+
+  for (int i = 0; i < N; ++i)
+    {
+      if (a[i] != 2)
+	abort ();
+      if (b[i] != 4)
+	abort ();
+      if (c[i] != 4)
+	abort ();
+      if (d[i] != 1)
+	abort ();
+      if (e[i] != 11)
+	abort ();
+    }
+
+  free (a);
+  free (b);
+  free (c);
+  free (d);
+  free (e);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-2.c
new file mode 100644
index 00000000000..5298e4c54f7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-2.c
@@ -0,0 +1,61 @@
+/* { dg-do run } */
+
+/* Based on asyncwait-2.f90.  */
+
+#include <stdlib.h>
+
+#define N 64
+
+int *a, *b, *c;
+
+int
+main (void)
+{
+  a = (int *)malloc (N * sizeof (*a));
+  b = (int *)malloc (N * sizeof (*b));
+  c = (int *)malloc (N * sizeof (*c));
+
+#pragma acc parallel copy (a[0:N]) async (0)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    a[i] = 1;
+
+#pragma acc parallel copy (b[0:N]) async (1)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    b[i] = 1;
+
+#pragma acc parallel copy (a[0:N], b[0:N], c[0:N]) wait (0, 1)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    c[i] = a[i] + b[i];
+
+  for (int i = 0; i < N; ++i)
+    if (c[i] != 2)
+      abort ();
+
+#if 1
+#pragma acc kernels copy (a[0:N]) async (0)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    a[i] = 1;
+
+#pragma acc kernels copy (b[0:N]) async (1)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    b[i] = 1;
+
+#pragma acc kernels copy (a[0:N], b[0:N], c[0:N]) wait (0, 1)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    c[i] = a[i] + b[i];
+
+  for (int i = 0; i < N; ++i)
+    if (c[i] != 2)
+      abort ();
+#endif
+
+  free (a);
+  free (b);
+  free (c);
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-3.c
new file mode 100644
index 00000000000..319eea61dc7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-3.c
@@ -0,0 +1,63 @@
+/* { dg-do run } */
+
+/* Based on asyncwait-3.f90.  */
+
+#include <stdlib.h>
+
+#define N 64
+
+int
+main (void)
+{
+  int *a, *b, *c;
+
+  a = (int *)malloc (N * sizeof (*a));
+  b = (int *)malloc (N * sizeof (*b));
+  c = (int *)malloc (N * sizeof (*c));
+
+#pragma acc parallel copy (a[0:N]) async (0)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    a[i] = 1;
+
+#pragma acc parallel copy (b[0:N]) async (1)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    b[i] = 1;
+
+#pragma acc wait (0, 1)
+
+#pragma acc parallel copy (a[0:N], b[0:N], c[0:N])
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    c[i] = a[i] + b[i];
+
+  for (int i = 0; i < N; ++i)
+    if (c[i] != 2)
+      abort ();
+
+#pragma acc kernels copy (a[0:N]) async (0)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    a[i] = 1;
+
+#pragma acc kernels copy (b[0:N]) async (1)
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    b[i] = 1;
+
+#pragma acc wait (0, 1)
+
+#pragma acc kernels copy (a[0:N], b[0:N], c[0:N])
+#pragma acc loop
+  for (int i = 0; i < N; ++i)
+    c[i] = a[i] + b[i];
+
+  for (int i = 0; i < N; ++i)
+    if (c[i] != 2)
+      abort ();
+
+  free (a);
+  free (b);
+  free (c);
+}