diff mbox

nvptx libgcc changes

Message ID alpine.LNX.2.20.1604202005590.14803@monopod.intra.ispras.ru
State New
Headers show

Commit Message

Alexander Monakov April 20, 2016, 5:06 p.m. UTC 
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  <amonakov@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  <amonakov@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 mbox

Patch

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