diff --git libgcc/ChangeLog.gomp libgcc/ChangeLog.gomp
index 085bfda..7de8361 100644
--- libgcc/ChangeLog.gomp
+++ libgcc/ChangeLog.gomp
@@ -1,3 +1,13 @@
+2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* config.host [nvptx-*] (extra_parts): Don't add "libgomp.a", and
+	"libgomp.spec".
+	* config/nvptx/gomp-acc_on_device.c: Remove file.
+	* config/nvptx/gomp-atomic.asm: Likewise.
+	* config/nvptx/t-nvptx (OBJS_libgomp): Don't set.
+	(gomp-acc_on_device.o, gomp-atomic.o, libgomp.a, libgomp.spec):
+	Remove targets.
+
 2015-08-03  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* config/nvptx/gomp-acc_on_device.c: Don't include
diff --git libgcc/config.host libgcc/config.host
index ee7ce03..3a2c75d 100644
--- libgcc/config.host
+++ libgcc/config.host
@@ -1304,11 +1304,7 @@ mep*-*-*)
 	;;
 nvptx-*)
 	tmake_file="$tmake_file nvptx/t-nvptx"
-	if test "x${enable_as_accelerator_for}" != x; then
-		extra_parts="crt0.o libgomp.a libgomp.spec"
-	else
-		extra_parts="crt0.o"
-	fi
+	extra_parts="crt0.o"
 	;;
 *)
 	echo "*** Configuration ${host} not supported" 1>&2
diff --git libgcc/config/nvptx/gomp-acc_on_device.c libgcc/config/nvptx/gomp-acc_on_device.c
deleted file mode 100644
index db94350..0000000
--- libgcc/config/nvptx/gomp-acc_on_device.c
+++ /dev/null
@@ -1,15 +0,0 @@
-/* The compiler always attempts to expand acc_on_device, but if the
-   user disables the builtin, or calls it via a pointer, we have this
-   version.  */
-
-int
-acc_on_device (int dev)
-{
-  /* Just rely on the compiler builtin.  */
-  return __builtin_acc_on_device (dev);
-}
-
-int acc_on_device_h_(int *d)
-{
-  return acc_on_device(*d);
-}
diff --git libgcc/config/nvptx/gomp-atomic.asm libgcc/config/nvptx/gomp-atomic.asm
deleted file mode 100644
index ae9d925..0000000
--- libgcc/config/nvptx/gomp-atomic.asm
+++ /dev/null
@@ -1,37 +0,0 @@
-
-// BEGIN PREAMBLE
-	.version	3.1
-	.target	sm_30
-	.address_size 64
-	.extern .shared .u8 sdata[];
-// END PREAMBLE
-
-// BEGIN VAR DEF: libgomp_ptx_lock
-.global .align 4 .u32 libgomp_ptx_lock;
-
-// BEGIN GLOBAL FUNCTION DECL: GOMP_atomic_start
-.visible .func GOMP_atomic_start;
-// BEGIN GLOBAL FUNCTION DEF: GOMP_atomic_start
-.visible .func GOMP_atomic_start
-{
-	.reg .pred 	%p<2>;
-	.reg .s32 	%r<2>;
-	.reg .s64 	%rd<2>;
-BB5_1:
-	mov.u64 	%rd1, libgomp_ptx_lock;
-	atom.global.cas.b32 	%r1, [%rd1], 0, 1;
-	setp.ne.s32	%p1, %r1, 0;
-	@%p1 bra 	BB5_1;
-	ret;
-	}
-// BEGIN GLOBAL FUNCTION DECL: GOMP_atomic_end
-.visible .func GOMP_atomic_end;
-// BEGIN GLOBAL FUNCTION DEF: GOMP_atomic_end
-.visible .func GOMP_atomic_end
-{
-	.reg .s32 	%r<2>;
-	.reg .s64 	%rd<2>;
-	mov.u64 	%rd1, libgomp_ptx_lock;
-	atom.global.exch.b32 	%r1, [%rd1], 0;
-	ret;
-	}
diff --git libgcc/config/nvptx/t-nvptx libgcc/config/nvptx/t-nvptx
index c8741c4..0c2cea0 100644
--- libgcc/config/nvptx/t-nvptx
+++ libgcc/config/nvptx/t-nvptx
@@ -13,14 +13,3 @@ crt0.o: $(srcdir)/config/nvptx/crt0.s
 # support it, and it may cause the build to fail, because of alloca usage, for
 # example.
 INHIBIT_LIBC_CFLAGS = -Dinhibit_libc
-
-gomp-acc_on_device.o: $(srcdir)/config/nvptx/gomp-acc_on_device.c
-	$(gcc_compile) -c -fno-builtin-acc_on_device $<
-gomp-atomic.o: $(srcdir)/config/nvptx/gomp-atomic.asm
-	cp $< $@
-
-OBJS_libgomp= gomp-acc_on_device.o gomp-atomic.o
-libgomp.a: $(OBJS_libgomp)
-	$(AR_CREATE_FOR_TARGET) $@ $(OBJS_libgomp)
-libgomp.spec:
-	echo "*link_gomp: -lgomp" >$@
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 3898930..fa7eb75 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,9 @@
 2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* config/nvptx/critical.c: New file, replacing empty file.
+	* config/nvptx/oacc-init.c: Likewise.
+	* config/nvptx/openacc.f90: New file.
+
 	Backport trunk r226760:
 
 	2015-08-10  Thomas Schwinge  <thomas@codesourcery.com>
diff --git libgomp/config/nvptx/critical.c libgomp/config/nvptx/critical.c
index e69de29..1f55aad 100644
--- libgomp/config/nvptx/critical.c
+++ libgomp/config/nvptx/critical.c
@@ -0,0 +1,57 @@
+/* GOMP atomic 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/>.  */
+
+__asm__ ("// BEGIN VAR DEF: libgomp_ptx_lock\n"
+	 ".global .align 4 .u32 libgomp_ptx_lock;\n"
+	 "\n"
+	 "// BEGIN GLOBAL FUNCTION DECL: GOMP_atomic_start\n"
+	 ".visible .func GOMP_atomic_start;\n"
+	 "// BEGIN GLOBAL FUNCTION DEF: GOMP_atomic_start\n"
+	 ".visible .func GOMP_atomic_start\n"
+	 "{\n"
+	 "	.reg .pred 	%p<2>;\n"
+	 "	.reg .s32 	%r<2>;\n"
+	 "	.reg .s64 	%rd<2>;\n"
+	 "BB5_1:\n"
+	 "	mov.u64 	%rd1, libgomp_ptx_lock;\n"
+	 "	atom.global.cas.b32 	%r1, [%rd1], 0, 1;\n"
+	 "	setp.ne.s32	%p1, %r1, 0;\n"
+	 "	@%p1 bra 	BB5_1;\n"
+	 "	ret;\n"
+	 "	}\n"
+	 "// BEGIN GLOBAL FUNCTION DECL: GOMP_atomic_end\n"
+	 ".visible .func GOMP_atomic_end;\n"
+	 "// BEGIN GLOBAL FUNCTION DEF: GOMP_atomic_end\n"
+	 ".visible .func GOMP_atomic_end\n"
+	 "{\n"
+	 "	.reg .s32 	%r<2>;\n"
+	 "	.reg .s64 	%rd<2>;\n"
+	 "	mov.u64 	%rd1, libgomp_ptx_lock;\n"
+	 "	atom.global.exch.b32 	%r1, [%rd1], 0;\n"
+	 "	ret;\n"
+	 "	}");
diff --git libgomp/config/nvptx/oacc-init.c libgomp/config/nvptx/oacc-init.c
index e69de29..5b93e74 100644
--- libgomp/config/nvptx/oacc-init.c
+++ libgomp/config/nvptx/oacc-init.c
@@ -0,0 +1,40 @@
+/* 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"
+
+/* The compiler always attempts to expand acc_on_device, but if the
+   user disables the builtin, or calls it via a pointer, we have this
+   version.  */
+
+int
+acc_on_device (int dev)
+{
+  /* Just rely on the compiler builtin.  */
+  return __builtin_acc_on_device (dev);
+}
diff --git libgomp/config/nvptx/openacc.f90 libgomp/config/nvptx/openacc.f90
new file mode 100644
index 0000000..5dfc6a1
--- /dev/null
+++ libgomp/config/nvptx/openacc.f90
@@ -0,0 +1,101 @@
+!  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_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
