From a7cf26127a6a90f0ec4846ffc81c7fadeb77c3a0 Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Wed, 15 Nov 2017 13:40:58 +0000 Subject: [PATCH] Add libgomp.oacc-c-c++-common/f-asyncwait-{1,2,3}.c 2017-11-15 Tom de Vries * 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. From-SVN: r254769 --- libgomp/ChangeLog | 9 + .../libgomp.oacc-c-c++-common/f-asyncwait-1.c | 297 ++++++++++++++++++ .../libgomp.oacc-c-c++-common/f-asyncwait-2.c | 61 ++++ .../libgomp.oacc-c-c++-common/f-asyncwait-3.c | 63 ++++ 4 files changed, 430 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-3.c 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 + + * 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 * 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 + +#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 + +#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 + +#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); +} -- 2.30.2