openacc.c (acc_on_device): Add routine pragma for C++ wrapper.

	* openacc.c (acc_on_device): Add routine pragma for C++ wrapper.
	* testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c: New.

From-SVN: r232104

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 054e1b9..b958313 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2016-01-06  Nathan Sidwell  <nathan@acm.org>
+
+       * openacc.c (acc_on_device): Add routine pragma for C++ wrapper.
+       * testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c: New.
+
 2016-01-04  Jakub Jelinek  <jakub@redhat.com>
 
        Update copyright years.

diff --git a/libgomp/openacc.h b/libgomp/openacc.h
index 4595530..7ea8794 100644 (file)
--- a/libgomp/openacc.h
+++ b/libgomp/openacc.h
@@ -121,6 +121,7 @@ int acc_set_cuda_stream (int, void *) __GOACC_NOTHROW;
 
 /* Forwarding function with correctly typed arg.  */
 
+#pragma acc routine seq
 inline int acc_on_device (acc_device_t __arg) __GOACC_NOTHROW
 {
   return acc_on_device ((int) __arg);

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c
new file mode 100644 (file)
index 0000000..e5d9c36
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c
@@ -0,0 +1,23 @@
+/* { dg-additional-options "-O0" } */
+
+#include <openacc.h>
+
+/* acc_on_device might not be folded at -O0, but it should work. */
+
+int main ()
+{
+  int dev;
+  
+#pragma acc parallel copyout (dev)
+  {
+    dev = acc_on_device (acc_device_not_host);
+  }
+
+  int expect = 1;
+  
+#if  ACC_DEVICE_TYPE_host
+  expect = 0;
+#endif
+  
+  return dev != expect;
+}