This patch fixes acc_on_device's C++ wrapper when compiling at -O0. The wrapper
isn't inlined, and we need to mark the function as needing emission by the
device compiler too.
nathan
2016-01-04 Nathan Sidwell <nat...@codesourcery.com>
* openacc.c (acc_on_device): Add routine pragma for C++ wrapper.
* testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c: New.
Index: libgomp/openacc.h
===================================================================
--- libgomp/openacc.h (revision 232058)
+++ libgomp/openacc.h (working copy)
@@ -121,6 +121,7 @@ int acc_set_cuda_stream (int, void *) __
/* 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);
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c (working copy)
@@ -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;
+}
