From patchwork Mon Aug 10 16:31:33 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 505726 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 5C51214018C for ; Tue, 11 Aug 2015 02:32:05 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=jPLmaCq2; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=PCCws5qDfrO8LnfM 3DWnZgATCjWpxM8Ur8UZoObJI0U6EI0B58l1ZtF3ip3scKs1k9H9UZAvLzxLYbPc YVSErQoWePNnTgpWUvDOs9f2sc6Ceh+r5WXGv1zzqPb53rWnhnKVUjqb3rtwfdj+ t5ZMotGXCV/2BUfILZLsf9wnb4I= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=NUi6574/XvNf+rKXzc+qOk EZNzg=; b=jPLmaCq2Mqs1Rc4THQlS144vctF9g4H2VG/B3trR45CJ6QkIsXCnsE R3cssoJ8wBjAA/+lPIThMmjNgJei/cHySvpd5ENCMzuDsGmOQyCXsoL8ebH8w1iT VOmxjYJS5UL1g12O5k8gpI+pBlNJA7k/Cb8kQLHl7hkjK09mm4WGI= Received: (qmail 10277 invoked by alias); 10 Aug 2015 16:31:58 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 10260 invoked by uid 89); 10 Aug 2015 16:31:58 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.3 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 10 Aug 2015 16:31:55 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1ZOpzL-0000o5-S7 from Thomas_Schwinge@mentor.com ; Mon, 10 Aug 2015 09:31:52 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Mon, 10 Aug 2015 17:31:50 +0100 From: Thomas Schwinge To: CC: Julian Brown , Tobias Burnus , Ilya Verbin , Jakub Jelinek Subject: [gomp4] [nvptx] Move GOMP stuff from libgcc to libgomp (was: [WIP] OpenMP 4 NVPTX support) In-Reply-To: <87twswyy9n.fsf@schwinge.name> References: <20150421155839.GZ1725@tucnak.redhat.com> <87twswyy9n.fsf@schwinge.name> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (x86_64-pc-linux-gnu) Date: Mon, 10 Aug 2015 18:31:33 +0200 Message-ID: <87h9o7umu2.fsf@schwinge.name> MIME-Version: 1.0 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(-) Grüße, Thomas 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