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
