libgomp: Fix race condition data-2{,-lib}.c testcase

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.

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 e9d1edaba7f6fddf2e3bc42bdb667d3580aba65e..2af8afc006d7f58a163907c05299aa439da6d0e5 100644 (file)
--- 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 2fc4a598e8f6b64e6f3c44f3696a04c1f8d8efa8..b974277eecf8ed2e8a04048e7199e6e2a3ca3e04 100644 (file)
--- 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++)