[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

3 files changed, 84 insertions, 2 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index c1f98d7..2914066 100644
--- 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 1e08af7..89b6b6f 100644
--- 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
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;
+}