crt0.s: Delete.
authorNathan Sidwell <nathan@gcc.gnu.org>
Wed, 25 May 2016 12:25:01 +0000 (12:25 +0000)
committerNathan Sidwell <nathan@gcc.gnu.org>
Wed, 25 May 2016 12:25:01 +0000 (12:25 +0000)
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
gcc/testsuite/gcc.c-torture/execute/921110-1.c
libgcc/ChangeLog
libgcc/config/nvptx/crt0.c [new file with mode: 0644]
libgcc/config/nvptx/crt0.s [deleted file]
libgcc/config/nvptx/t-nvptx

index abdfdd81baae46cd837671fd9b956dff2913fe7a..d41af4e00c48a3c39998ca400262bc7d9664b278 100644 (file)
@@ -1,3 +1,7 @@
+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.
 
        * 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>
 
index 27ff96f18fb8b76b87c6b69bf0f06abfa144d5f1..b04b84a65f9b92f8d00a0b72785392020aa97e25 100644 (file)
@@ -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 6b20f9bd2f9c78e347422dfd10f1d4a6d4e9358d..0b07f0854ee1ee41be871ee77e6b063f96806f3c 100644 (file)
@@ -1,3 +1,9 @@
+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 
diff --git a/libgcc/config/nvptx/crt0.c b/libgcc/config/nvptx/crt0.c
new file mode 100644 (file)
index 0000000..3b7382d
--- /dev/null
@@ -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));
+}
diff --git a/libgcc/config/nvptx/crt0.s b/libgcc/config/nvptx/crt0.s
deleted file mode 100644 (file)
index 38327ed..0000000
+++ /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;
-}
index e66188f87225e7c2d095c38970436b344d118cbc..6001021ff2155ce10e2a3922774bbc9cd054f448 100644 (file)
@@ -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