Add the nvptx port.

	* configure.ac: Handle nvptx-*-*.
	* configure: Regenerate.

	gcc/
	* config/nvptx/nvptx.c: New file.
	* config/nvptx/nvptx.h: New file.
	* config/nvptx/nvptx-protos.h: New file.
	* config/nvptx/nvptx.md: New file.
	* config/nvptx/t-nvptx: New file.
	* config/nvptx/nvptx.opt: New file.
	* common/config/nvptx/nvptx-common.c: New file.
	* config.gcc: Handle nvptx-*-*.

	libgcc/
	* config.host: Handle nvptx-*-*.
	* shared-object.mk (as-flags-$o): Define.
	($(base)$(objext), $(base)_s$(objext)): Use it instead of
	-xassembler-with-cpp.
	* static-object.mk: Identical changes.
	* config/nvptx/t-nvptx: New file.
	* config/nvptx/crt0.s: New file.
	* config/nvptx/free.asm: New file.
	* config/nvptx/malloc.asm: New file.
	* config/nvptx/realloc.c: New file.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@217295 138bc75d-0d04-0410-961f-82ee72b054a4

10 files changed, 261 insertions, 4 deletions

diff --git a/libgcc/ChangeLog b/libgcc/ChangeLog
index 073ed11d321..17c8bb17678 100644
--- a/libgcc/ChangeLog
+++ b/libgcc/ChangeLog
@@ -1,3 +1,16 @@
+2014-11-06  Bernd Schmidt  <bernds@codesourcery.com>
+
+	* config.host: Handle nvptx-*-*.
+	* shared-object.mk (as-flags-$o): Define.
+	($(base)$(objext), $(base)_s$(objext)): Use it instead of
+	-xassembler-with-cpp.
+	* static-object.mk: Identical changes.
+	* config/nvptx/t-nvptx: New file.
+	* config/nvptx/crt0.s: New file.
+	* config/nvptx/free.asm: New file.
+	* config/nvptx/malloc.asm: New file.
+	* config/nvptx/realloc.c: New file.
+
 2014-10-30  Joseph Myers  <joseph@codesourcery.com>
 
 	* Makefile.in (libgcc.map.in): New target.
diff --git a/libgcc/config.host b/libgcc/config.host
index f3cc2765583..9903d153792 100644
--- a/libgcc/config.host
+++ b/libgcc/config.host
@@ -1256,6 +1256,10 @@ mep*-*-*)
 	tmake_file="mep/t-mep t-fdpbit"
 	extra_parts="crtbegin.o crtend.o"
 	;;
+nvptx-*)
+	tmake_file="$tmake_file nvptx/t-nvptx"
+	extra_parts="crt0.o"
+	;;
 *)
 	echo "*** Configuration ${host} not supported" 1>&2
 	exit 1
diff --git a/libgcc/config/nvptx/crt0.s b/libgcc/config/nvptx/crt0.s
new file mode 100644
index 00000000000..38327edcebd
--- /dev/null
+++ b/libgcc/config/nvptx/crt0.s
@@ -0,0 +1,45 @@
+	.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
new file mode 100644
index 00000000000..c7c56cf09b4
--- /dev/null
+++ b/libgcc/config/nvptx/free.asm
@@ -0,0 +1,50 @@
+// A wrapper around free to enable a realloc implementation.
+
+// Copyright (C) 2014 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/malloc.asm b/libgcc/config/nvptx/malloc.asm
new file mode 100644
index 00000000000..9d9db10a9e5
--- /dev/null
+++ b/libgcc/config/nvptx/malloc.asm
@@ -0,0 +1,55 @@
+// A wrapper around malloc to enable a realloc implementation.
+
+// Copyright (C) 2014 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/nvptx-malloc.h b/libgcc/config/nvptx/nvptx-malloc.h
new file mode 100644
index 00000000000..137d73c6d53
--- /dev/null
+++ b/libgcc/config/nvptx/nvptx-malloc.h
@@ -0,0 +1,26 @@
+/* Declarations for the malloc wrappers.
+
+   Copyright (C) 2014 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/>.  */
+
+extern void __nvptx_free (void *);
+extern void *__nvptx_malloc (size_t);
+extern void *__nvptx_realloc (void *, size_t);
diff --git a/libgcc/config/nvptx/realloc.c b/libgcc/config/nvptx/realloc.c
new file mode 100644
index 00000000000..41cf55443ff
--- /dev/null
+++ b/libgcc/config/nvptx/realloc.c
@@ -0,0 +1,51 @@
+/* Implement realloc with the help of the malloc and free wrappers.
+
+   Copyright (C) 2014 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 <stdlib.h>
+#include <string.h>
+#include "nvptx-malloc.h"
+
+void *
+__nvptx_realloc (void *ptr, size_t newsz)
+{
+  if (newsz == 0)
+    {
+      __nvptx_free (ptr);
+      return NULL;
+    }
+  void *newptr = __nvptx_malloc (newsz);
+
+  size_t oldsz;
+  if (ptr == NULL)
+    oldsz = 0;
+  else
+    {
+      size_t *sp = __extension__ (size_t *)(ptr - 8);
+      oldsz = *sp;
+    }
+  if (oldsz != 0)
+    memcpy (newptr, ptr, oldsz > newsz ? newsz : oldsz);
+
+  __nvptx_free (ptr);
+  return newptr;
+}
diff --git a/libgcc/config/nvptx/t-nvptx b/libgcc/config/nvptx/t-nvptx
new file mode 100644
index 00000000000..08d3a67189e
--- /dev/null
+++ b/libgcc/config/nvptx/t-nvptx
@@ -0,0 +1,9 @@
+LIB2ADD=$(srcdir)/config/nvptx/malloc.asm \
+	$(srcdir)/config/nvptx/free.asm \
+	$(srcdir)/config/nvptx/realloc.c
+
+LIB2ADDEH=
+LIB2FUNCS_EXCLUDE=__main
+
+crt0.o: $(srcdir)/config/nvptx/crt0.s
+	cp $< $@
diff --git a/libgcc/shared-object.mk b/libgcc/shared-object.mk
index d9ee922c96b..efac7973b2c 100644
--- a/libgcc/shared-object.mk
+++ b/libgcc/shared-object.mk
@@ -24,13 +24,15 @@ $(error Unsupported file type: $o)
 endif
 endif
 
+as_flags-$o := -xassembler$(if $(filter .S,$(suffix $o)),-with-cpp)
+
 $(base)$(objext): $o $(base).vis
-	$(gcc_compile) -c -xassembler-with-cpp -include $*.vis $<
+	$(gcc_compile) -c $(as_flags-$<) -include $*.vis $<
 
 $(base).vis: $(base)_s$(objext)
 	$(gen-hide-list)
 
 $(base)_s$(objext): $o
-	$(gcc_s_compile) -c -xassembler-with-cpp $<
+	$(gcc_s_compile) -c $(as_flags-$<) $<
 
 endif
diff --git a/libgcc/static-object.mk b/libgcc/static-object.mk
index 4f536369f0d..891787e85ac 100644
--- a/libgcc/static-object.mk
+++ b/libgcc/static-object.mk
@@ -24,13 +24,15 @@ $(error Unsupported file type: $o)
 endif
 endif
 
+as_flags-$o := -xassembler$(if $(filter .S,$(suffix $o)),-with-cpp)
+
 $(base)$(objext): $o $(base).vis
-	$(gcc_compile) -c -xassembler-with-cpp -include $*.vis $<
+	$(gcc_compile) -c $(as_flags-$<) -include $*.vis $<
 
 $(base).vis: $(base)_s$(objext)
 	$(gen-hide-list)
 
 $(base)_s$(objext): $o
-	$(gcc_s_compile) -c -xassembler-with-cpp $<
+	$(gcc_s_compile) -c $(as_flags-$<) $<
 
 endif