@@ -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.
@@ -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
new file mode 100644
@@ -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;
+}
Hi! Committed to trunk in r267151: commit 44b7d2b9c1b1a535212b8312c6dc76dd1570db45 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Fri Dec 14 20:42:50 2018 +0000 [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. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@267151 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 4 ++ libgomp/oacc-parallel.c | 4 +- .../libgomp.oacc-c-c++-common/asyncwait-nop-1.c | 78 ++++++++++++++++++++++ 3 files changed, 84 insertions(+), 2 deletions(-) Grüße Thomas