Alright, fixed those. I've also noticed the same issue on N in the
interim, so this revision adds those waits also.
Also did the same testing as before.
---------- >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' and 'N', as they read
those.
Previously, the testcases exited 'a' on queue 10 and 'N' on queue 15,
meaning that it was possible for the aforementioned kernels to execute
and to have 'a' and 'N' 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, and the same on
'N'.
libgomp/ChangeLog:
* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c (explanatory
header): Fix typo.
(main): Insert waits on kernels reading 'a' into queue 10 before
exiting 'a', and waits on kernels reading 'N' into queue 15
before exiting 'N'.
* testsuite/libgomp.oacc-c-c++-common/data-2.c: Ditto.
---
.../libgomp.oacc-c-c++-common/data-2-lib.c | 14 +++++++++++++-
.../testsuite/libgomp.oacc-c-c++-common/data-2.c | 9 ++++++---
2 files changed, 19 insertions(+), 4 deletions(-)
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..2af8afc006d7 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
@@ -1,4 +1,4 @@
-/* Test asynchronous, unstructed data regions, runtime library variant. */
+/* Test asynchronous, unstructured data regions, runtime library variant. */
/* See also data-2.c. */
#include <stdlib.h>
@@ -155,11 +155,23 @@ main (int argc, char **argv)
for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+ /* The kernels above use `a', so wait for them to finish with it before
+ exiting that array. */
+ 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);
acc_copyout_async (d, nbytes, 13);
acc_copyout_async (e, nbytes, 14);
+
+ /* As for `a', same goes for `N'. */
+ acc_wait_async (11, 15);
+ acc_wait_async (12, 15);
+ acc_wait_async (13, 15);
+ acc_wait_async (14, 15);
acc_delete_async (&N, sizeof (int), 15);
acc_wait_all ();
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..b974277eecf8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
@@ -1,4 +1,4 @@
-/* Test asynchronous, unstructed data regions, directives variant. */
+/* Test asynchronous, unstructured data regions, directives variant. */
/* See also data-2-lib.c. */
#include <stdlib.h>
@@ -149,12 +149,15 @@ 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)
+ /* The kernels above use `a', so wait for them to finish with it before
+ exiting that array. */
+#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)
#pragma acc exit data copyout (e[0:N]) async (14)
-#pragma acc exit data delete (N) async (15)
+ /* As for `a`, same goes for `N'. */
+#pragma acc exit data delete (N) async (15) wait (11) wait (12) wait (13)
wait (14)
#pragma acc wait
for (i = 0; i < N; i++)
--
2.51.2
