Add oacc kernels test in libgomp
2015-11-09 Tom de Vries <tom@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-2.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-enter-exit-2.c:
Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-enter-exit.c:
Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data-update.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-data.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop.c: Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-parallel-loop-data-enter-exit.c:
Same.
* testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c: Same.
---
.../libgomp.oacc-c-c++-common/kernels-loop-2.c | 47 ++++++++++++++++++++++
.../libgomp.oacc-c-c++-common/kernels-loop-3.c | 34 ++++++++++++++++
.../kernels-loop-and-seq-2.c | 36 +++++++++++++++++
.../kernels-loop-and-seq-3.c | 37 +++++++++++++++++
.../kernels-loop-and-seq-4.c | 36 +++++++++++++++++
.../kernels-loop-and-seq-5.c | 37 +++++++++++++++++
.../kernels-loop-and-seq-6.c | 36 +++++++++++++++++
.../kernels-loop-and-seq.c | 37 +++++++++++++++++
.../kernels-loop-collapse.c | 40 ++++++++++++++++++
.../libgomp.oacc-c-c++-common/kernels-loop-g.c | 5 +++
.../kernels-loop-mod-not-zero.c | 41 +++++++++++++++++++
.../libgomp.oacc-c-c++-common/kernels-loop-n.c | 47 ++++++++++++++++++++++
.../libgomp.oacc-c-c++-common/kernels-loop-nest.c | 26 ++++++++++++
.../libgomp.oacc-c-c++-common/kernels-loop.c | 41 +++++++++++++++++++
.../libgomp.oacc-c-c++-common/kernels-reduction.c | 37 +++++++++++++++++
15 files changed, 537 insertions(+)
new file mode 100644
@@ -0,0 +1,47 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N (1024 * 512)
+#define COUNTERTYPE unsigned int
+
+int
+main (void)
+{
+ unsigned int *__restrict a;
+ unsigned int *__restrict b;
+ unsigned int *__restrict c;
+
+ a = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+ b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+ c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+
+#pragma acc kernels copyout (a[0:N])
+ {
+ for (COUNTERTYPE i = 0; i < N; i++)
+ a[i] = i * 2;
+ }
+
+#pragma acc kernels copyout (b[0:N])
+ {
+ for (COUNTERTYPE i = 0; i < N; i++)
+ b[i] = i * 4;
+ }
+
+#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
+ {
+ for (COUNTERTYPE ii = 0; ii < N; ii++)
+ c[ii] = a[ii] + b[ii];
+ }
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ if (c[i] != a[i] + b[i])
+ abort ();
+
+ free (a);
+ free (b);
+ free (c);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,34 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N (1024 * 512)
+#define COUNTERTYPE unsigned int
+
+int
+main (void)
+{
+ unsigned int i;
+
+ unsigned int *__restrict c;
+
+ c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ c[i] = i * 2;
+
+#pragma acc kernels copy (c[0:N])
+ {
+ for (COUNTERTYPE ii = 0; ii < N; ii++)
+ c[ii] = c[ii] + ii + 1;
+ }
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ if (c[i] != i * 2 + i + 1)
+ abort ();
+
+ free (c);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+#pragma acc kernels copy (a[0:N])
+ {
+ a[0] = a[0] + 1;
+
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 1)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+
+#pragma acc kernels copy (a[0:N])
+ {
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+
+ a[0] = 2;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 2)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+#pragma acc kernels copy (a[0:N])
+ {
+ a[0] = 2;
+
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 1)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+ int r;
+#pragma acc kernels copyout(r) copy (a[0:N])
+ {
+ r = a[0];
+
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+ }
+
+ return r;
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 0)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+#pragma acc kernels copy (a[0:N])
+ {
+ int r = a[0];
+
+ for (int i = 0; i < n; i++)
+ a[i] = 1 + r;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 1)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+
+#pragma acc kernels copy (a[0:N])
+ {
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+
+ a[0] = a[0] + 1;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 2)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,40 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 100
+
+int a[N][N];
+
+void __attribute__((noinline, noclone))
+foo (int m, int n)
+{
+ int i, j;
+ #pragma acc kernels
+ {
+#pragma acc loop collapse(2)
+ for (i = 0; i < m; i++)
+ for (j = 0; j < n; j++)
+ a[i][j] = 1;
+ }
+}
+
+int
+main (void)
+{
+ int i, j;
+
+ for (i = 0; i < N; i++)
+ for (j = 0; j < N; j++)
+ a[i][j] = 0;
+
+ foo (N, N);
+
+ for (i = 0; i < N; i++)
+ for (j = 0; j < N; j++)
+ if (a[i][j] != 1)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,5 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+/* { dg-additional-options "-g" } */
+
+#include "kernels-loop.c"
new file mode 100644
@@ -0,0 +1,41 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N ((1024 * 512) + 1)
+#define COUNTERTYPE unsigned int
+
+int
+main (void)
+{
+ unsigned int *__restrict a;
+ unsigned int *__restrict b;
+ unsigned int *__restrict c;
+
+ a = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+ b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+ c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ a[i] = i * 2;
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ b[i] = i * 4;
+
+#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
+ {
+ for (COUNTERTYPE ii = 0; ii < N; ii++)
+ c[ii] = a[ii] + b[ii];
+ }
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ if (c[i] != a[i] + b[i])
+ abort ();
+
+ free (a);
+ free (b);
+ free (c);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,47 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N ((1024 * 512) + 1)
+#define COUNTERTYPE unsigned int
+
+static int __attribute__((noinline,noclone))
+foo (COUNTERTYPE n)
+{
+ unsigned int *__restrict a;
+ unsigned int *__restrict b;
+ unsigned int *__restrict c;
+
+ a = (unsigned int *__restrict)malloc (n * sizeof (unsigned int));
+ b = (unsigned int *__restrict)malloc (n * sizeof (unsigned int));
+ c = (unsigned int *__restrict)malloc (n * sizeof (unsigned int));
+
+ for (COUNTERTYPE i = 0; i < n; i++)
+ a[i] = i * 2;
+
+ for (COUNTERTYPE i = 0; i < n; i++)
+ b[i] = i * 4;
+
+#pragma acc kernels copyin (a[0:n], b[0:n]) copyout (c[0:n])
+ {
+ for (COUNTERTYPE ii = 0; ii < n; ii++)
+ c[ii] = a[ii] + b[ii];
+ }
+
+ for (COUNTERTYPE i = 0; i < n; i++)
+ if (c[i] != a[i] + b[i])
+ abort ();
+
+ free (a);
+ free (b);
+ free (c);
+
+ return 0;
+}
+
+int
+main (void)
+{
+ return foo (N);
+}
new file mode 100644
@@ -0,0 +1,26 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 1000
+
+int
+main (void)
+{
+ int x[N][N];
+
+#pragma acc kernels copyout (x)
+ {
+ for (int ii = 0; ii < N; ii++)
+ for (int jj = 0; jj < N; jj++)
+ x[ii][jj] = ii + jj + 3;
+ }
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ if (x[i][j] != i + j + 3)
+ abort ();
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,41 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N (1024 * 512)
+#define COUNTERTYPE unsigned int
+
+int
+main (void)
+{
+ unsigned int *__restrict a;
+ unsigned int *__restrict b;
+ unsigned int *__restrict c;
+
+ a = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+ b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+ c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ a[i] = i * 2;
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ b[i] = i * 4;
+
+#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
+ {
+ for (COUNTERTYPE ii = 0; ii < N; ii++)
+ c[ii] = a[ii] + b[ii];
+ }
+
+ for (COUNTERTYPE i = 0; i < N; i++)
+ if (c[i] != a[i] + b[i])
+ abort ();
+
+ free (a);
+ free (b);
+ free (c);
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define n 10000
+
+unsigned int a[n];
+
+void __attribute__((noinline,noclone))
+foo (void)
+{
+ int i;
+ unsigned int sum = 1;
+
+#pragma acc kernels copyin (a[0:n]) copy (sum)
+ {
+ for (i = 0; i < n; ++i)
+ sum += a[i];
+ }
+
+ if (sum != 5001)
+ abort ();
+}
+
+int
+main ()
+{
+ int i;
+
+ for (i = 0; i < n; ++i)
+ a[i] = i % 2;
+
+ foo ();
+
+ return 0;
+}