2015-10-29  Nathan Sidwell  <nathan@codesourcery.com>

	libgomp/
	* openacc.h (enum acc_device_t): Ensure layout compatibility.
	(acc_on_device): Declare compatible with builtin and provide C++
	wrapper.
	* oacc-init.c (acc_on_device): Change arg type.
	* config/nvptx/oacc-init.c (acc_on_device): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc-on-device.c: New.

	gcc/testsuite/
	* c-c++-common/goacc/acc_on_device-2-off.c: Delete.
	* c-c++-common/goacc/acc_on_device-2.c: Delete.
	* c-c++-common/goacc/kernels-acc-on-device-2.c: Delete.
	* c-c++-common/goacc/kernels-acc-on-device.c: Delete.
	* c-c++-common/goacc/openacc.h: Delete.

Index: libgomp/oacc-init.c
===================================================================
--- libgomp/oacc-init.c	(revision 229502)
+++ libgomp/oacc-init.c	(working copy)
@@ -646,7 +646,7 @@ ialias (acc_set_device_num)
    this, rather than generating infinitely recursive code.  */
 
 int __attribute__ ((__optimize__ ("O2")))
-acc_on_device (int dev)
+acc_on_device (acc_device_t dev)
 {
   return __builtin_acc_on_device (dev);
 }
Index: libgomp/openacc.h
===================================================================
--- libgomp/openacc.h	(revision 229502)
+++ libgomp/openacc.h	(working copy)
@@ -56,7 +56,10 @@ typedef enum acc_device_t
     /* acc_device_host_nonshm = 3 removed.  */
     acc_device_not_host = 4,
     acc_device_nvidia = 5,
-    _ACC_device_hwm
+    _ACC_device_hwm,
+    /* Ensure enumeration is layout compatible with int.  */
+    _ACC_highest = __INT_MAX__,
+    _ACC_neg = -1
   } acc_device_t;
 
 typedef enum acc_async_t
@@ -79,11 +82,11 @@ void acc_wait_all (void) __GOACC_NOTHROW
 void acc_wait_all_async (int) __GOACC_NOTHROW;
 void acc_init (acc_device_t) __GOACC_NOTHROW;
 void acc_shutdown (acc_device_t) __GOACC_NOTHROW;
-/* Library function declaration.  Although it should take an
-   acc_device_t argument, that causes problems with matching the
-   builtin, which takes an int (to avoid declaring the enumeration
-   inside the compiler).  */
-int acc_on_device (int) __GOACC_NOTHROW;
+#ifdef __cplusplus
+int acc_on_device (int __arg) __GOACC_NOTHROW;
+#else
+int acc_on_device (acc_device_t __arg) __GOACC_NOTHROW;
+#endif
 void *acc_malloc (size_t) __GOACC_NOTHROW;
 void acc_free (void *) __GOACC_NOTHROW;
 /* Some of these would be more correct with const qualifiers, but
@@ -117,6 +120,10 @@ int acc_set_cuda_stream (int, void *) __
 
 #ifdef __cplusplus
 }
+inline int acc_on_device (acc_device_t __arg) __GOACC_NOTHROW
+{
+  return acc_on_device ((int) __arg);
+}
 #endif
 
 #endif /* _OPENACC_H */
Index: libgomp/config/nvptx/oacc-init.c
===================================================================
--- libgomp/config/nvptx/oacc-init.c	(revision 229502)
+++ libgomp/config/nvptx/oacc-init.c	(working copy)
@@ -36,7 +36,7 @@
    this, rather than generating infinitely recursive code.  */
 
 int __attribute__ ((__optimize__ ("O2")))
-acc_on_device (int dev)
+acc_on_device (acc_device_t dev)
 {
   return __builtin_acc_on_device (dev);
 }
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device.c	(revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device.c	(working copy)
@@ -0,0 +1,11 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-O2" } */
+
+#include <openacc.h>
+
+int Foo (acc_device_t x)
+{
+  return acc_on_device (x);
+}
+
+/* { dg-final { scan-assembler-not "acc_on_device" } } */
Index: gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c	(revision 229535)
+++ gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c	(working copy)
@@ -1,17 +0,0 @@
-/* Have to enable optimizations, as otherwise builtins won't be expanded.  */
-/* { dg-additional-options "-O -fdump-rtl-expand -fno-openacc" } */
-
-/* Duplicate parts of libgomp/openacc.h, because we can't include it here.  */
-
-#include "openacc.h"
-
-int
-f (void)
-{
-  const acc_device_t dev = acc_device_X;
-  return acc_on_device (dev);
-}
-
-/* Without -fopenacc, we're expecting one call.
-   { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]* acc_on_device" 1 "expand" } } */
-
Index: gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c	(revision 229535)
+++ gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c	(working copy)
@@ -1,19 +0,0 @@
-/* Have to enable optimizations, as otherwise builtins won't be expanded.  */
-/* { dg-additional-options "-O -fdump-rtl-expand" } */
-
-#include "openacc.h"
-
-int
-f (void)
-{
-  const acc_device_t dev = acc_device_X;
-  return acc_on_device (dev);
-}
-
-/* With -fopenacc, we're expecting the builtin to be expanded, so no calls.
-   TODO: in C++, the use of enum acc_device_t for acc_on_device's parameter
-   perturbs expansion as a builtin, which expects an int parameter.  It's fine
-   when changing acc_device_t to plain int, but that's not necessarily what a
-   user will be doing.
-
-   { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]* acc_on_device" 0 "expand" { xfail c++ } } } */
Index: gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c	(revision 229535)
+++ gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c	(working copy)
@@ -1,37 +0,0 @@
-/* { dg-additional-options "-O2" } */
-
-#include "openacc.h"
-
-#define N 32
-
-void
-foo (float *a, float *b)
-{
-#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
-  {
-    int ii;
-    int on_host = acc_on_device (acc_device_X);
-
-    for (ii = 0; ii < N; ii++)
-      {
-	if (on_host)
-	  b[ii] = a[ii] + 1;
-	else
-	  b[ii] = a[ii];
-      }
-  }
-
-#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
-  {
-    int ii;
-    int on_host = acc_on_device (acc_device_X);
-
-    for (ii = 0; ii < N; ii++)
-      {
-	if (on_host)
-	  b[ii] = a[ii] + 2;
-	else
-	  b[ii] = a[ii];
-      }
-  }
-}
Index: gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c	(revision 229535)
+++ gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c	(working copy)
@@ -1,35 +0,0 @@
-/* { dg-additional-options "-O2" } */
-
-#include "openacc.h"
-
-#define N 32
-
-void
-foo (float *a, float *b)
-{
-#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
-  {
-    int ii;
-
-    for (ii = 0; ii < N; ii++)
-      {
-	if (acc_on_device (acc_device_X))
-	  b[ii] = a[ii] + 1;
-	else
-	  b[ii] = a[ii];
-      }
-  }
-
-#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
-  {
-    int ii;
-
-    for (ii = 0; ii < N; ii++)
-      {
-	if (acc_on_device (acc_device_X))
-	  b[ii] = a[ii] + 2;
-	else
-	  b[ii] = a[ii];
-      }
-  }
-}
Index: gcc/testsuite/c-c++-common/goacc/openacc.h
===================================================================
--- gcc/testsuite/c-c++-common/goacc/openacc.h	(revision 229535)
+++ gcc/testsuite/c-c++-common/goacc/openacc.h	(working copy)
@@ -1,18 +0,0 @@
-#if __cplusplus
-extern "C" {
-#endif
-
-#if __cplusplus >= 201103
-# define __GOACC_NOTHROW noexcept
-#elif __cplusplus
-# define __GOACC_NOTHROW throw ()
-#else /* Not C++ */
-# define __GOACC_NOTHROW __attribute__ ((__nothrow__))
-#endif
-
-typedef enum acc_device_t { acc_device_X = 123 } acc_device_t;
-int acc_on_device (int) __GOACC_NOTHROW;
-
-#if __cplusplus
-}
-#endif
