new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
@@ -0,0 +1,12 @@
+/* This header is used during the build process to find the size and
+ alignment of the public OpenMP locks, so that we can export data
+ structures without polluting the namespace.
+
+ When using the Linux futex primitive, non-recursive locks require
+ one int. Recursive locks require we identify the owning task
+ and so require in addition one int and a pointer. */
+
+typedef int omp_lock_t;
+typedef struct { int lock, count; void *owner; } omp_nest_lock_t;
+typedef int omp_lock_25_t;
+typedef struct { int owner, count; } omp_nest_lock_25_t;
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
new file mode 100644
Next, we can then (on gomp-4_0-branch) move the libgcc code into libgomp:
commit d8d75d17630d7633be4f1733fd195a104cb2ccc4
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed Jul 22 13:05:16 2015 +0200
[nvptx] Move GOMP stuff from libgcc to libgomp
---
libgcc/config.host | 6 +---
libgcc/config/nvptx/gomp-acc_on_device.c | 9 -----
libgcc/config/nvptx/gomp-atomic.asm | 37 ---------------------
libgcc/config/nvptx/t-nvptx | 11 ------
libgomp/config/nvptx/critical.c | 57 ++++++++++++++++++++++++++++++++
libgomp/config/nvptx/oacc-init.c | 35 ++++++++++++++++++++
6 files changed, 93 insertions(+), 62 deletions(-)
@@ -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,9 +0,0 @@
-int acc_on_device(int d)
-{
- return __builtin_acc_on_device(d);
-}
-
-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" >$@
@@ -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,35 @@
+/* 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"
+
+int
+acc_on_device (acc_device_t d)
+{
+ return __builtin_acc_on_device (d);
+}