This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp-nvptx 6/9] nvptx libgcc: rewrite in C
- From: Alexander Monakov <amonakov at ispras dot ru>
- To: gcc-patches at gcc dot gnu dot org
- Cc: Jakub Jelinek <jakub at redhat dot com>, Bernd Schmidt <bschmidt at redhat dot com>, Dmitry Melnik <dm at ispras dot ru>
- Date: Tue, 1 Dec 2015 18:28:24 +0300
- Subject: [gomp-nvptx 6/9] nvptx libgcc: rewrite in C
- Authentication-results: sourceware.org; auth=none
- References: <1448983707-18854-1-git-send-email-amonakov at ispras dot ru>
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
+ <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);
+}
+
+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
-// <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..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
+ <http://www.gnu.org/licenses/>. */
+
+#include <stddef.h>
+#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
-// <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..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
+ <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 + 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
<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 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
+ <http://www.gnu.org/licenses/>. */
+
+/* __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