@@ -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.
@@ -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:
@@ -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. */
@@ -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
@@ -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
}
@@ -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)
@@ -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)
@@ -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.
new file mode 100644
@@ -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" } } */
new file mode 100644
@@ -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" } } */
new file mode 100644
@@ -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" } } */
new file mode 100644
@@ -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" } }
new file mode 100644
@@ -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" } }
new file mode 100644
@@ -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" } }
@@ -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>
@@ -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);
+}
@@ -234,6 +234,9 @@ GOMP_4.0.1 {
} GOMP_4.0;
OACC_2.0 {
+ global:
+ acc_on_device;
+ acc_on_device_;
};
GOACC_2.0 {
@@ -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)
@@ -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
@@ -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
@@ -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
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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