Merge trunk into svelinaro-dev/sve

12 files changed, 505 insertions, 2 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index f68604ca424..7331d412a82 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,33 @@
+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
+	non-nvidia devices.
+
+2017-11-07  Jakub Jelinek  <jakub@redhat.com>
+
+	PR c++/82835
+	* testsuite/libgomp.c++/pr82835.C: New test.
+
+2017-11-06  Martin Liska  <mliska@suse.cz>
+
+	* testsuite/libgomp.c++/loop-2.C: Return a value
+	for functions with non-void return type, or change type to void,
+	or add -Wno-return-type for test.
+	* testsuite/libgomp.c++/loop-4.C: Likewise.
+	* testsuite/libgomp.c++/parallel-1.C: Likewise.
+	* testsuite/libgomp.c++/shared-1.C: Likewise.
+	* testsuite/libgomp.c++/single-1.C: Likewise.
+	* testsuite/libgomp.c++/single-2.C: Likewise.
+
 2017-10-31  Tom de Vries  <tom@codesourcery.com>
 
 	* plugin/plugin-hsa.c (HSA_LOG): Remove semicolon after
diff --git a/libgomp/testsuite/libgomp.c++/loop-2.C b/libgomp/testsuite/libgomp.c++/loop-2.C
index ea3dc588afd..77144b8e7c6 100644
--- a/libgomp/testsuite/libgomp.c++/loop-2.C
+++ b/libgomp/testsuite/libgomp.c++/loop-2.C
@@ -15,6 +15,7 @@ void parloop (int *a)
     a[i] = i + 3;
 }
 
+int
 main()
 {
   int i, a[N];
diff --git a/libgomp/testsuite/libgomp.c++/loop-4.C b/libgomp/testsuite/libgomp.c++/loop-4.C
index 731f2345021..a940854c637 100644
--- a/libgomp/testsuite/libgomp.c++/loop-4.C
+++ b/libgomp/testsuite/libgomp.c++/loop-4.C
@@ -1,5 +1,6 @@
 extern "C" void abort (void);
 
+int
 main()
 {
   int i, a;
diff --git a/libgomp/testsuite/libgomp.c++/parallel-1.C b/libgomp/testsuite/libgomp.c++/parallel-1.C
index 3c931471328..ce338d0ddf9 100644
--- a/libgomp/testsuite/libgomp.c++/parallel-1.C
+++ b/libgomp/testsuite/libgomp.c++/parallel-1.C
@@ -8,6 +8,7 @@ foo (void)
   return 10;
 }
 
+int
 main ()
 {
   int A = 0;
diff --git a/libgomp/testsuite/libgomp.c++/pr82835.C b/libgomp/testsuite/libgomp.c++/pr82835.C
new file mode 100644
index 00000000000..df64ecfb1cf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr82835.C
@@ -0,0 +1,34 @@
+// PR c++/82835
+// { dg-do run }
+
+int a, b;
+
+template <class>
+struct C {
+  C (int x = a) : c (5) { if (x != 137) __builtin_abort (); }
+  int c;
+};
+
+struct D {
+  void foo ();
+  int d;
+};
+
+void
+D::foo ()
+{
+  C<int> c;
+#pragma omp for private (c)
+  for (b = 0; b < d; b++)
+    c.c++;
+}
+
+int
+main ()
+{
+  a = 137;
+  D d;
+  d.d = 16;
+  d.foo ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/shared-1.C b/libgomp/testsuite/libgomp.c++/shared-1.C
index 334a553ce23..2f61daa1213 100644
--- a/libgomp/testsuite/libgomp.c++/shared-1.C
+++ b/libgomp/testsuite/libgomp.c++/shared-1.C
@@ -53,6 +53,7 @@ parallel (int a, int b)
     abort ();
 }
 
+int
 main()
 {
   parallel (1, 2);
diff --git a/libgomp/testsuite/libgomp.c++/single-1.C b/libgomp/testsuite/libgomp.c++/single-1.C
index e318a48ca5c..221236f24f2 100644
--- a/libgomp/testsuite/libgomp.c++/single-1.C
+++ b/libgomp/testsuite/libgomp.c++/single-1.C
@@ -1,5 +1,6 @@
 extern "C" void abort (void);
 
+int
 main()
 {
   int i = 0;
diff --git a/libgomp/testsuite/libgomp.c++/single-2.C b/libgomp/testsuite/libgomp.c++/single-2.C
index c2dd228568d..d24b1d85e66 100644
--- a/libgomp/testsuite/libgomp.c++/single-2.C
+++ b/libgomp/testsuite/libgomp.c++/single-2.C
@@ -7,6 +7,7 @@ struct X
   int c;
 };
 
+int
 main()
 {
   int i = 0;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
index d478ce2eef5..e780845a793 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
@@ -1,9 +1,11 @@
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-/* { dg-additional-options "-lcuda" } */
+/* { dg-do run } */
+/* { dg-additional-options "-lcuda" { target openacc_nvidia_accel_selected } } */
 
 #include <openacc.h>
 #include <stdlib.h>
+#if defined ACC_DEVICE_TYPE_nvidia
 #include "cuda.h"
+#endif
 
 #include <stdio.h>
 #include <sys/time.h>
@@ -11,14 +13,18 @@
 int
 main (int argc, char **argv)
 {
+#if defined ACC_DEVICE_TYPE_nvidia
     CUresult r;
     CUstream stream1;
+#endif
     int N = 128; //1024 * 1024;
     float *a, *b, *c, *d, *e;
     int i;
     int nbytes;
 
+#if defined ACC_DEVICE_TYPE_nvidia
     acc_init (acc_device_nvidia);
+#endif
 
     nbytes = N * sizeof (float);
 
@@ -210,6 +216,7 @@ main (int argc, char **argv)
     }
 
 
+#if defined ACC_DEVICE_TYPE_nvidia
     r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING);
     if (r != CUDA_SUCCESS)
     {
@@ -218,6 +225,7 @@ main (int argc, char **argv)
     }
 
     acc_set_cuda_stream (1, stream1);
+#endif
 
     for (i = 0; i < N; i++)
     {
@@ -642,6 +650,7 @@ main (int argc, char **argv)
     }
 
 
+#if defined ACC_DEVICE_TYPE_nvidia
     r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING);
     if (r != CUDA_SUCCESS)
     {
@@ -650,6 +659,7 @@ main (int argc, char **argv)
     }
 
     acc_set_cuda_stream (1, stream1);
+#endif
 
     for (i = 0; i < N; i++)
     {
@@ -892,7 +902,9 @@ main (int argc, char **argv)
             abort ();
     }
 
+#if defined ACC_DEVICE_TYPE_nvidia
     acc_shutdown (acc_device_nvidia);
+#endif
 
     return 0;
 }
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);
+}