@@ -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
@@ -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
deleted file mode 100644
@@ -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);
-}
deleted file mode 100644
@@ -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;
- }
@@ -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" >$@
@@ -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>
@@ -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"
+ " }");
@@ -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);
+}
new file mode 100644
@@ -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