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(-)
diff --git libgomp/ChangeLog libgomp/ChangeLog
index c1f98d76e013..2914066f7532 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,5 +1,9 @@
2018-12-14 Thomas Schwinge <tho...@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 libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index 1e08af70b4da..89b6b6f6fc2b 100644
--- libgomp/oacc-parallel.c
+++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
new file mode 100644
index 000000000000..e4f627d38bc2
--- /dev/null
+++ 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;
+}
Grüße
Thomas
