Several OpenACC tests accidentally abuse async semantics, leading to
race conditions & test failures. This patch fixes those tests.
Tested with offloading to AMD GCN. I can probably self-approve this as
a testcase change only, unless anyone objects.
Thanks,
Julian
2021-06-29 Julian Brown <jul...@codesourcery.com>
libgomp/
* testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: Fix async
behaviour and increase number of iterations.
* testsuite/libgomp.oacc-fortran/lib-16-2.f90: Fix async behaviour.
* testsuite/libgomp.oacc-fortran/lib-16.f90: Likewise.
---
.../libgomp.oacc-c-c++-common/deep-copy-10.c | 14 ++++++++------
.../testsuite/libgomp.oacc-fortran/lib-16-2.f90 | 5 +++++
libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90 | 5 +++++
3 files changed, 18 insertions(+), 6 deletions(-)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
index 573a8214bf0..dadb6d37942 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
@@ -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 ();
}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/lib-16-2.f90
b/libgomp/testsuite/libgomp.oacc-fortran/lib-16-2.f90
index ddd557d3be0..e2e47c967fa 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/lib-16-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-16-2.f90
@@ -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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90
b/libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90
index ccd1ce6ee18..ef9a6f6626c 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-16.f90
@@ -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
--
2.29.2
