From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 17649 invoked by alias); 1 Dec 2015 15:28:57 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 17560 invoked by uid 89); 1 Dec 2015 15:28:56 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.1 required=5.0 tests=AWL,BAYES_50,KAM_LAZY_DOMAIN_SECURITY autolearn=no version=3.3.2 X-HELO: eggs.gnu.org Received: from eggs.gnu.org (HELO eggs.gnu.org) (208.118.235.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-SHA encrypted) ESMTPS; Tue, 01 Dec 2015 15:28:38 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1a3mr3-0006B4-P3 for gcc-patches@gcc.gnu.org; Tue, 01 Dec 2015 10:28:36 -0500 Received: from smtp.ispras.ru ([83.149.199.79]:49331) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1a3mr3-0006AU-9J for gcc-patches@gcc.gnu.org; Tue, 01 Dec 2015 10:28:33 -0500 Received: from condor.intra.ispras.ru (unknown [83.149.199.91]) by smtp.ispras.ru (Postfix) with ESMTP id BE49B20419; Tue, 1 Dec 2015 18:28:28 +0300 (MSK) Received: by condor.intra.ispras.ru (Postfix, from userid 23246) id 3ABE41225DEA; Tue, 1 Dec 2015 18:28:28 +0300 (MSK) From: Alexander Monakov To: gcc-patches@gcc.gnu.org Cc: Jakub Jelinek , Bernd Schmidt , Dmitry Melnik Subject: [gomp-nvptx 6/9] nvptx libgcc: rewrite in C Date: Tue, 01 Dec 2015 15:28:00 -0000 Message-Id: <1448983707-18854-7-git-send-email-amonakov@ispras.ru> In-Reply-To: <1448983707-18854-1-git-send-email-amonakov@ispras.ru> References: <1448983707-18854-1-git-send-email-amonakov@ispras.ru> X-detected-operating-system: by eggs.gnu.org: GNU/Linux 3.x X-Received-From: 83.149.199.79 X-IsSubscribed: yes X-SW-Source: 2015-12/txt/msg00120.txt.bz2 To easily build libgcc for -mgomp multilib, I've rewritten libgcc routines from asm to C. En passant, I've fixed a bug in malloc and realloc wrappers where they failed to handle out-of-memory conditions. I'm assuming it wasn't intentional. I also use a patch for Newlib that rewrites its nvptx-specific 'printf' implementation in C. * config/nvptx/crt0.c: New, rewritten in C from ... * config/nvptx/crt0.s: ...this. Delete. * config/nvptx/free.c: New, rewritten in C from ... * config/nvptx/free.asm: ...this. Delete. * config/nvptx/malloc.c: New, rewritten in C from ... * config/nvptx/malloc.asm: ...this. Delete. * config/nvptx/realloc.c: Handle out-of-memory condition. * config/nvptx/nvptx-malloc.h (__nvptx_real_free, __nvptx_real_malloc): Declare. * config/nvptx/stacks.c: New. * config/nvptx/t-nvptx: Adjust. --- libgcc/config/nvptx/crt0.c | 61 ++++++++++++++++++++++++++++++++++++++ libgcc/config/nvptx/crt0.s | 54 --------------------------------- libgcc/config/nvptx/free.asm | 50 ------------------------------- libgcc/config/nvptx/free.c | 34 +++++++++++++++++++++ libgcc/config/nvptx/malloc.asm | 55 ---------------------------------- libgcc/config/nvptx/malloc.c | 35 ++++++++++++++++++++++ libgcc/config/nvptx/nvptx-malloc.h | 5 ++++ libgcc/config/nvptx/realloc.c | 2 ++ libgcc/config/nvptx/stacks.c | 30 +++++++++++++++++++ libgcc/config/nvptx/t-nvptx | 11 +++---- 10 files changed, 173 insertions(+), 164 deletions(-) create mode 100644 libgcc/config/nvptx/crt0.c delete mode 100644 libgcc/config/nvptx/crt0.s delete mode 100644 libgcc/config/nvptx/free.asm create mode 100644 libgcc/config/nvptx/free.c delete mode 100644 libgcc/config/nvptx/malloc.asm create mode 100644 libgcc/config/nvptx/malloc.c create mode 100644 libgcc/config/nvptx/stacks.c diff --git a/libgcc/config/nvptx/crt0.c b/libgcc/config/nvptx/crt0.c new file mode 100644 index 0000000..74483c4 --- /dev/null +++ b/libgcc/config/nvptx/crt0.c @@ -0,0 +1,61 @@ +/* Startup routine for standalone execution. + + Copyright (C) 2015 Free Software Foundation, Inc. + + This file 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. + + This file 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 + . */ + +void exit (int); +void abort (void); +void __attribute__((kernel)) __main (int *, int, char *[]); + +static int *__exitval; + +void +exit (int arg) +{ + *__exitval = arg; + asm volatile ("exit;"); + __builtin_unreachable (); +} + +void +abort (void) +{ + exit (255); +} + +asm ("// BEGIN GLOBAL VAR DECL: __nvptx_stacks"); +asm (".extern .shared .u64 __nvptx_stacks[32];"); +asm ("// BEGIN GLOBAL VAR DECL: __nvptx_uni"); +asm (".extern .shared .u32 __nvptx_uni[32];"); + +extern int main (int argc, char *argv[]); + +void __attribute__((kernel)) +__main (int *__retval, int __argc, char *__argv[]) +{ + __exitval = __retval; + + static char gstack[131072] __attribute__((aligned(8))); + asm ("st.shared.u64 [__nvptx_stacks], %0;" : : "r" (gstack + sizeof gstack)); + asm ("st.shared.u32 [__nvptx_uni], %0;" : : "r" (0)); + + exit (main (__argc, __argv)); +} diff --git a/libgcc/config/nvptx/crt0.s b/libgcc/config/nvptx/crt0.s deleted file mode 100644 index 1ac69a5..0000000 --- a/libgcc/config/nvptx/crt0.s +++ /dev/null @@ -1,54 +0,0 @@ - .version 3.1 - .target sm_30 - .address_size 64 - -.global .u64 %__exitval; -// BEGIN GLOBAL FUNCTION DEF: abort -.visible .func abort -{ - .reg .u64 %rd1; - ld.global.u64 %rd1,[%__exitval]; - st.u32 [%rd1], 255; - exit; -} -// BEGIN GLOBAL FUNCTION DEF: exit -.visible .func exit (.param .u32 %arg) -{ - .reg .u64 %rd1; - .reg .u32 %val; - ld.param.u32 %val,[%arg]; - ld.global.u64 %rd1,[%__exitval]; - st.u32 [%rd1], %val; - exit; -} - -.visible .shared .u64 __nvptx_stacks[1]; -.global .align 8 .u8 %__softstack[131072]; - -.extern .func (.param.u32 retval) main (.param.u32 argc, .param.u64 argv); - -.visible .entry __main (.param .u64 __retval, .param.u32 __argc, .param.u64 __argv) -{ - .reg .u32 %r<3>; - .reg .u64 %rd<3>; - .param.u32 %argc; - .param.u64 %argp; - .param.u32 %mainret; - ld.param.u64 %rd0, [__retval]; - st.global.u64 [%__exitval], %rd0; - - .reg .u64 %stackptr; - mov.u64 %stackptr, %__softstack; - cvta.global.u64 %stackptr, %stackptr; - add.u64 %stackptr, %stackptr, 131072; - st.shared.u64 [__nvptx_stacks], %stackptr; - - ld.param.u32 %r1, [__argc]; - ld.param.u64 %rd1, [__argv]; - st.param.u32 [%argc], %r1; - st.param.u64 [%argp], %rd1; - call.uni (%mainret), main, (%argc, %argp); - ld.param.u32 %r1,[%mainret]; - st.s32 [%rd0], %r1; - exit; -} diff --git a/libgcc/config/nvptx/free.asm b/libgcc/config/nvptx/free.asm deleted file mode 100644 index 251d733..0000000 --- a/libgcc/config/nvptx/free.asm +++ /dev/null @@ -1,50 +0,0 @@ -// A wrapper around free to enable a realloc implementation. - -// Copyright (C) 2014-2015 Free Software Foundation, Inc. - -// This file 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. - -// This file 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 -// . - - .version 3.1 - .target sm_30 - .address_size 64 - -.extern .func free(.param.u64 %in_ar1); - -// BEGIN GLOBAL FUNCTION DEF: __nvptx_free -.visible .func __nvptx_free(.param.u64 %in_ar1) -{ - .reg.u64 %ar1; - .reg.u64 %hr10; - .reg.u64 %r23; - .reg.pred %r25; - .reg.u64 %r27; - ld.param.u64 %ar1, [%in_ar1]; - mov.u64 %r23, %ar1; - setp.eq.u64 %r25,%r23,0; - @%r25 bra $L1; - add.u64 %r27, %r23, -8; - { - .param.u64 %out_arg0; - st.param.u64 [%out_arg0], %r27; - call free, (%out_arg0); - } -$L1: - ret; - } diff --git a/libgcc/config/nvptx/free.c b/libgcc/config/nvptx/free.c new file mode 100644 index 0000000..90699c7 --- /dev/null +++ b/libgcc/config/nvptx/free.c @@ -0,0 +1,34 @@ +/* Implement free wrapper to help support realloc. + + Copyright (C) 2015 Free Software Foundation, Inc. + + This file 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. + + This file 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 +#include "nvptx-malloc.h" + +void +__nvptx_free (void *ptr) +{ + if (ptr == NULL) + return; + + __nvptx_real_free ((char *)ptr - 8); +} diff --git a/libgcc/config/nvptx/malloc.asm b/libgcc/config/nvptx/malloc.asm deleted file mode 100644 index 9f36715..0000000 --- a/libgcc/config/nvptx/malloc.asm +++ /dev/null @@ -1,55 +0,0 @@ -// A wrapper around malloc to enable a realloc implementation. - -// Copyright (C) 2014-2015 Free Software Foundation, Inc. - -// This file 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. - -// This file 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 -// . - - .version 3.1 - .target sm_30 - .address_size 64 - -.extern .func (.param.u64 %out_retval) malloc(.param.u64 %in_ar1); - -// BEGIN GLOBAL FUNCTION DEF: __nvptx_malloc -.visible .func (.param.u64 %out_retval) __nvptx_malloc(.param.u64 %in_ar1) -{ - .reg.u64 %ar1; -.reg.u64 %retval; - .reg.u64 %hr10; - .reg.u64 %r26; - .reg.u64 %r28; - .reg.u64 %r29; - .reg.u64 %r31; - ld.param.u64 %ar1, [%in_ar1]; - mov.u64 %r26, %ar1; - add.u64 %r28, %r26, 8; - { - .param.u64 %retval_in; - .param.u64 %out_arg0; - st.param.u64 [%out_arg0], %r28; - call (%retval_in), malloc, (%out_arg0); - ld.param.u64 %r29, [%retval_in]; - } - st.u64 [%r29], %r26; - add.u64 %r31, %r29, 8; - mov.u64 %retval, %r31; - st.param.u64 [%out_retval], %retval; - ret; -} diff --git a/libgcc/config/nvptx/malloc.c b/libgcc/config/nvptx/malloc.c new file mode 100644 index 0000000..2de995c --- /dev/null +++ b/libgcc/config/nvptx/malloc.c @@ -0,0 +1,35 @@ +/* Implement malloc wrapper to help support realloc. + + Copyright (C) 2015 Free Software Foundation, Inc. + + This file 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. + + This file 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 +#include "nvptx-malloc.h" + +void * +__nvptx_malloc (size_t sz) +{ + size_t *ptr = __nvptx_real_malloc (sz + 8); + if (!ptr) + return NULL; + *ptr = sz; + return ptr + 1; +} diff --git a/libgcc/config/nvptx/nvptx-malloc.h b/libgcc/config/nvptx/nvptx-malloc.h index d0ce65a..437f8b3 100644 --- a/libgcc/config/nvptx/nvptx-malloc.h +++ b/libgcc/config/nvptx/nvptx-malloc.h @@ -21,6 +21,11 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see . */ +/* malloc/realloc/free are remapped to these by the NVPTX backend. */ extern void __nvptx_free (void *); extern void *__nvptx_malloc (size_t); extern void *__nvptx_realloc (void *, size_t); + +/* And these are remapped back to "real" malloc/free. */ +extern void __nvptx_real_free (void *); +extern void *__nvptx_real_malloc (size_t); diff --git a/libgcc/config/nvptx/realloc.c b/libgcc/config/nvptx/realloc.c index 136f010..dba429e 100644 --- a/libgcc/config/nvptx/realloc.c +++ b/libgcc/config/nvptx/realloc.c @@ -33,6 +33,8 @@ __nvptx_realloc (void *ptr, size_t newsz) return NULL; } void *newptr = __nvptx_malloc (newsz); + if (!newptr) + return NULL; size_t oldsz; if (ptr == NULL) diff --git a/libgcc/config/nvptx/stacks.c b/libgcc/config/nvptx/stacks.c new file mode 100644 index 0000000..c597cd1 --- /dev/null +++ b/libgcc/config/nvptx/stacks.c @@ -0,0 +1,30 @@ +/* Define shared memory arrays for -msoft-stack and -munified-simt. + + Copyright (C) 2015 Free Software Foundation, Inc. + + This file 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. + + This file 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 + . */ + +/* __shared__ char *__nvptx_stacks[32]; */ +asm ("// BEGIN GLOBAL VAR DEF: __nvptx_stacks"); +asm (".visible .shared .u64 __nvptx_stacks[32];"); + +/* __shared__ unsigned __nvptx_uni[32]; */ +asm ("// BEGIN GLOBAL VAR DEF: __nvptx_uni"); +asm (".visible .shared .u32 __nvptx_uni[32];"); diff --git a/libgcc/config/nvptx/t-nvptx b/libgcc/config/nvptx/t-nvptx index 34d68cc..e302494 100644 --- a/libgcc/config/nvptx/t-nvptx +++ b/libgcc/config/nvptx/t-nvptx @@ -1,12 +1,13 @@ -LIB2ADD=$(srcdir)/config/nvptx/malloc.asm \ - $(srcdir)/config/nvptx/free.asm \ - $(srcdir)/config/nvptx/realloc.c +LIB2ADD=$(srcdir)/config/nvptx/malloc.c \ + $(srcdir)/config/nvptx/free.c \ + $(srcdir)/config/nvptx/realloc.c \ + $(srcdir)/config/nvptx/stacks.c LIB2ADDEH= LIB2FUNCS_EXCLUDE=__main -crt0.o: $(srcdir)/config/nvptx/crt0.s - cp $< $@ +crt0.o: $(srcdir)/config/nvptx/crt0.c + $(gcc_compile) -c $< # Prevent building "advanced" stuff (for example, gcov support). We don't # support it, and it may cause the build to fail, because of alloca usage, for