crt0.s: Delete.

	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.
	add missing 2016-05-20  Nathan Sidwell  <nathan@acm.org> entry

From-SVN: r236702

gcc/testsuite/ChangeLog

+2016-05-25  Nathan Sidwell  <nathan@acm.org>
+
+	* gcc.c-torture/execute/921110-1.c: Fix abort decl.
+
 2016-05-25  Richard Biener  <rguenther@suse.de>
 
 	PR tree-optimization/71261
@@ -86,8 +90,8 @@
 
 2016-05-24  Ilya Verbin  <ilya.verbin@intel.com>
 
-	* gcc.target/i386/avx-ceil-sfix-2-vec.c: Define __NO_MATH_INLINES before
-	math.h is included.
+	* gcc.target/i386/avx-ceil-sfix-2-vec.c: Define __NO_MATH_INLINES
+	before math.h is included.
 	* gcc.target/i386/avx-floor-sfix-2-vec.c: Likewise.
 	* gcc.target/i386/avx-rint-sfix-2-vec.c: Likewise.
 	* gcc.target/i386/avx-round-sfix-2-vec.c: Likewise.
@@ -346,8 +350,14 @@
 
 	* gcc.target/powerpc/vec-adde.c: Change dejagnu options, fix code
 	formatting.
-	* gcc.target/powerpc/vec-adde-int128.c: Change dejagnu options, fix code
-	formatting.
+	* gcc.target/powerpc/vec-adde-int128.c: Change dejagnu options,
+	fix code formatting.
+
+2016-05-20  Nathan Sidwell  <nathan@acm.org>
+
+	* gcc.target/nvptx/uninit-decl.c: Force common storage,  add
+	non-common cases.
+	* gcc.dg/tree-ssa/ssa-store-ccp-2.c: Add -fcommon.
 
 2016-05-20  David Malcolm  <dmalcolm@redhat.com>
 

gcc/testsuite/gcc.c-torture/execute/921110-1.c

-extern int abort();
-typedef int (*frob)();
+extern void abort(void);
+typedef void (*frob)();
 frob f[] = {abort};
-main()
+
+int main(void)
 {
   exit(0);
 }

libgcc/ChangeLog

+2016-05-25  Nathan Sidwell  <nathan@acm.org>
+
+	* config/nvptx/crt0.s: Delete.
+	* config/nvptx/crt0.c: New.
+	* t-nvptx: Update.
+
 2016-05-19  Sandra Loosemore  <sandra@codesourcery.com>
 
 	* config.host [x86_64-*-cygwin*]: Handle tmake_eh_file for mixed 

libgcc/config/nvptx/crt0.c

+/* 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));
+}

libgcc/config/nvptx/crt0.s

-	.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;
-}

libgcc/config/nvptx/t-nvptx

@@ -6,8 +6,8 @@ LIB2ADD=$(srcdir)/config/nvptx/malloc.asm \
 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