Rewrite libgcc/config/nvptx from PTX assembly to C, to allow building it for -mgomp multilib.
Previously posted here: [gomp-nvptx 6/9] nvptx libgcc: rewrite in C https://gcc.gnu.org/ml/gcc-patches/2015-12/msg00120.html [gomp-nvptx 2/7] nvptx libgcc: use attribute shared https://gcc.gnu.org/ml/gcc-patches/2016-03/msg01106.html 2016-03-10 Alexander Monakov <amona...@ispras.ru> * config/nvptx/crt0.c (__nvptx_stacks): Define in C. Use it... (__nvptx_uni): Ditto. (__main): ...here instead of inline asm. * config/nvptx/stacks.c (__nvptx_stacks): Define in C. (__nvptx_uni): Ditto. 2015-12-09 Alexander Monakov <amona...@ispras.ru> * 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. diff --git a/libgcc/config/nvptx/crt0.c b/libgcc/config/nvptx/crt0.c new file mode 100644 index 0000000..9e9a25e --- /dev/null +++ b/libgcc/config/nvptx/crt0.c @@ -0,0 +1,59 @@ +/* Startup routine for standalone execution. + + Copyright (C) 2015-2016 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 + <http://www.gnu.org/licenses/>. */ + +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); +} + +extern char *__nvptx_stacks[32] __attribute__((shared)); +extern unsigned __nvptx_uni[32] __attribute__((shared)); + +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))); + __nvptx_stacks[0] = gstack + sizeof gstack; + __nvptx_uni[0] = 0; + + exit (main (__argc, __argv)); +} diff --git a/libgcc/config/nvptx/crt0.s b/libgcc/config/nvptx/crt0.s deleted file mode 100644 index 38327ed..0000000 --- a/libgcc/config/nvptx/crt0.s +++ /dev/null @@ -1,45 +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; -} - -.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; - - 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 3b8e39e..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-2016 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 -// <http://www.gnu.org/licenses/>. - - .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..dde2e1f --- /dev/null +++ b/libgcc/config/nvptx/free.c @@ -0,0 +1,34 @@ +/* Implement free wrapper to help support realloc. + + Copyright (C) 2015-2016 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 + <http://www.gnu.org/licenses/>. */ + +#include <stddef.h> +#include "nvptx-malloc.h" + +void +__nvptx_free (void *ptr) +{ + if (ptr == NULL) + return; + + __nvptx_real_free ((size_t *)ptr - 1); +} diff --git a/libgcc/config/nvptx/malloc.asm b/libgcc/config/nvptx/malloc.asm deleted file mode 100644 index 3d60981..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-2016 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 -// <http://www.gnu.org/licenses/>. - - .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..f00eb48 --- /dev/null +++ b/libgcc/config/nvptx/malloc.c @@ -0,0 +1,35 @@ +/* Implement malloc wrapper to help support realloc. + + Copyright (C) 2015-2016 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 + <http://www.gnu.org/licenses/>. */ + +#include <stddef.h> +#include "nvptx-malloc.h" + +void * +__nvptx_malloc (size_t sz) +{ + size_t *ptr = __nvptx_real_malloc (sz + sizeof *ptr); + 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 9cbd846..1675f1c 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 <http://www.gnu.org/licenses/>. */ +/* 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 82d6a02..4f43f3c 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..4640fc9 --- /dev/null +++ b/libgcc/config/nvptx/stacks.c @@ -0,0 +1,25 @@ +/* Define shared memory arrays for -msoft-stack and -munified-simt. + + Copyright (C) 2015-2016 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 + <http://www.gnu.org/licenses/>. */ + +char *__nvptx_stacks[32] __attribute__((shared)) = { 0 }; +unsigned __nvptx_uni[32] __attribute__((shared)) = { 0 }; diff --git a/libgcc/config/nvptx/t-nvptx b/libgcc/config/nvptx/t-nvptx index e66188f..930bfd2 100644 --- a/libgcc/config/nvptx/t-nvptx +++ b/libgcc/config/nvptx/t-nvptx @@ -1,13 +1,14 @@ -LIB2ADD=$(srcdir)/config/nvptx/malloc.asm \ - $(srcdir)/config/nvptx/free.asm \ +LIB2ADD=$(srcdir)/config/nvptx/malloc.c \ + $(srcdir)/config/nvptx/free.c \ $(srcdir)/config/nvptx/realloc.c \ - $(srcdir)/config/nvptx/reduction.c + $(srcdir)/config/nvptx/reduction.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
