[gomp4] OpenACC acc_on_device (was: various OpenACC/PTX built-ins and a reduction tweak)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index ef9a81d..0c25a27 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,11 @@ 
+2014-09-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* builtins.def (DEF_GOACC_BUILTIN_COMPILER): New macro.
+	* oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin.
+	* builtins.c (expand_builtin_acc_on_device): New function.
+	(expand_builtin): Use it to handle BUILT_IN_ACC_ON_DEVICE.
+	(is_inexpensive_builtin): Handle BUILT_IN_ACC_ON_DEVICE.
+
 2014-09-08  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* configure.ac (offload_targets): Remove.
diff --git gcc/builtins.c gcc/builtins.c
index 975f696..5b2ebcc 100644
--- gcc/builtins.c
+++ gcc/builtins.c
@@ -5747,6 +5747,49 @@  expand_stack_save (void)
   return ret;
 }
 
+
+/* Expand OpenACC acc_on_device.
+
+   This has to happen late (that is, not in early folding; expand_builtin_*,
+   rather than fold_builtin_*), as we have to act differently for host and
+   acceleration device (ACCEL_COMPILER conditional).  */
+
+static rtx
+expand_builtin_acc_on_device (tree exp, rtx target ATTRIBUTE_UNUSED)
+{
+  if (!validate_arglist (exp, INTEGER_TYPE, VOID_TYPE))
+    return NULL_RTX;
+
+  tree arg, v1, v2, ret;
+  location_t loc;
+
+  arg = CALL_EXPR_ARG (exp, 0);
+  arg = builtin_save_expr (arg);
+  loc = EXPR_LOCATION (exp);
+
+  /* Build: (arg == v1 || arg == v2) ? 1 : 0.  */
+
+#ifdef ACCEL_COMPILER
+  v1 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_not_host */ 3);
+  v2 = build_int_cst (TREE_TYPE (arg), ACCEL_COMPILER_acc_device);
+#else
+  v1 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_none */ 0);
+  v2 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_host */ 2);
+#endif
+
+  v1 = fold_build2_loc (loc, EQ_EXPR, integer_type_node, arg, v1);
+  v2 = fold_build2_loc (loc, EQ_EXPR, integer_type_node, arg, v2);
+
+  /* Can't use TRUTH_ORIF_EXPR, as that is not supported by
+     expand_expr_real*.  */
+  ret = fold_build3_loc (loc, COND_EXPR, integer_type_node, v1, v1, v2);
+  ret = fold_build3_loc (loc, COND_EXPR, integer_type_node,
+			 ret, integer_one_node, integer_zero_node);
+
+  return expand_normal (ret);
+}
+
+
 /* Expand an expression EXP that calls a built-in function,
    with result going to TARGET if that's convenient
    (and in mode MODE if that's convenient).
@@ -6816,6 +6859,12 @@  expand_builtin (tree exp, rtx target, rtx subtarget, enum machine_mode mode,
       expand_builtin_cilk_pop_frame (exp);
       return const0_rtx;
 
+    case BUILT_IN_ACC_ON_DEVICE:
+      target = expand_builtin_acc_on_device (exp, target);
+      if (target)
+	return target;
+      break;
+
     default:	/* just do library call, if unknown builtin */
       break;
     }
@@ -12748,6 +12797,7 @@  is_inexpensive_builtin (tree decl)
       case BUILT_IN_LABS:
       case BUILT_IN_LLABS:
       case BUILT_IN_PREFETCH:
+      case BUILT_IN_ACC_ON_DEVICE:
 	return true;
 
       default:
diff --git gcc/builtins.def gcc/builtins.def
index 2ef896e..b9b8e74 100644
--- gcc/builtins.def
+++ gcc/builtins.def
@@ -146,12 +146,16 @@  along with GCC; see the file COPYING3.  If not see
   DEF_BUILTIN (ENUM, NAME, BUILT_IN_NORMAL, BT_LAST, BT_LAST, false, false, \
 	       false, ATTR_LAST, false, false)
 
-/* Builtin used by the implementation of GNU OpenACC.  None of these are
-   actually implemented in the compiler; they're all in libgomp.  */
+/* Builtin used by the implementation of GNU OpenACC.  Few of these are
+   actually implemented in the compiler; most are in libgomp.  */
 #undef DEF_GOACC_BUILTIN
 #define DEF_GOACC_BUILTIN(ENUM, NAME, TYPE, ATTRS) \
   DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE,    \
                false, true, true, ATTRS, false, flag_openacc)
+#undef DEF_GOACC_BUILTIN_COMPILER
+#define DEF_GOACC_BUILTIN_COMPILER(ENUM, NAME, TYPE, ATTRS) \
+  DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE,    \
+               flag_openacc, true, true, ATTRS, false, true)
 
 /* Builtin used by the implementation of GNU OpenMP.  None of these are
    actually implemented in the compiler; they're all in libgomp.  */
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index cc7c888..df86db7 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,8 @@ 
+2014-09-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* f95-lang.c (DEF_GOACC_BUILTIN_COMPILER): New macro.
+	* types.def (BT_FN_INT_INT): New type.
+
 2014-09-08  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* gfortran.h (enum OMP_LIST_FIRST, OMP_LIST_LAST): New
diff --git gcc/fortran/f95-lang.c gcc/fortran/f95-lang.c
index e7c64b7..1b017b1 100644
--- gcc/fortran/f95-lang.c
+++ gcc/fortran/f95-lang.c
@@ -1093,7 +1093,12 @@  gfc_init_builtin_functions (void)
 #define DEF_GOACC_BUILTIN(code, name, type, attr) \
       gfc_define_builtin ("__builtin_" name, builtin_types[type], \
 			  code, name, attr);
+#undef DEF_GOACC_BUILTIN_COMPILER
+      /* TODO: this is not doing the right thing.  */
+#define DEF_GOACC_BUILTIN_COMPILER(code, name, type, attr) \
+      gfc_define_builtin (name, builtin_types[type], code, name, attr);
 #include "../oacc-builtins.def"
+#undef DEF_GOACC_BUILTIN_COMPILER
 #undef DEF_GOACC_BUILTIN
     }
 
diff --git gcc/fortran/types.def gcc/fortran/types.def
index 59ac4c3..6c2fdc0 100644
--- gcc/fortran/types.def
+++ gcc/fortran/types.def
@@ -82,6 +82,7 @@  DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_INT_INT, BT_INT, BT_INT)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
 DEF_FUNCTION_TYPE_1 (BT_FN_PTR_PTR, BT_PTR, BT_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_INT, BT_VOID, BT_INT)
diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def
index dfb688c..e4bc756 100644
--- gcc/oacc-builtins.def
+++ gcc/oacc-builtins.def
@@ -39,3 +39,5 @@  DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
 		   ATTR_NOTHROW_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update",
 		   BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device",
+			    BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 4427521..e210c6b 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,12 @@ 
+2014-09-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/acc_on_device-1.c: New file.
+	* c-c++-common/goacc/acc_on_device-2.c: Likewise.
+	* c-c++-common/goacc/acc_on_device-2-off.c: Likewise.
+	* gfortran.dg/goacc/acc_on_device-1.f95: Likewise.
+	* gfortran.dg/goacc/acc_on_device-2.f95: Likewise.
+	* gfortran.dg/goacc/acc_on_device-2-off.f95: Likewise.
+
 2014-09-08  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* gfortran.dg/goacc/private-1.f95: New test.
diff --git gcc/testsuite/c-c++-common/goacc/acc_on_device-1.c gcc/testsuite/c-c++-common/goacc/acc_on_device-1.c
new file mode 100644
index 0000000..d0e137b
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/acc_on_device-1.c
@@ -0,0 +1,20 @@ 
+/* Have to enable optimizations, as otherwise builtins won't be expanded.  */
+/* { dg-additional-options "-O -fdump-rtl-expand" } */
+
+int
+f (void)
+{
+  int r = 0;
+
+  r |= acc_on_device ();
+  r |= acc_on_device (1, 2);
+  r |= acc_on_device (3.14);
+  r |= acc_on_device ("hello");
+
+  return r;
+}
+
+/* Unsuitable to be handled as a builtin, so we're expecting four calls.
+   { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 4 "expand" } } */
+
+/* { dg-final { cleanup-rtl-dump "expand" } } */
diff --git gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c
new file mode 100644
index 0000000..ddc43ab
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c
@@ -0,0 +1,17 @@ 
+/* Have to enable optimizations, as otherwise builtins won't be expanded.  */
+/* { dg-additional-options "-O -fdump-rtl-expand -fno-openacc" } */
+
+typedef enum acc_device_t { acc_device_X = 123 } acc_device_t;
+extern int acc_on_device (acc_device_t);
+
+int
+f (void)
+{
+  const int 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" } } */
+
+/* { dg-final { cleanup-rtl-dump "expand" } } */
diff --git gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c
new file mode 100644
index 0000000..65b4ae6
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c
@@ -0,0 +1,17 @@ 
+/* Have to enable optimizations, as otherwise builtins won't be expanded.  */
+/* { dg-additional-options "-O -fdump-rtl-expand" } */
+
+typedef enum acc_device_t { acc_device_X = 123 } acc_device_t;
+extern int acc_on_device (acc_device_t);
+
+int
+f (void)
+{
+  const int dev = acc_device_X;
+  return acc_on_device (dev);
+}
+
+/* With -fopenacc, we're expecting the builtin to be expanded, so no calls.
+   { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 0 "expand" } } */
+
+/* { dg-final { cleanup-rtl-dump "expand" } } */
diff --git gcc/testsuite/gfortran.dg/goacc/acc_on_device-1.f95 gcc/testsuite/gfortran.dg/goacc/acc_on_device-1.f95
new file mode 100644
index 0000000..9dfde26
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/acc_on_device-1.f95
@@ -0,0 +1,22 @@ 
+! Have to enable optimizations, as otherwise builtins won't be expanded.
+! { dg-additional-options "-O -fdump-rtl-expand" }
+
+logical function f ()
+  implicit none
+
+  external acc_on_device
+  logical (4) acc_on_device
+
+  f = .false.
+  f = f .or. acc_on_device ()
+  f = f .or. acc_on_device (1, 2)
+  f = f .or. acc_on_device (3.14)
+  f = f .or. acc_on_device ("hello")
+
+  return
+end function f
+
+! Unsuitable to be handled as a builtin, so we're expecting four calls.
+! { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 4 "expand" } }
+
+! { dg-final { cleanup-rtl-dump "expand" } }
diff --git gcc/testsuite/gfortran.dg/goacc/acc_on_device-2-off.f95 gcc/testsuite/gfortran.dg/goacc/acc_on_device-2-off.f95
new file mode 100644
index 0000000..cf28264
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/acc_on_device-2-off.f95
@@ -0,0 +1,39 @@ 
+! Have to enable optimizations, as otherwise builtins won't be expanded.
+! { dg-additional-options "-O -fdump-rtl-expand -fno-openacc" }
+
+module openacc_kinds
+  implicit none
+
+  integer, parameter :: acc_device_kind = 4
+
+end module openacc_kinds
+
+module openacc
+  use openacc_kinds
+  implicit none
+
+  integer (acc_device_kind), parameter :: acc_device_host = 2
+
+  interface
+     function acc_on_device (dev)
+       use openacc_kinds
+       logical (4) :: acc_on_device
+       integer (acc_device_kind), intent (in) :: dev
+     end function acc_on_device
+  end interface
+end module openacc
+
+logical (4) function f ()
+  use openacc
+  implicit none
+
+  integer (4), parameter :: dev = 2
+
+  f = acc_on_device (dev)
+  return
+end function f
+
+! Without -fopenacc, we're expecting one call.
+! { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 1 "expand" } }
+
+! { dg-final { cleanup-rtl-dump "expand" } }
diff --git gcc/testsuite/gfortran.dg/goacc/acc_on_device-2.f95 gcc/testsuite/gfortran.dg/goacc/acc_on_device-2.f95
new file mode 100644
index 0000000..7730a60
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/acc_on_device-2.f95
@@ -0,0 +1,40 @@ 
+! Have to enable optimizations, as otherwise builtins won't be expanded.
+! { dg-additional-options "-O -fdump-rtl-expand" }
+
+module openacc_kinds
+  implicit none
+
+  integer, parameter :: acc_device_kind = 4
+
+end module openacc_kinds
+
+module openacc
+  use openacc_kinds
+  implicit none
+
+  integer (acc_device_kind), parameter :: acc_device_host = 2
+
+  interface
+     function acc_on_device (dev)
+       use openacc_kinds
+       logical (4) :: acc_on_device
+       integer (acc_device_kind), intent (in) :: dev
+     end function acc_on_device
+  end interface
+end module openacc
+
+logical (4) function f ()
+  use openacc
+  implicit none
+
+  integer (4), parameter :: dev = 2
+
+  f = acc_on_device (dev)
+  return
+end function f
+
+! With -fopenacc, we're expecting the builtin to be expanded, so no calls.
+! TODO: not working.
+! { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 0 "expand" { xfail *-*-* } } }
+
+! { dg-final { cleanup-rtl-dump "expand" } }
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 5b2a39d..8d774ee 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,25 @@ 
+2014-09-18  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* libgomp.map (OACC_2.0): Add acc_on_device, acc_on_device_.
+	* fortran.c: Include "openacc.h".
+	(acc_on_device_): New function.
+	* oacc-parallel.c: Include "openacc.h".
+	(acc_on_device): New function.
+	* openacc.f90 (acc_device_kind, acc_device_none)
+	(acc_device_default, acc_device_host, acc_device_not_host): New
+	parameters.
+	(acc_on_device): New function declaration.
+	* openacc_lib.h (acc_device_kind, acc_device_none)
+	(acc_device_default, acc_device_host, acc_device_not_host): New
+	parameters.
+	(acc_on_device): New function declaration.
+	* openacc.h (acc_device_t): New enum.
+	(acc_on_device): New function declaration.
+	* testsuite/libgomp.oacc-c/acc_on_device-1.c: New file.
+	* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
+	* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
+	* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
+
 2014-07-09  Thomas Schwinge  <thomas@codesourcery.com>
 	    Jakub Jelinek  <jakub@redhat.com>
 
diff --git libgomp/fortran.c libgomp/fortran.c
index 1f30c51..28c83cc 100644
--- libgomp/fortran.c
+++ libgomp/fortran.c
@@ -26,6 +26,7 @@ 
 
 #include "libgomp.h"
 #include "libgomp_f.h"
+#include "openacc.h"
 #include <stdlib.h>
 #include <limits.h>
 
@@ -73,6 +74,7 @@  ialias_redirect (omp_get_num_devices)
 ialias_redirect (omp_get_num_teams)
 ialias_redirect (omp_get_team_num)
 ialias_redirect (omp_is_initial_device)
+ialias_redirect (acc_on_device)
 #endif
 
 #ifndef LIBGOMP_GNU_SYMBOL_VERSIONING
@@ -492,3 +494,9 @@  omp_is_initial_device_ (void)
 {
   return omp_is_initial_device ();
 }
+
+int32_t
+acc_on_device_ (const int32_t *dev)
+{
+  return acc_on_device (*dev);
+}
diff --git libgomp/libgomp.map libgomp/libgomp.map
index c575be3..69a4d83 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -234,6 +234,9 @@  GOMP_4.0.1 {
 } GOMP_4.0;
 
 OACC_2.0 {
+  global:
+	acc_on_device;
+	acc_on_device_;
 };
 
 GOACC_2.0 {
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index 79b6254..02fbb12 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -27,6 +27,7 @@ 
 
 #include "libgomp.h"
 #include "libgomp_g.h"
+#include "openacc.h"
 
 void
 GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target,
@@ -128,3 +129,12 @@  GOACC_update (int device, const void *openmp_target, size_t mapnum,
     }
   GOMP_target_update (device, openmp_target, mapnum, hostaddrs, sizes, kinds_);
 }
+
+/* TODO: Move elsewhere.  */
+int
+acc_on_device (acc_device_t dev)
+{
+  /* Just rely on the compiler builtin.  */
+  return __builtin_acc_on_device (dev);
+}
+ialias (acc_on_device)
diff --git libgomp/openacc.f90 libgomp/openacc.f90
index b2a79f6..70b58d6 100644
--- libgomp/openacc.f90
+++ libgomp/openacc.f90
@@ -1,6 +1,6 @@ 
 !  OpenACC Runtime Library Definitions.
 
-!  Copyright (C) 2013 Free Software Foundation, Inc.
+!  Copyright (C) 2013-2014 Free Software Foundation, Inc.
 
 !  Contributed by Thomas Schwinge <thomas@codesourcery.com>.
 
@@ -28,6 +28,8 @@ 
 module openacc_kinds
   implicit none
 
+  integer, parameter :: acc_device_kind = 4
+
 end module openacc_kinds
 
 module openacc
@@ -36,4 +38,17 @@  module openacc
 
   integer, parameter :: openacc_version = 201306
 
+  integer (acc_device_kind), parameter :: acc_device_none = 0
+  integer (acc_device_kind), parameter :: acc_device_default = 1
+  integer (acc_device_kind), parameter :: acc_device_host = 2
+  integer (acc_device_kind), parameter :: acc_device_not_host = 3
+
+  interface
+     function acc_on_device (dev)
+       use openacc_kinds
+       logical (4) :: acc_on_device
+       integer (acc_device_kind), intent (in) :: dev
+     end function acc_on_device
+  end interface
+
 end module openacc
diff --git libgomp/openacc.h libgomp/openacc.h
index a6f7ec94..cde7429 100644
--- libgomp/openacc.h
+++ libgomp/openacc.h
@@ -37,7 +37,18 @@  extern "C" {
 #else
 # define __GOACC_NOTHROW __attribute__ ((__nothrow__))
 #endif
-  
+
+typedef enum acc_device_t
+  {
+    acc_device_none = 0,
+    acc_device_default, /* This has to be a distinct value, as no
+			   return value can match it.  */
+    acc_device_host = 2,
+    acc_device_not_host = 3
+  } acc_device_t;
+
+int acc_on_device (acc_device_t __dev) __GOACC_NOTHROW;
+
 #ifdef __cplusplus
 }
 #endif
diff --git libgomp/openacc_lib.h libgomp/openacc_lib.h
index d19c95c..be49100 100644
--- libgomp/openacc_lib.h
+++ libgomp/openacc_lib.h
@@ -1,6 +1,6 @@ 
 !  OpenACC Runtime Library Definitions.                   -*- mode: fortran -*-
 
-!  Copyright (C) 2013 Free Software Foundation, Inc.
+!  Copyright (C) 2013-2014 Free Software Foundation, Inc.
 
 !  Contributed by Thomas Schwinge <thomas@codesourcery.com>.
 
@@ -27,3 +27,17 @@ 
 
       integer openacc_version
       parameter (openacc_version = 201306)
+
+      integer acc_device_kind
+      parameter (acc_device_kind = 4)
+      integer (acc_device_kind) acc_device_none
+      parameter (acc_device_none = 0)
+      integer (acc_device_kind) acc_device_default
+      parameter (acc_device_default = 1)
+      integer (acc_device_kind) acc_device_host
+      parameter (acc_device_host = 2)
+      integer (acc_device_kind) acc_device_not_host
+      parameter (acc_device_not_host = 3)
+
+      external acc_on_device
+      logical (4) acc_on_device
diff --git libgomp/testsuite/libgomp.oacc-c/acc_on_device-1.c libgomp/testsuite/libgomp.oacc-c/acc_on_device-1.c
new file mode 100644
index 0000000..f216587
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/acc_on_device-1.c
@@ -0,0 +1,54 @@ 
+/* Disable the acc_on_device builtin; we want to test the libgomp library
+   function.  */
+/* TODO: Remove -DACC_DEVICE_TYPE_host once that is set by the test harness.  */
+/* { dg-additional-options "-fno-builtin-acc_on_device -DACC_DEVICE_TYPE_host" } */
+
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+  /* Host.  */
+
+  {
+    if (!acc_on_device (acc_device_none))
+      abort ();
+    if (!acc_on_device (acc_device_host))
+      abort ();
+    if (acc_on_device (acc_device_not_host))
+      abort ();
+  }
+
+
+  /* Host via offloading fallback mode.  */
+
+#pragma acc parallel if(0)
+  {
+    if (!acc_on_device (acc_device_none))
+      abort ();
+    if (!acc_on_device (acc_device_host))
+      abort ();
+    if (acc_on_device (acc_device_not_host))
+      abort ();
+  }
+
+
+#if !ACC_DEVICE_TYPE_host
+
+  /* Offloaded.  */
+
+#pragma acc parallel
+  {
+    if (acc_on_device (acc_device_none))
+      abort ();
+    if (acc_on_device (acc_device_host))
+      abort ();
+    if (!acc_on_device (acc_device_not_host))
+      abort ();
+  }
+
+#endif
+
+  return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
new file mode 100644
index 0000000..c4597a6
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
@@ -0,0 +1,39 @@ 
+! TODO: Remove -DACC_DEVICE_TYPE_host once that is set by the test harness.
+! { dg-additional-options "-cpp -DACC_DEVICE_TYPE_host" }
+! TODO: Have to disable the acc_on_device builtin for we want to test the
+! libgomp library function?  The command line option
+! '-fno-builtin-acc_on_device' is valid for C/C++/ObjC/ObjC++ but not for
+! Fortran.
+
+use openacc
+implicit none
+
+! Host.
+
+if (.not. acc_on_device (acc_device_none)) call abort
+if (.not. acc_on_device (acc_device_host)) call abort
+if (acc_on_device (acc_device_not_host)) call abort
+
+
+! Host via offloading fallback mode.
+
+!$acc parallel if(.false.)
+if (.not. acc_on_device (acc_device_none)) call abort
+if (.not. acc_on_device (acc_device_host)) call abort
+if (acc_on_device (acc_device_not_host)) call abort
+!$acc end parallel
+
+
+#if !ACC_DEVICE_TYPE_host
+
+! Offloaded.
+
+!$acc parallel
+if (acc_on_device (acc_device_none)) call abort
+if (acc_on_device (acc_device_host)) call abort
+if (.not. acc_on_device (acc_device_not_host)) call abort
+!$acc end parallel
+
+#endif
+
+end
diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
new file mode 100644
index 0000000..3787e1e
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
@@ -0,0 +1,39 @@ 
+! TODO: Remove -DACC_DEVICE_TYPE_host once that is set by the test harness.
+! { dg-additional-options "-cpp -DACC_DEVICE_TYPE_host" }
+! TODO: Have to disable the acc_on_device builtin for we want to test
+! the libgomp library function?  The command line option
+! '-fno-builtin-acc_on_device' is valid for C/C++/ObjC/ObjC++ but not
+! for Fortran.
+
+      USE OPENACC
+      IMPLICIT NONE
+
+!Host.
+
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
+      IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
+
+
+!Host via offloading fallback mode.
+
+!$ACC PARALLEL IF(.FALSE.)
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
+      IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
+!$ACC END PARALLEL
+
+
+#if !ACC_DEVICE_TYPE_host
+
+! Offloaded.
+
+!$ACC PARALLEL
+      IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
+      IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
+!$ACC END PARALLEL
+
+#endif
+
+      END
diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
new file mode 100644
index 0000000..1ee5926
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
@@ -0,0 +1,39 @@ 
+! TODO: Remove -DACC_DEVICE_TYPE_host once that is set by the test harness.
+! { dg-additional-options "-cpp -DACC_DEVICE_TYPE_host" }
+! TODO: Have to disable the acc_on_device builtin for we want to test
+! the libgomp library function?  The command line option
+! '-fno-builtin-acc_on_device' is valid for C/C++/ObjC/ObjC++ but not
+! for Fortran.
+
+      IMPLICIT NONE
+      INCLUDE "openacc_lib.h"
+
+!Host.
+
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
+      IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
+
+
+!Host via offloading fallback mode.
+
+!$ACC PARALLEL IF(.FALSE.)
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
+      IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
+!$ACC END PARALLEL
+
+
+#if !ACC_DEVICE_TYPE_host
+
+! Offloaded.
+
+!$ACC PARALLEL
+      IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT
+      IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT
+      IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT
+!$ACC END PARALLEL
+
+#endif
+
+      END