Tested on x86_64-linux-gnu with gfx908 and gfx90a accelerators. Also,
as extra testing, ran these testcases 500 times per testcase per each
accelerator.
OK for trunk?
TIA, have a lovely day.
---------- >8 ----------
In the testcases, the kernels scheduled on queues 11, 12, 13, 14 have
data dependencies on, respectively, 'b', 'c', 'd', and 'e', as they
write to them.
However, they also have a data dependency on 'a', as they read from it.
Previously, the testcases exited 'a' on queue 10, meaning that it was
possible for the aforementioned kernels to execute and to have 'a'
pulled under their feet.
This patch adds waits for each of the kernels onto queue 10 before
freeing 'a', guaranteeing that 'a' outlives the kernels.
libgomp/ChangeLog:
* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Insert waits
on kernels reading 'a' into queue 10 before exiting 'a'.
* testsuite/libgomp.oacc-c-c++-common/data-2.c: Ditto.
---
libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c | 4 ++++
libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c | 2 +-
2 files changed, 5 insertions(+), 1 deletion(-)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c
index e9d1edaba7f6..1ad21e955956 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c
@@ -155,6 +155,10 @@ main (int argc, char **argv)
for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+ acc_wait_async (11, 10);
+ acc_wait_async (12, 10);
+ acc_wait_async (13, 10);
+ acc_wait_async (14, 10);
acc_copyout_async (a, nbytes, 10);
acc_copyout_async (b, nbytes, 11);
acc_copyout_async (c, nbytes, 12);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
index 2fc4a598e8f6..c1cc8bd49da6 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
@@ -149,7 +149,7 @@ main (int argc, char **argv)
for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
-#pragma acc exit data copyout (a[0:N]) async (10)
+#pragma acc exit data copyout (a[0:N]) async (10) wait (11) wait (12) wait
(13) wait (14)
#pragma acc exit data copyout (b[0:N]) async (11)
#pragma acc exit data copyout (c[0:N]) async (12)
#pragma acc exit data copyout (d[0:N]) async (13)
--
2.51.2
