[PR88484] OpenACC wait directive without wait argument but with async clause

We don't correctly handle "#pragma acc wait async (a)" for "a >= 0", handling
as a no-op whereas it should enqueue the appropriate wait operations on
"async (a)".

	libgomp/
	PR libgomp/88484
	* oacc-parallel.c (GOACC_wait): Correct handling for "async >= 0".
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: New file.

From-SVN: r267151

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index c1f98d76e0137e47760f94c949c9f6a18a475133..2914066f7532833e241c232f181f1159fca4d74d 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,9 @@
 2018-12-14  Thomas Schwinge  <thomas@codesourcery.com>
 
+       PR libgomp/88484
+       * oacc-parallel.c (GOACC_wait): Correct handling for "async >= 0".
+       * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: New file.
+
        PR libgomp/88407
        * plugin/plugin-nvptx.c (nvptx_async_test, nvptx_wait)
        (nvptx_wait_async): Unseen async-argument is a no-op.

diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 1e08af70b4daaceab1870bc0f7f3fa25694ba520..89b6b6f6fc2b1c3b092617404ef41ea38808fbda 100644 (file)
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -630,8 +630,8 @@ GOACC_wait (int async, int num_waits, ...)
     }
   else if (async == acc_async_sync)
     acc_wait_all ();
-  else if (async == acc_async_noval)
-    goacc_thread ()->dev->openacc.async_wait_all_async_func (acc_async_noval);
+  else
+    acc_wait_all_async (async);
 }
 
 int

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
new file mode 100644 (file)
index 0000000..e4f627d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
@@ -0,0 +1,78 @@
+/* Several of the async/wait combinations invoked here are no-ops -- they don't
+   effect anything, but are still valid.
+
+   This doesn't verify that the asynchronous operations synchronize correctly,
+   but just verifies that we don't refuse any variants.  */
+
+#undef NDEBUG
+#include <assert.h>
+#include <openacc.h>
+
+int values[] = { acc_async_sync,
+                acc_async_noval,
+                0,
+                1,
+                2,
+                36,
+                1982, };
+const size_t values_n = sizeof values / sizeof values[0];
+
+int
+main ()
+{
+  /* Explicitly initialize: it's not clear whether the following OpenACC
+     runtime library calls implicitly initialize;
+     <https://github.com/OpenACC/openacc-spec/issues/102>.  */
+  acc_device_t d;
+#if defined ACC_DEVICE_TYPE_nvidia
+  d = acc_device_nvidia;
+#elif defined ACC_DEVICE_TYPE_host
+  d = acc_device_host;
+#else
+# error Not ported to this ACC_DEVICE_TYPE
+#endif
+  acc_init (d);
+
+
+  for (size_t i = 0; i < values_n; ++i)
+    assert (acc_async_test (values[i]) == 1);
+
+
+  for (size_t i = 0; i < values_n; ++i)
+    {
+#pragma acc parallel wait (values[i])
+      ;
+#pragma acc wait (values[i])
+      acc_wait (values[i]);
+    }
+
+
+  for (size_t i = 0; i < values_n; ++i)
+    {
+      for (size_t j = 0; j < values_n; ++j)
+       {
+         if (values[i] == values[j])
+           continue;
+
+#pragma acc parallel wait (values[i]) async (values[j])
+         ;
+#pragma acc wait (values[i]) async (values[j])
+         acc_wait_async (values[i], values[j]);
+       }
+    }
+
+
+  for (size_t i = 0; i < values_n; ++i)
+    {
+#pragma acc parallel wait async (values[i])
+      ;
+#pragma acc wait async (values[i])
+      acc_wait_all_async (values[i]);
+    }
+
+
+  /* Clean up.  */
+  acc_wait_all ();
+
+  return 0;
+}