https://gcc.gnu.org/g:10d9df13f99471702e4d4c7d59096b803dc01e6d
commit r16-5485-g10d9df13f99471702e4d4c7d59096b803dc01e6d Author: Arsen Arsenović <[email protected]> Date: Fri Nov 21 12:17:42 2025 +0100 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: --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c | 14 +++++++++++++- libgomp/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++)
