Skip to content

R
riscv-gcc-1
Commit 7de562ee by Thomas Schwinge Committed by Thomas Schwinge
Options

Revise libgomp.oacc-c-c++-common/data-2-lib.c, libgomp.oacc-c-c++-common/data-2.c 

These are meant to be functionally equivalent (but no longer are), just using
different means.  Also, use the OpenACC "*_async" functions recently added.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Revise.
	* testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.

From-SVN: r267149
parent 17469af7
Showing with 125 additions and 157 deletions
libgomp/ChangeLog
2018-12-14 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Revise.
* testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.
2018-12-14 Chung-Lin Tang <cltang@codesourcery.com> 2018-12-14 Chung-Lin Tang <cltang@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Adjust.
......
libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c
/* This test is similar to data-2.c, but it uses acc_* library functions /* Test asynchronous, unstructed data regions, runtime library variant. */
to move data. */ /* See also data-2.c. */
/* { dg-do run } */
#include <stdlib.h> #include <stdlib.h>
#undef NDEBUG
#include <assert.h> #include <assert.h>
#include <openacc.h> #include <openacc.h>
int int
main (int argc, char **argv) main (int argc, char **argv)
{ {
int N = 128; //1024 * 1024; int N = 12345;
float *a, *b, *c, *d, *e; float *a, *b, *c, *d, *e;
void *d_a, *d_b, *d_c, *d_d; void *d_a, *d_b, *d_c, *d_d;
int i; int i;
...@@ -30,19 +29,21 @@ main (int argc, char **argv) ...@@ -30,19 +29,21 @@ main (int argc, char **argv)
b[i] = 0.0; b[i] = 0.0;
} }
d_a = acc_copyin (a, nbytes); acc_copyin_async (a, nbytes, acc_async_noval);
d_b = acc_copyin (b, nbytes); acc_copyin_async (b, nbytes, acc_async_noval);
acc_copyin (&N, sizeof (int)); acc_copyin_async (&N, sizeof (int), acc_async_noval);
#pragma acc parallel present (a[0:N], b[0:N], N) async wait #pragma acc parallel present (a[0:N], b[0:N], N) async
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
b[i] = a[i]; b[i] = a[i];
acc_wait_all (); d_a = acc_deviceptr (a);
acc_memcpy_from_device_async (a, d_a, nbytes, acc_async_noval);
d_b = acc_deviceptr (b);
acc_memcpy_from_device_async (b, d_b, nbytes, acc_async_noval);
acc_memcpy_from_device (a, d_a, nbytes); acc_wait (acc_async_noval);
acc_memcpy_from_device (b, d_b, nbytes);
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
{ {
...@@ -56,19 +57,19 @@ main (int argc, char **argv) ...@@ -56,19 +57,19 @@ main (int argc, char **argv)
b[i] = 0.0; b[i] = 0.0;
} }
acc_update_device (a, nbytes); acc_update_device_async (a, nbytes, 1);
acc_update_device (b, nbytes); acc_update_device_async (b, nbytes, 1);
#pragma acc parallel present (a[0:N], b[0:N], N) async (1) #pragma acc parallel present (a[0:N], b[0:N], N) async (1)
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
b[i] = a[i]; b[i] = a[i];
acc_memcpy_from_device_async (a, d_a, nbytes, 1);
acc_memcpy_from_device_async (b, d_b, nbytes, 1);
acc_wait (1); acc_wait (1);
acc_memcpy_from_device (a, d_a, nbytes);
acc_memcpy_from_device (b, d_b, nbytes);
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
{ {
assert (a[i] == 2.0); assert (a[i] == 2.0);
...@@ -83,46 +84,42 @@ main (int argc, char **argv) ...@@ -83,46 +84,42 @@ main (int argc, char **argv)
d[i] = 0.0; d[i] = 0.0;
} }
acc_update_device (a, nbytes); acc_update_device_async (a, nbytes, 0);
acc_update_device (b, nbytes); acc_update_device_async (b, nbytes, 1);
d_c = acc_copyin (c, nbytes); acc_copyin_async (c, nbytes, 2);
d_d = acc_copyin (d, nbytes); acc_copyin_async (d, nbytes, 3);
#pragma acc parallel present (a[0:N], b[0:N], N) async (1) #pragma acc parallel present (a[0:N], b[0:N], N) wait (0) async (1)
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
b[i] = (a[i] * a[i] * a[i]) / a[i]; b[i] = (a[i] * a[i] * a[i]) / a[i];
#pragma acc parallel present (a[0:N], c[0:N], N) async (2) #pragma acc parallel present (a[0:N], c[0:N], N) wait (0) async (2)
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
#pragma acc parallel present (a[0:N], d[0:N], N) async (3) #pragma acc parallel present (a[0:N], d[0:N], N) wait (0) async (3)
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
acc_wait_all (); acc_memcpy_from_device_async (a, d_a, nbytes, 0);
acc_memcpy_from_device_async (b, d_b, nbytes, 1);
d_c = acc_deviceptr (c);
acc_memcpy_from_device_async (c, d_c, nbytes, 2);
d_d = acc_deviceptr (d);
acc_memcpy_from_device_async (d, d_d, nbytes, 3);
acc_memcpy_from_device (a, d_a, nbytes); acc_wait_all_async (0);
acc_memcpy_from_device (b, d_b, nbytes); acc_wait (0);
acc_memcpy_from_device (c, d_c, nbytes);
acc_memcpy_from_device (d, d_d, nbytes);
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
{ {
if (a[i] != 3.0) assert (a[i] == 3.0);
abort (); assert (b[i] == 9.0);
assert (c[i] == 4.0);
if (b[i] != 9.0) assert (d[i] == 1.0);
abort ();
if (c[i] != 4.0)
abort ();
if (d[i] != 1.0)
abort ();
} }
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
...@@ -134,53 +131,43 @@ main (int argc, char **argv) ...@@ -134,53 +131,43 @@ main (int argc, char **argv)
e[i] = 0.0; e[i] = 0.0;
} }
acc_update_device (a, nbytes); acc_update_device_async (a, nbytes, 10);
acc_update_device (b, nbytes); acc_update_device_async (b, nbytes, 11);
acc_update_device (c, nbytes); acc_update_device_async (c, nbytes, 12);
acc_update_device (d, nbytes); acc_update_device_async (d, nbytes, 13);
acc_copyin (e, nbytes); acc_copyin_async (e, nbytes, 14);
#pragma acc parallel present (a[0:N], b[0:N], N) async (1) #pragma acc parallel present (a[0:N], b[0:N], N) wait (10) async (11)
for (int ii = 0; ii < N; ii++) for (int ii = 0; ii < N; ii++)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
#pragma acc parallel present (a[0:N], c[0:N], N) async (2) #pragma acc parallel present (a[0:N], c[0:N], N) wait (10) async (12)
for (int ii = 0; ii < N; ii++) for (int ii = 0; ii < N; ii++)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
#pragma acc parallel present (a[0:N], d[0:N], N) async (3) #pragma acc parallel present (a[0:N], d[0:N], N) wait (10) async (13)
for (int ii = 0; ii < N; ii++) for (int ii = 0; ii < N; ii++)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N], N) \ #pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N], N) wait (11) wait (12) wait (13) async (14)
wait (1, 2, 3) async (4)
for (int ii = 0; ii < N; ii++) for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
acc_copyout_async (a, nbytes, 10);
acc_copyout_async (b, nbytes, 11);
acc_copyout_async (c, nbytes, 12);
acc_copyout_async (d, nbytes, 13);
acc_copyout_async (e, nbytes, 14);
acc_delete_async (&N, sizeof (int), 15);
acc_wait_all (); acc_wait_all ();
acc_copyout (a, nbytes);
acc_copyout (b, nbytes);
acc_copyout (c, nbytes);
acc_copyout (d, nbytes);
acc_copyout (e, nbytes);
acc_delete (&N, sizeof (int));
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
{ {
if (a[i] != 2.0) assert (a[i] == 2.0);
abort (); assert (b[i] == 4.0);
assert (c[i] == 4.0);
if (b[i] != 4.0) assert (d[i] == 1.0);
abort (); assert (e[i] == 11.0);
if (c[i] != 4.0)
abort ();
if (d[i] != 1.0)
abort ();
if (e[i] != 11.0)
abort ();
} }
return 0; return 0;
......
libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
/* Test 'acc enter/exit data' regions. */ /* Test asynchronous, unstructed data regions, directives variant. */
/* See also data-2-lib.c. */
/* { dg-do run } */
/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } } */
#include <stdlib.h> #include <stdlib.h>
#undef NDEBUG
#include <assert.h>
int int
main (int argc, char **argv) main (int argc, char **argv)
{ {
int N = 128; //1024 * 1024; int N = 12345;
float *a, *b, *c, *d, *e; float *a, *b, *c, *d, *e;
int i; int i;
int nbytes; int nbytes;
...@@ -27,48 +27,24 @@ main (int argc, char **argv) ...@@ -27,48 +27,24 @@ main (int argc, char **argv)
b[i] = 0.0; b[i] = 0.0;
} }
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async #pragma acc enter data copyin (a[0:N]) async
#pragma acc parallel present (a[0:N], b[0:N]) async wait #pragma acc enter data copyin (b[0:N]) async
#pragma acc loop #pragma acc enter data copyin (N) async
for (i = 0; i < N; i++)
b[i] = a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
#pragma acc wait
for (i = 0; i < N; i++)
{
if (a[i] != 3.0)
abort ();
if (b[i] != 3.0)
abort ();
}
for (i = 0; i < N; i++)
{
a[i] = 3.0;
b[i] = 0.0;
}
#pragma acc enter data copyin (a[0:N]) async #pragma acc parallel present (a[0:N], b[0:N], N) async
#pragma acc enter data copyin (b[0:N]) async wait
#pragma acc enter data copyin (N) async wait
#pragma acc parallel async wait
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
b[i] = a[i]; b[i] = a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) delete (N) wait async #pragma acc update self (a[0:N]) async
#pragma acc update self (b[0:N]) async
#pragma acc wait #pragma acc wait
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
{ {
if (a[i] != 3.0) assert (a[i] == 3.0);
abort (); assert (b[i] == 3.0);
if (b[i] != 3.0)
abort ();
} }
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
...@@ -77,22 +53,23 @@ main (int argc, char **argv) ...@@ -77,22 +53,23 @@ main (int argc, char **argv)
b[i] = 0.0; b[i] = 0.0;
} }
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1) #pragma acc update device (a[0:N]) async (1)
#pragma acc parallel present (a[0:N], b[0:N]) async (1) #pragma acc update device (b[0:N]) async (1)
#pragma acc parallel present (a[0:N], b[0:N], N) async (1)
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
b[i] = a[i]; b[i] = a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait (1) async (1) #pragma acc update self (a[0:N]) async (1)
#pragma acc update self (b[0:N]) async (1)
#pragma acc wait (1) #pragma acc wait (1)
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
{ {
if (a[i] != 2.0) assert (a[i] == 2.0);
abort (); assert (b[i] == 2.0);
if (b[i] != 2.0)
abort ();
} }
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
...@@ -103,39 +80,40 @@ main (int argc, char **argv) ...@@ -103,39 +80,40 @@ main (int argc, char **argv)
d[i] = 0.0; d[i] = 0.0;
} }
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1) #pragma acc update device (a[0:N]) async (0)
#pragma acc update device (b[0:N]) async (1)
#pragma acc enter data copyin (c[0:N]) async (2)
#pragma acc enter data copyin (d[0:N]) async (3)
#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) #pragma acc parallel present (a[0:N], b[0:N], N) wait (0) async (1)
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
b[i] = (a[i] * a[i] * a[i]) / a[i]; b[i] = (a[i] * a[i] * a[i]) / a[i];
#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) #pragma acc parallel present (a[0:N], c[0:N], N) wait (0) async (2)
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) #pragma acc parallel present (a[0:N], d[0:N], N) wait (0) async (3)
#pragma acc loop #pragma acc loop
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1) #pragma acc update self (a[0:N]) async (0)
#pragma acc wait (1) #pragma acc update self (b[0:N]) async (1)
#pragma acc update self (c[0:N]) async (2)
#pragma acc update self (d[0:N]) async (3)
#pragma acc wait async (0)
#pragma acc wait (0)
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
{ {
if (a[i] != 3.0) assert (a[i] == 3.0);
abort (); assert (b[i] == 9.0);
assert (c[i] == 4.0);
if (b[i] != 9.0) assert (d[i] == 1.0);
abort ();
if (c[i] != 4.0)
abort ();
if (d[i] != 1.0)
abort ();
} }
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
...@@ -147,45 +125,43 @@ main (int argc, char **argv) ...@@ -147,45 +125,43 @@ main (int argc, char **argv)
e[i] = 0.0; e[i] = 0.0;
} }
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1) #pragma acc update device (a[0:N]) async (10)
#pragma acc update device (b[0:N]) async (11)
#pragma acc update device (c[0:N]) async (12)
#pragma acc update device (d[0:N]) async (13)
#pragma acc enter data copyin (e[0:N]) async (14)
#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) #pragma acc parallel present (a[0:N], b[0:N], N) wait (10) async (11)
for (int ii = 0; ii < N; ii++) for (int ii = 0; ii < N; ii++)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) #pragma acc parallel present (a[0:N], c[0:N], N) wait (10) async (12)
for (int ii = 0; ii < N; ii++) for (int ii = 0; ii < N; ii++)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) #pragma acc parallel present (a[0:N], d[0:N], N) wait (10) async (13)
for (int ii = 0; ii < N; ii++) for (int ii = 0; ii < N; ii++)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \ #pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N], N) wait (11) wait (12) wait (13) async (14)
wait (1, 2, 3) async (4)
for (int ii = 0; ii < N; ii++) for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \ #pragma acc exit data copyout (a[0:N]) async (10)
copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) #pragma acc exit data copyout (b[0:N]) async (11)
#pragma acc wait (1) #pragma acc exit data copyout (c[0:N]) async (12)
#pragma acc exit data copyout (d[0:N]) async (13)
#pragma acc exit data copyout (e[0:N]) async (14)
#pragma acc exit data delete (N) async (15)
#pragma acc wait
for (i = 0; i < N; i++) for (i = 0; i < N; i++)
{ {
if (a[i] != 2.0) assert (a[i] == 2.0);
abort (); assert (b[i] == 4.0);
assert (c[i] == 4.0);
if (b[i] != 4.0) assert (d[i] == 1.0);
abort (); assert (e[i] == 11.0);
if (c[i] != 4.0)
abort ();
if (d[i] != 1.0)
abort ();
if (e[i] != 11.0)
abort ();
} }
return 0; return 0;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment