[PTX] crt0

[Nathan Sidwell <nathan at acm dot org>]

This  patch reimplements crt0  as C rather than assembly.  That means it;ll be 
good with 32 bit addresses.  I've also completely moved abort and exit into 
newlib (a fork on github), which had it's own copies that failed to set the exit 
code.  Things happened to work because they were never found, being hidden by 
the ones provided in crt0.s.  You'll need to update your newlib sources for exit 
codes to continue to work.

Alex, this should make it simple to add a  check in __main for a single thread 
soft stack allocation.  I imagine a check on a weakly declared __soft_stack 
symbol or something?  Or is that something nvptx-run should check for and 
initialize?

nathan


2016-05-25  Nathan Sidwell  <nathan@acm.org>

	libgcc/
	* config/nvptx/crt0.s: Delete.
	* config/nvptx/crt0.c: New.
	* t-nvptx: Update.

	gcc/testsuite/
	* gcc.c-torture/execute/921110-1.c: Fix abort decl.

Index: gcc/testsuite/gcc.c-torture/execute/921110-1.c
===================================================================
--- gcc/testsuite/gcc.c-torture/execute/921110-1.c	(revision 236531)
+++ gcc/testsuite/gcc.c-torture/execute/921110-1.c	(working copy)
@@ -1,7 +1,8 @@
-extern int abort();
-typedef int (*frob)();
+extern void abort(void);
+typedef void (*frob)();
 frob f[] = {abort};
-main()
+
+int main(void)
 {
   exit(0);
 }
Index: libgcc/config/nvptx/crt0.c
===================================================================
--- libgcc/config/nvptx/crt0.c	(nonexistent)
+++ libgcc/config/nvptx/crt0.c	(working copy)
@@ -0,0 +1,37 @@
+/* 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/>.  */
+
+int *__exitval_ptr;
+
+extern void __attribute__((noreturn)) exit (int status);
+extern int main (int, void **);
+
+void __attribute__((kernel))
+__main (int *rval_ptr, int argc, void **argv)
+{
+  __exitval_ptr = rval_ptr;
+  /* Store something non-zero, so the host knows something went wrong,
+     if we fail to reach exit properly.   */
+  if (rval_ptr)
+    *rval_ptr = 255;
+
+  exit (main (argc, argv));
+}
Index: libgcc/config/nvptx/crt0.s
===================================================================
--- libgcc/config/nvptx/crt0.s	(revision 236531)
+++ libgcc/config/nvptx/crt0.s	(nonexistent)
@@ -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;
-}
Index: libgcc/config/nvptx/t-nvptx
===================================================================
--- libgcc/config/nvptx/t-nvptx	(revision 236531)
+++ libgcc/config/nvptx/t-nvptx	(working copy)
@@ -6,8 +6,8 @@ LIB2ADD=$(srcdir)/config/nvptx/malloc.as
 LIB2ADDEH=
 LIB2FUNCS_EXCLUDE=__main
 
-crt0.o: $(srcdir)/config/nvptx/crt0.s
-	cp $< $@
+crt0.o: $(srcdir)/config/nvptx/crt0.c
+	$(crt_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