Fix OpenACC deep-copy-10.c and lib-16-2.f90 test async-safety issues (PR93030)

Message ID 20200103004232.6c46b079@squid.athome
State New
Headers show
Series
  • Fix OpenACC deep-copy-10.c and lib-16-2.f90 test async-safety issues (PR93030)
Related show

Commit Message

Julian Brown Jan. 3, 2020, 12:42 a.m.
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.

Patch

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.

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