[gomp4] acc_on_device

[Nathan Sidwell <nathan at acm dot org>]

I've  committed this to gomp4 branch.  It resolves a problem with 
builtin_acc_on_device and C++.  The builtin's arg is  an int, but the 
acc_on_device fn shoud take enum acc_device_t.  In C++ a prototype
 int Foo (enum X);
fails to match up with a builtin of type
 int Foo (int);

We'd worked around this on gomp4 by making the fn decl in openacc.h take an int. 
 This patch resolves things in a different manner, preserving the expected type 
in the header file.

For C, we simply declare it as having enum type, and C matches it up with the 
builtin.   For C++ we declare it as taking an int,  which then matches the 
builtin too.  We also provide an inline forwarding function, taking the enum type.

Because I;m paranoid, I added entries to the enum, to ensure it's layout 
compatible with int.

The test cases in the gcc testsuite were hiding the problem by providing part of 
openacc.h in the test directory, and this  had diverged from the openacc.h we 
actually have.  I deleted those tests and inserted one in the libgomp testsuite, 
which correctly picks up the openacc.h of the tool under test, (rather than one 
in system includes).


nathan

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