From: Tom de Vries Date: Mon, 18 Jan 2016 12:53:03 +0000 (+0000) Subject: Add oacc kernels test in libgomp X-Git-Tag: basepoints/gcc-7~1531 X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=5d7804a940dca2bddbb34978d45e06ac7ef13c80;p=thirdparty%2Fgcc.git Add oacc kernels test in libgomp 2016-01-18 Tom de Vries * 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. From-SVN: r232515 --- diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index a626cb82e752..2efc516624aa 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,30 @@ +2016-01-18 Tom de Vries + + * 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. + 2016-01-15 Jakub Jelinek * task.c (GOMP_PLUGIN_target_task_completion): Add missing return. diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c new file mode 100644 index 000000000000..13e57bd98778 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c new file mode 100644 index 000000000000..f61a74a00fb2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-3.c @@ -0,0 +1,34 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c new file mode 100644 index 000000000000..2e4100fa7a2f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c @@ -0,0 +1,36 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c new file mode 100644 index 000000000000..b3e736bf3589 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c new file mode 100644 index 000000000000..8b9affabedc8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c @@ -0,0 +1,36 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c new file mode 100644 index 000000000000..83d4e7faf8f4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c new file mode 100644 index 000000000000..01d5e5e4c24a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c @@ -0,0 +1,36 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c new file mode 100644 index 000000000000..61d1283a36a2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c new file mode 100644 index 000000000000..f7f04cb9ed5c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-collapse.c @@ -0,0 +1,40 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c new file mode 100644 index 000000000000..96b6e4eaf455 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c @@ -0,0 +1,5 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ +/* { dg-additional-options "-g" } */ + +#include "kernels-loop.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c new file mode 100644 index 000000000000..1433cb203b8b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-mod-not-zero.c @@ -0,0 +1,41 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c new file mode 100644 index 000000000000..fd0d5b193a81 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-n.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c new file mode 100644 index 000000000000..21d2599f0aa1 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-nest.c @@ -0,0 +1,26 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop.c new file mode 100644 index 000000000000..3762e5a8b9ac --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop.c @@ -0,0 +1,41 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c new file mode 100644 index 000000000000..511e25f882aa --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ + +#include + +#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; +}