nvptx offloading linking

	gcc/
	* config/nvptx/mkoffload.c (Kind, Vis): Remove enums.
	(Token, Stmt): Remove structs.
	(decls, vars, fns): Remove variables.
	(alloc_comment, append_stmt, is_keyword): Remove macros.
	(tokenize, write_token, write_tokens, alloc_stmt, rev_stmts)
	(write_stmt, write_stmts, parse_insn, parse_list_nosemi)
	(parse_init, parse_file): Remove functions.
	(read_file): Accept a pointer to a length and store into it.
	(process): Don't try to parse the input file, just write it out as
	a string, but looking for maps.  Also write out the length.
	(main): Don't use "-S" to compile PTX code.

	libgomp/
	* oacc-ptx.h: Remove file, moving its content into...
	* config/nvptx/fortran.c: ... here...
	* config/nvptx/oacc-init.c: ..., here...
	* config/nvptx/oacc-parallel.c: ..., and here.
	* config/nvptx/openacc.f90: New file.
	* plugin/plugin-nvptx.c: Don't include "oacc-ptx.h".
	(link_ptx): Don't link in predefined bits of PTX code.

Co-Authored-By: Bernd Schmidt <bernds@codesourcery.com>

From-SVN: r228418

4 files changed, 542 insertions, 0 deletions

diff --git a/libgomp/config/nvptx/fortran.c b/libgomp/config/nvptx/fortran.c
index e69de29..58ca790 100644
--- a/libgomp/config/nvptx/fortran.c
+++ b/libgomp/config/nvptx/fortran.c
@@ -0,0 +1,40 @@
+/* OpenACC Runtime Fortran wrapper routines
+
+   Copyright (C) 2014-2015 Free Software Foundation, Inc.
+
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+/* Temporary hack; this will be provided by libgfortran.  */
+
+extern void _gfortran_abort (void);
+
+__asm__ ("// BEGIN GLOBAL FUNCTION DECL: _gfortran_abort\n"
+	 ".visible .func _gfortran_abort;\n"
+	 "// BEGIN GLOBAL FUNCTION DEF: _gfortran_abort\n"
+	 ".visible .func _gfortran_abort\n"
+	 "{\n"
+	 "trap;\n"
+	 "ret;\n"
+	 "}\n");
diff --git a/libgomp/config/nvptx/oacc-init.c b/libgomp/config/nvptx/oacc-init.c
index e69de29..c57a3f3 100644
--- a/libgomp/config/nvptx/oacc-init.c
+++ b/libgomp/config/nvptx/oacc-init.c
@@ -0,0 +1,42 @@
+/* OpenACC Runtime initialization routines
+
+   Copyright (C) 2014-2015 Free Software Foundation, Inc.
+
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "openacc.h"
+
+/* For -O and higher, the compiler always attempts to expand acc_on_device, but
+   if the user disables the builtin, or calls it via a pointer, we'll need this
+   version.
+
+   Compile this with optimization, so that the compiler expands
+   this, rather than generating infinitely recursive code.  */
+
+int __attribute__ ((__optimize__ ("O2")))
+acc_on_device (acc_device_t dev)
+{
+  return __builtin_acc_on_device (dev);
+}
diff --git a/libgomp/config/nvptx/oacc-parallel.c b/libgomp/config/nvptx/oacc-parallel.c
index e69de29..b971256 100644
--- a/libgomp/config/nvptx/oacc-parallel.c
+++ b/libgomp/config/nvptx/oacc-parallel.c
@@ -0,0 +1,358 @@
+/* OpenACC constructs
+
+   Copyright (C) 2014-2015 Free Software Foundation, Inc.
+
+   Contributed by Mentor Embedded.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+#include "libgomp_g.h"
+
+__asm__ (".visible .func (.param .u32 %out_retval) GOACC_tid (.param .u32 %in_ar1);\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_ntid (.param .u32 %in_ar1);\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_ctaid (.param .u32 %in_ar1);\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_nctaid (.param .u32 %in_ar1);\n"
+	 "// BEGIN GLOBAL FUNCTION DECL: GOACC_get_num_threads\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads;\n"
+	 "// BEGIN GLOBAL FUNCTION DECL: GOACC_get_thread_num\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num;\n"
+	 "// BEGIN GLOBAL FUNCTION DECL: abort\n"
+	 ".extern .func abort;\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_tid (.param .u32 %in_ar1)\n"
+	 "{\n"
+	 ".reg .u32 %ar1;\n"
+	 ".reg .u32 %retval;\n"
+	 ".reg .u64 %hr10;\n"
+	 ".reg .u32 %r22;\n"
+	 ".reg .u32 %r23;\n"
+	 ".reg .u32 %r24;\n"
+	 ".reg .u32 %r25;\n"
+	 ".reg .u32 %r26;\n"
+	 ".reg .u32 %r27;\n"
+	 ".reg .u32 %r28;\n"
+	 ".reg .u32 %r29;\n"
+	 ".reg .pred %r30;\n"
+	 ".reg .u32 %r31;\n"
+	 ".reg .pred %r32;\n"
+	 ".reg .u32 %r33;\n"
+	 ".reg .pred %r34;\n"
+	 ".local .align 8 .b8 %frame[4];\n"
+	 "ld.param.u32 %ar1,[%in_ar1];\n"
+	 "mov.u32 %r27,%ar1;\n"
+	 "st.local.u32 [%frame],%r27;\n"
+	 "ld.local.u32 %r28,[%frame];\n"
+	 "mov.u32 %r29,1;\n"
+	 "setp.eq.u32 %r30,%r28,%r29;\n"
+	 "@%r30 bra $L4;\n"
+	 "mov.u32 %r31,2;\n"
+	 "setp.eq.u32 %r32,%r28,%r31;\n"
+	 "@%r32 bra $L5;\n"
+	 "mov.u32 %r33,0;\n"
+	 "setp.eq.u32 %r34,%r28,%r33;\n"
+	 "@!%r34 bra $L8;\n"
+	 "mov.u32 %r23,%tid.x;\n"
+	 "mov.u32 %r22,%r23;\n"
+	 "bra $L7;\n"
+	 "$L4:\n"
+	 "mov.u32 %r24,%tid.y;\n"
+	 "mov.u32 %r22,%r24;\n"
+	 "bra $L7;\n"
+	 "$L5:\n"
+	 "mov.u32 %r25,%tid.z;\n"
+	 "mov.u32 %r22,%r25;\n"
+	 "bra $L7;\n"
+	 "$L8:\n"
+	 "{\n"
+	 "{\n"
+	 "call abort;\n"
+	 "}\n"
+	 "}\n"
+	 "$L7:\n"
+	 "mov.u32 %r26,%r22;\n"
+	 "mov.u32 %retval,%r26;\n"
+	 "st.param.u32 [%out_retval],%retval;\n"
+	 "ret;\n"
+	 "}\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_ntid (.param .u32 %in_ar1)\n"
+	 "{\n"
+	 ".reg .u32 %ar1;\n"
+	 ".reg .u32 %retval;\n"
+	 ".reg .u64 %hr10;\n"
+	 ".reg .u32 %r22;\n"
+	 ".reg .u32 %r23;\n"
+	 ".reg .u32 %r24;\n"
+	 ".reg .u32 %r25;\n"
+	 ".reg .u32 %r26;\n"
+	 ".reg .u32 %r27;\n"
+	 ".reg .u32 %r28;\n"
+	 ".reg .u32 %r29;\n"
+	 ".reg .pred %r30;\n"
+	 ".reg .u32 %r31;\n"
+	 ".reg .pred %r32;\n"
+	 ".reg .u32 %r33;\n"
+	 ".reg .pred %r34;\n"
+	 ".local .align 8 .b8 %frame[4];\n"
+	 "ld.param.u32 %ar1,[%in_ar1];\n"
+	 "mov.u32 %r27,%ar1;\n"
+	 "st.local.u32 [%frame],%r27;\n"
+	 "ld.local.u32 %r28,[%frame];\n"
+	 "mov.u32 %r29,1;\n"
+	 "setp.eq.u32 %r30,%r28,%r29;\n"
+	 "@%r30 bra $L11;\n"
+	 "mov.u32 %r31,2;\n"
+	 "setp.eq.u32 %r32,%r28,%r31;\n"
+	 "@%r32 bra $L12;\n"
+	 "mov.u32 %r33,0;\n"
+	 "setp.eq.u32 %r34,%r28,%r33;\n"
+	 "@!%r34 bra $L15;\n"
+	 "mov.u32 %r23,%ntid.x;\n"
+	 "mov.u32 %r22,%r23;\n"
+	 "bra $L14;\n"
+	 "$L11:\n"
+	 "mov.u32 %r24,%ntid.y;\n"
+	 "mov.u32 %r22,%r24;\n"
+	 "bra $L14;\n"
+	 "$L12:\n"
+	 "mov.u32 %r25,%ntid.z;\n"
+	 "mov.u32 %r22,%r25;\n"
+	 "bra $L14;\n"
+	 "$L15:\n"
+	 "{\n"
+	 "{\n"
+	 "call abort;\n"
+	 "}\n"
+	 "}\n"
+	 "$L14:\n"
+	 "mov.u32 %r26,%r22;\n"
+	 "mov.u32 %retval,%r26;\n"
+	 "st.param.u32 [%out_retval],%retval;\n"
+	 "ret;\n"
+	 "}\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_ctaid (.param .u32 %in_ar1)\n"
+	 "{\n"
+	 ".reg .u32 %ar1;\n"
+	 ".reg .u32 %retval;\n"
+	 ".reg .u64 %hr10;\n"
+	 ".reg .u32 %r22;\n"
+	 ".reg .u32 %r23;\n"
+	 ".reg .u32 %r24;\n"
+	 ".reg .u32 %r25;\n"
+	 ".reg .u32 %r26;\n"
+	 ".reg .u32 %r27;\n"
+	 ".reg .u32 %r28;\n"
+	 ".reg .u32 %r29;\n"
+	 ".reg .pred %r30;\n"
+	 ".reg .u32 %r31;\n"
+	 ".reg .pred %r32;\n"
+	 ".reg .u32 %r33;\n"
+	 ".reg .pred %r34;\n"
+	 ".local .align 8 .b8 %frame[4];\n"
+	 "ld.param.u32 %ar1,[%in_ar1];\n"
+	 "mov.u32 %r27,%ar1;\n"
+	 "st.local.u32 [%frame],%r27;\n"
+	 "ld.local.u32 %r28,[%frame];\n"
+	 "mov.u32 %r29,1;\n"
+	 "setp.eq.u32 %r30,%r28,%r29;\n"
+	 "@%r30 bra $L18;\n"
+	 "mov.u32 %r31,2;\n"
+	 "setp.eq.u32 %r32,%r28,%r31;\n"
+	 "@%r32 bra $L19;\n"
+	 "mov.u32 %r33,0;\n"
+	 "setp.eq.u32 %r34,%r28,%r33;\n"
+	 "@!%r34 bra $L22;\n"
+	 "mov.u32 %r23,%ctaid.x;\n"
+	 "mov.u32 %r22,%r23;\n"
+	 "bra $L21;\n"
+	 "$L18:\n"
+	 "mov.u32 %r24,%ctaid.y;\n"
+	 "mov.u32 %r22,%r24;\n"
+	 "bra $L21;\n"
+	 "$L19:\n"
+	 "mov.u32 %r25,%ctaid.z;\n"
+	 "mov.u32 %r22,%r25;\n"
+	 "bra $L21;\n"
+	 "$L22:\n"
+	 "{\n"
+	 "{\n"
+	 "call abort;\n"
+	 "}\n"
+	 "}\n"
+	 "$L21:\n"
+	 "mov.u32 %r26,%r22;\n"
+	 "mov.u32 %retval,%r26;\n"
+	 "st.param.u32 [%out_retval],%retval;\n"
+	 "ret;\n"
+	 "}\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_nctaid (.param .u32 %in_ar1)\n"
+	 "{\n"
+	 ".reg .u32 %ar1;\n"
+	 ".reg .u32 %retval;\n"
+	 ".reg .u64 %hr10;\n"
+	 ".reg .u32 %r22;\n"
+	 ".reg .u32 %r23;\n"
+	 ".reg .u32 %r24;\n"
+	 ".reg .u32 %r25;\n"
+	 ".reg .u32 %r26;\n"
+	 ".reg .u32 %r27;\n"
+	 ".reg .u32 %r28;\n"
+	 ".reg .u32 %r29;\n"
+	 ".reg .pred %r30;\n"
+	 ".reg .u32 %r31;\n"
+	 ".reg .pred %r32;\n"
+	 ".reg .u32 %r33;\n"
+	 ".reg .pred %r34;\n"
+	 ".local .align 8 .b8 %frame[4];\n"
+	 "ld.param.u32 %ar1,[%in_ar1];\n"
+	 "mov.u32 %r27,%ar1;\n"
+	 "st.local.u32 [%frame],%r27;\n"
+	 "ld.local.u32 %r28,[%frame];\n"
+	 "mov.u32 %r29,1;\n"
+	 "setp.eq.u32 %r30,%r28,%r29;\n"
+	 "@%r30 bra $L25;\n"
+	 "mov.u32 %r31,2;\n"
+	 "setp.eq.u32 %r32,%r28,%r31;\n"
+	 "@%r32 bra $L26;\n"
+	 "mov.u32 %r33,0;\n"
+	 "setp.eq.u32 %r34,%r28,%r33;\n"
+	 "@!%r34 bra $L29;\n"
+	 "mov.u32 %r23,%nctaid.x;\n"
+	 "mov.u32 %r22,%r23;\n"
+	 "bra $L28;\n"
+	 "$L25:\n"
+	 "mov.u32 %r24,%nctaid.y;\n"
+	 "mov.u32 %r22,%r24;\n"
+	 "bra $L28;\n"
+	 "$L26:\n"
+	 "mov.u32 %r25,%nctaid.z;\n"
+	 "mov.u32 %r22,%r25;\n"
+	 "bra $L28;\n"
+	 "$L29:\n"
+	 "{\n"
+	 "{\n"
+	 "call abort;\n"
+	 "}\n"
+	 "}\n"
+	 "$L28:\n"
+	 "mov.u32 %r26,%r22;\n"
+	 "mov.u32 %retval,%r26;\n"
+	 "st.param.u32 [%out_retval],%retval;\n"
+	 "ret;\n"
+	 "}\n"
+	 "// BEGIN GLOBAL FUNCTION DEF: GOACC_get_num_threads\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads\n"
+	 "{\n"
+	 ".reg .u32 %retval;\n"
+	 ".reg .u64 %hr10;\n"
+	 ".reg .u32 %r22;\n"
+	 ".reg .u32 %r23;\n"
+	 ".reg .u32 %r24;\n"
+	 ".reg .u32 %r25;\n"
+	 ".reg .u32 %r26;\n"
+	 ".reg .u32 %r27;\n"
+	 ".reg .u32 %r28;\n"
+	 ".reg .u32 %r29;\n"
+	 "mov.u32 %r26,0;\n"
+	 "{\n"
+	 ".param .u32 %retval_in;\n"
+	 "{\n"
+	 ".param .u32 %out_arg0;\n"
+	 "st.param.u32 [%out_arg0],%r26;\n"
+	 "call (%retval_in),GOACC_ntid,(%out_arg0);\n"
+	 "}\n"
+	 "ld.param.u32 %r27,[%retval_in];\n"
+	 "}\n"
+	 "mov.u32 %r22,%r27;\n"
+	 "mov.u32 %r28,0;\n"
+	 "{\n"
+	 ".param .u32 %retval_in;\n"
+	 "{\n"
+	 ".param .u32 %out_arg0;\n"
+	 "st.param.u32 [%out_arg0],%r28;\n"
+	 "call (%retval_in),GOACC_nctaid,(%out_arg0);\n"
+	 "}\n"
+	 "ld.param.u32 %r29,[%retval_in];\n"
+	 "}\n"
+	 "mov.u32 %r23,%r29;\n"
+	 "mul.lo.u32 %r24,%r22,%r23;\n"
+	 "mov.u32 %r25,%r24;\n"
+	 "mov.u32 %retval,%r25;\n"
+	 "st.param.u32 [%out_retval],%retval;\n"
+	 "ret;\n"
+	 "}\n"
+	 "// BEGIN GLOBAL FUNCTION DEF: GOACC_get_thread_num\n"
+	 ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num\n"
+	 "{\n"
+	 ".reg .u32 %retval;\n"
+	 ".reg .u64 %hr10;\n"
+	 ".reg .u32 %r22;\n"
+	 ".reg .u32 %r23;\n"
+	 ".reg .u32 %r24;\n"
+	 ".reg .u32 %r25;\n"
+	 ".reg .u32 %r26;\n"
+	 ".reg .u32 %r27;\n"
+	 ".reg .u32 %r28;\n"
+	 ".reg .u32 %r29;\n"
+	 ".reg .u32 %r30;\n"
+	 ".reg .u32 %r31;\n"
+	 ".reg .u32 %r32;\n"
+	 ".reg .u32 %r33;\n"
+	 "mov.u32 %r28,0;\n"
+	 "{\n"
+	 ".param .u32 %retval_in;\n"
+	 "{\n"
+	 ".param .u32 %out_arg0;\n"
+	 "st.param.u32 [%out_arg0],%r28;\n"
+	 "call (%retval_in),GOACC_ntid,(%out_arg0);\n"
+	 "}\n"
+	 "ld.param.u32 %r29,[%retval_in];\n"
+	 "}\n"
+	 "mov.u32 %r22,%r29;\n"
+	 "mov.u32 %r30,0;\n"
+	 "{\n"
+	 ".param .u32 %retval_in;\n"
+	 "{\n"
+	 ".param .u32 %out_arg0;\n"
+	 "st.param.u32 [%out_arg0],%r30;\n"
+	 "call (%retval_in),GOACC_ctaid,(%out_arg0);\n"
+	 "}\n"
+	 "ld.param.u32 %r31,[%retval_in];\n"
+	 "}\n"
+	 "mov.u32 %r23,%r31;\n"
+	 "mul.lo.u32 %r24,%r22,%r23;\n"
+	 "mov.u32 %r32,0;\n"
+	 "{\n"
+	 ".param .u32 %retval_in;\n"
+	 "{\n"
+	 ".param .u32 %out_arg0;\n"
+	 "st.param.u32 [%out_arg0],%r32;\n"
+	 "call (%retval_in),GOACC_tid,(%out_arg0);\n"
+	 "}\n"
+	 "ld.param.u32 %r33,[%retval_in];\n"
+	 "}\n"
+	 "mov.u32 %r25,%r33;\n"
+	 "add.u32 %r26,%r24,%r25;\n"
+	 "mov.u32 %r27,%r26;\n"
+	 "mov.u32 %retval,%r27;\n"
+	 "st.param.u32 [%out_retval],%retval;\n"
+	 "ret;\n"
+	 "}\n");
diff --git a/libgomp/config/nvptx/openacc.f90 b/libgomp/config/nvptx/openacc.f90
new file mode 100644
index 0000000..d8b5c06
--- /dev/null
+++ b/libgomp/config/nvptx/openacc.f90
@@ -0,0 +1,102 @@
+!  OpenACC Runtime Library Definitions.
+
+!  Copyright (C) 2014-2015 Free Software Foundation, Inc.
+
+!  Contributed by Tobias Burnus <burnus@net-b.de>
+!              and Mentor Embedded.
+
+!  This file is part of the GNU Offloading and Multi Processing Library
+!  (libgomp).
+
+!  Libgomp is free software; you can redistribute it and/or modify it
+!  under the terms of the GNU General Public License as published by
+!  the Free Software Foundation; either version 3, or (at your option)
+!  any later version.
+
+!  Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+!  WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+!  FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+!  more details.
+
+!  Under Section 7 of GPL version 3, you are granted additional
+!  permissions described in the GCC Runtime Library Exception, version
+!  3.1, as published by the Free Software Foundation.
+
+!  You should have received a copy of the GNU General Public License and
+!  a copy of the GCC Runtime Library Exception along with this program;
+!  see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+!  <http://www.gnu.org/licenses/>.
+
+! Wrapper functions will be built from openacc.f90.  We use a separate file
+! here, because for using ../../openacc.f90, implementations are required for
+! all the functions that it wraps, which we currently don't provide, so linking
+! would fail.
+
+module openacc_kinds
+  use iso_fortran_env, only: int32
+  implicit none
+
+  private :: int32
+  public :: acc_device_kind
+
+  integer, parameter :: acc_device_kind = int32
+
+  public :: acc_device_none, acc_device_default, acc_device_host
+  public :: acc_device_not_host, acc_device_nvidia
+
+  ! Keep in sync with include/gomp-constants.h.
+  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_host_nonshm = 3 removed.
+  integer (acc_device_kind), parameter :: acc_device_not_host = 4
+  integer (acc_device_kind), parameter :: acc_device_nvidia = 5
+
+end module
+
+module openacc_internal
+  use openacc_kinds
+  implicit none
+
+  interface
+    function acc_on_device_h (d)
+      import
+      integer (acc_device_kind) d
+      logical acc_on_device_h
+    end function
+  end interface
+
+  interface
+    function acc_on_device_l (d) &
+        bind (C, name = "acc_on_device")
+      use iso_c_binding, only: c_int
+      integer (c_int) :: acc_on_device_l
+      integer (c_int), value :: d
+    end function
+  end interface
+end module
+
+module openacc
+  use openacc_kinds
+  use openacc_internal
+  implicit none
+
+  public :: acc_on_device
+
+  interface acc_on_device
+    procedure :: acc_on_device_h
+  end interface
+
+end module openacc
+
+function acc_on_device_h (d)
+  use openacc_internal, only: acc_on_device_l
+  use openacc_kinds
+  integer (acc_device_kind) d
+  logical acc_on_device_h
+  if (acc_on_device_l (d) .eq. 1) then
+    acc_on_device_h = .TRUE.
+  else
+    acc_on_device_h = .FALSE.
+  end if
+end function