This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

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


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  <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 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


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]