aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorNathan Sidwell <nathan@acm.org>2016-01-06 15:47:14 +0000
committerNathan Sidwell <nathan@gcc.gnu.org>2016-01-06 15:47:14 +0000
commite9b596c13252df2f1cea3f27183e9969ede54fbd (patch)
tree945fd4e724f1c10ee1609968b449d1da268975fe
parent7c961392a68dcf032ec75aa38a41b6ec7d8ec477 (diff)
downloadgcc-e9b596c13252df2f1cea3f27183e9969ede54fbd.zip
gcc-e9b596c13252df2f1cea3f27183e9969ede54fbd.tar.gz
gcc-e9b596c13252df2f1cea3f27183e9969ede54fbd.tar.bz2
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
-rw-r--r--libgomp/ChangeLog5
-rw-r--r--libgomp/openacc.h1
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c23
3 files changed, 29 insertions, 0 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 054e1b9..b958313 100644
--- 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
--- 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
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;
+}