diff mbox

[gomp4] Fix acc_on_device for C++

Message ID 568AB69F.4020502@acm.org
State New
Headers show

Commit Message

Nathan Sidwell Jan. 4, 2016, 6:14 p.m. UTC 
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
diff mbox

Patch

2016-01-04  Nathan Sidwell  <nathan@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;
+}