Hi! On Wed, 22 Jul 2015 18:04:20 +0200, I wrote: > On Tue, 21 Apr 2015 17:58:39 +0200, Jakub Jelinek 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 > 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 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 + + * 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 * 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 + * 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 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 + . */ + +__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 + . */ + +#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 +! 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 +! . + +! 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