commit 990d954e84ff393e4d07661d82979e9acce6c9a0
Author: Julian Brown <julian@codesourcery.com>
Date: Thu Jan 2 16:14:17 2020 -0800
Fix deep-copy-10.c and lib-16-2.f90 async-safety issues
PR libgomp/93030
libgomp/
* testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: Fix unsafe
async usage.
* testsuite/libgomp.oacc-fortran/lib-16-2.f90: Add async waits.
@@ -1,6 +1,8 @@
#include <stdlib.h>
-/* Test asyncronous attach and detach operation. */
+#define ITERATIONS 1023
+
+/* Test asynchronous attach and detach operation. */
typedef struct {
int *a;
@@ -25,13 +27,13 @@ main (int argc, char* argv[])
#pragma acc enter data copyin(m)
- for (int i = 0; i < 99; i++)
+ for (int i = 0; i < ITERATIONS; i++)
{
int j;
-#pragma acc parallel loop copy(m.a[0:N]) async(i % 2)
+#pragma acc parallel loop copy(m.a[0:N]) async(0)
for (j = 0; j < N; j++)
m.a[j]++;
-#pragma acc parallel loop copy(m.b[0:N]) async((i + 1) % 2)
+#pragma acc parallel loop copy(m.b[0:N]) async(1)
for (j = 0; j < N; j++)
m.b[j]++;
}
@@ -40,9 +42,9 @@ main (int argc, char* argv[])
for (i = 0; i < N; i++)
{
- if (m.a[i] != 99)
+ if (m.a[i] != ITERATIONS)
abort ();
- if (m.b[i] != 99)
+ if (m.b[i] != ITERATIONS)
abort ();
}
@@ -27,6 +27,9 @@ program main
if (acc_is_present (h) .neqv. .TRUE.) stop 1
+ ! We must wait for the update to be done.
+ call acc_wait (async)
+
h(:) = 0
call acc_copyout_async (h, sizeof (h), async)
@@ -45,6 +48,8 @@ program main
if (acc_is_present (h) .neqv. .TRUE.) stop 3
+ call acc_wait (async)
+
do i = 1, N
if (h(i) /= i + i) stop 4
end do
Hi, This patch fixes some unsafe usage of asynchronous operations in a couple of OpenACC tests. I have extracted it from the og9 version of the patch posted here concerning "ephemeral" asynchronous host-to-device copies: https://gcc.gnu.org/ml/gcc-patches/2019-11/msg01482.html The deep-copy-10.c changes weren't part of that submission because manual deep-copy support hadn't landed on trunk yet, but the lib-16-2.f90 changes were posted in the same series here: https://gcc.gnu.org/ml/gcc-patches/2019-11/msg01485.html Anyway: these fixes are probably desirable independent from the middle-end worker-partitioning support patches, so can be committed separately. The deep-copy-10.c test won't actually pass on AMD GCN without the ephemeral async host-to-device copy patch, though. OK? Thanks, Julian ChangeLog PR libgomp/93030 libgomp/ * testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: Fix unsafe async usage. * testsuite/libgomp.oacc-fortran/lib-16-2.f90: Add async waits.