This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[openacc] fix unoptimized acc_on_device
- From: Nathan Sidwell <nathan at acm dot org>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Date: Wed, 6 Jan 2016 10:48:52 -0500
- Subject: [openacc] fix unoptimized acc_on_device
- Authentication-results: sourceware.org; auth=none
I've committed this to trunk. C++ needs a wrapper function to deal with the
enumerated type argument. Usually that's inlined, but when not optimizing we
need to emit a definition of the wrapper. That means marking it as an openacc
routine.
nathan
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.
Index: openacc.h
===================================================================
--- openacc.h (revision 232103)
+++ 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: testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c
===================================================================
--- testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c (revision 0)
+++ 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;
+}