Hi! On Wed, 22 Jul 2015 18:04:20 +0200, I wrote: > On Tue, 21 Apr 2015 17:58:39 +0200, Jakub Jelinek <ja...@redhat.com> wrote: > > Attached is a minimal patch to get at least a trivial OpenMP 4.0 testcase > > offloading to NVPTX (the first patch). The second patch is WIP, just first > > few needed changes to make libgomp to build for NVPTX (several weeks of work > > at least). > > We're not in particular working on making nvptx offloading work for > OpenMP, but also for OpenACC offloading a tiny bit of code is required to > be shipped in an offloading device's runtime library -- code that > conceptually belongs into libgomp. (On gomp-4_0-branch, it currently > lives in libgcc because that was easier to do.) [...]
> Next, we can then (on gomp-4_0-branch) move the libgcc code into libgomp:
>
> commit d8d75d17630d7633be4f1733fd195a104cb2ccc4
> Author: Thomas Schwinge <tho...@codesourcery.com>
> Date: Wed Jul 22 13:05:16 2015 +0200
>
> [nvptx] Move GOMP stuff from libgcc to libgomp
Committed to gomp-4_0-branch in r226762:
commit c49a2b23a76591f26b4076401647011442df92df
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon Aug 10 16:26:46 2015 +0000
[nvptx] Move GOMP stuff from libgcc to libgomp
libgcc/
* 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.
libgomp/
* config/nvptx/critical.c: New file, replacing empty file.
* config/nvptx/oacc-init.c: Likewise.
* config/nvptx/openacc.f90: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@226762
138bc75d-0d04-0410-961f-82ee72b054a4
---
libgcc/ChangeLog.gomp | 10 +++
libgcc/config.host | 6 +-
libgcc/config/nvptx/gomp-acc_on_device.c | 15 -----
libgcc/config/nvptx/gomp-atomic.asm | 37 -----------
libgcc/config/nvptx/t-nvptx | 11 ----
libgomp/ChangeLog.gomp | 4 ++
libgomp/config/nvptx/critical.c | 57 +++++++++++++++++
libgomp/config/nvptx/oacc-init.c | 40 ++++++++++++
libgomp/config/nvptx/openacc.f90 | 101 ++++++++++++++++++++++++++++++
9 files changed, 213 insertions(+), 68 deletions(-)
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 <tho...@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 <tho...@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 <tho...@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 <tho...@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 <bur...@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
Grüße,
Thomas
pgpDzzXjbP5VB.pgp
Description: PGP signature
