From: Arsen Arsenović <aarsenovic@baylibre.com>
Date: Fri, 21 Nov 2025 11:17:42 +0000 (+0100)
Subject: libgomp: Fix race condition data-2{,-lib}.c testcase
X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=10d9df13f99471702e4d4c7d59096b803dc01e6d;p=thirdparty%2Fgcc.git

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 e9d1edaba7f..2af8afc006d 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 2fc4a598e8f..b974277eecf 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++)