This is the mail archive of the
newlib@sourceware.org
mailing list for the newlib project.
Re: [PATCH] AMD GCN Port
On 10/01/2019 19:22, Jeff Johnston wrote:
Try libc/sys/sh as a basis. You have to replace the creat.c in
configure.in <http://configure.in> with a file you have.
Don't specify crt0.o in Makefile.am and put in your source files.
Create an empty sys directory under
you libc/sys/xxxx dir to possibly use in the future to replace sys includes.
Other than that, most of your patch from machine dir goes there. The
libc/machine/xxxx dir should have anything
you are overriding in shared library (often optimized string routines,
but at minimum, setjmp/longjmp
(see libc/machine/fr30).
How about the attached?
I've moved the syscalls, but left the machine overrides in place, I think.
We don't have an implementation for setjmp/longjmp yet, although we'll
need to look into that when we build out more C++ support. The register
file on this architecture is "up to" ~64kB, depending on the thread
configuration, and there's no well-defined ABI yet (just what the
compiler implements), so it's trickier than average.
Once again, I've omitted the autogenerated files from the patch file.
I've also changed to a commit message more in keeping with the new style.
Andrew
AMD GCN Port
Add support for the AMD GCN GPU architecture. This is primarily intended for
use with OpenMP and OpenACC offloading. It can also be used for stand-alone
programs, but this is intended mostly for testing the compiler and is not
expected to be useful in general.
The GPU architecture is highly parallel, and therefore Newlib must be
configured to use dynamic re-entrancy, and thread-safe malloc.
The only I/O available is a via a shared-memory interface provided by libgomp
and the gcn-run tool included with GCC. At this time this is limited to
stdout, argc/argv, and the return code.
diff --git a/newlib/configure.host b/newlib/configure.host
index 6c49cb7..fa805d6 100644
--- a/newlib/configure.host
+++ b/newlib/configure.host
@@ -118,6 +118,10 @@ case "${host_cpu}" in
machine_dir=aarch64
libm_machine_dir=aarch64
;;
+ amdgcn*)
+ newlib_cflags="${newlib_cflags} -D__DYNAMIC_REENT__"
+ machine_dir=amdgcn
+ ;;
arc*)
machine_dir=arc
;;
@@ -442,6 +446,10 @@ case "${host}" in
aarch64*-*-*)
newlib_cflags="${newlib_cflags} -D_COMPILING_NEWLIB"
;;
+ amdgcn*)
+ sys_dir=amdgcn
+ have_crt0="no"
+ ;;
arm*-*-*)
newlib_cflags="${newlib_cflags} -D_COMPILING_NEWLIB"
sys_dir=arm
diff --git a/newlib/libc/include/machine/ieeefp.h b/newlib/libc/include/machine/ieeefp.h
index a409752..911eeb5 100644
--- a/newlib/libc/include/machine/ieeefp.h
+++ b/newlib/libc/include/machine/ieeefp.h
@@ -452,6 +452,10 @@
#define __IEEE_BIG_ENDIAN
#endif
+#ifdef __AMDGCN__
+#define __IEEE_LITTLE_ENDIAN
+#endif
+
#ifdef __CYGWIN__
#define __OBSOLETE_MATH_DEFAULT 0
#endif
diff --git a/newlib/libc/include/sys/config.h b/newlib/libc/include/sys/config.h
index 49b62eb..d746b15 100644
--- a/newlib/libc/include/sys/config.h
+++ b/newlib/libc/include/sys/config.h
@@ -8,6 +8,10 @@
#define MALLOC_ALIGNMENT 16
#endif
+#ifdef __AMDGCN__
+#define __DYNAMIC_REENT__
+#endif
+
/* exceptions first */
#if defined(__H8500__) || defined(__W65__)
#define __SMALL_BITFIELDS
diff --git a/newlib/libc/machine/amdgcn/Makefile.am b/newlib/libc/machine/amdgcn/Makefile.am
new file mode 100644
index 0000000..4d8d2d1
--- /dev/null
+++ b/newlib/libc/machine/amdgcn/Makefile.am
@@ -0,0 +1,15 @@
+## Process this file with automake to generate Makefile.in
+
+AUTOMAKE_OPTIONS = cygnus
+
+INCLUDES = $(NEWLIB_CFLAGS) $(CROSS_CFLAGS) $(TARGET_CFLAGS)
+
+AM_CCASFLAGS = $(INCLUDES)
+
+noinst_LIBRARIES = lib.a
+
+lib_a_SOURCES = abort.c exit.c atexit.c malloc_support.c getreent.c
+lib_a_CFLAGS = $(AM_CFLAGS)
+
+ACLOCAL_AMFLAGS = -I ../../.. -I ../../../..
+CONFIG_STATUS_DEPENDENCIES = $(newlib_basedir)/configure.host
diff --git a/newlib/libc/machine/amdgcn/abort.c b/newlib/libc/machine/amdgcn/abort.c
new file mode 100644
index 0000000..ccbca72
--- /dev/null
+++ b/newlib/libc/machine/amdgcn/abort.c
@@ -0,0 +1,25 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2014-2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <stdlib.h>
+#include <signal.h>
+#include "exit-value.h"
+
+void __attribute__((noreturn))
+abort (void)
+{
+ write (2, "GCN Kernel Aborted\n", 19);
+ exit_with_status_and_signal (0, SIGABRT);
+}
diff --git a/newlib/libc/machine/amdgcn/atexit.c b/newlib/libc/machine/amdgcn/atexit.c
new file mode 100644
index 0000000..6745714
--- /dev/null
+++ b/newlib/libc/machine/amdgcn/atexit.c
@@ -0,0 +1,25 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2014-2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <stdlib.h>
+
+int
+atexit (void (*function)(void))
+{
+ /* Our current implementation of exit does not run functions registered with
+ atexit, so fail here. */
+ abort ();
+ return 1;
+}
diff --git a/newlib/libc/machine/amdgcn/configure.in b/newlib/libc/machine/amdgcn/configure.in
new file mode 100644
index 0000000..028e9d7
--- /dev/null
+++ b/newlib/libc/machine/amdgcn/configure.in
@@ -0,0 +1,14 @@
+dnl This is the newlib/libc/machine/amdgcn configure.in file.
+dnl Process this file with autoconf to produce a configure script.
+
+AC_PREREQ(2.59)
+AC_INIT([newlib],[NEWLIB_VERSION])
+AC_CONFIG_SRCDIR([Makefile.am])
+
+dnl Can't be done in NEWLIB_CONFIGURE because that confuses automake.
+AC_CONFIG_AUX_DIR(../../../..)
+
+NEWLIB_CONFIGURE(../../..)
+
+AC_CONFIG_FILES([Makefile])
+AC_OUTPUT
diff --git a/newlib/libc/machine/amdgcn/exit-value.h b/newlib/libc/machine/amdgcn/exit-value.h
new file mode 100644
index 0000000..6e88625
--- /dev/null
+++ b/newlib/libc/machine/amdgcn/exit-value.h
@@ -0,0 +1,48 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#ifndef _AMDGCN_EXIT_VALUE_H_
+#define _AMDGCN_EXIT_VALUE_H_
+
+static inline void __attribute__((noreturn))
+exit_with_int (int val)
+{
+ /* Write the exit value to the conventional place. */
+ int *return_value;
+ asm ("s_load_dwordx2 %0, s[8:9], 16 glc\n\t"
+ "s_waitcnt 0" : "=Sg"(return_value));
+ *return_value = val;
+
+ /* Terminate the current kernel. */
+ asm ("s_dcache_wb");
+ asm ("s_endpgm");
+ __builtin_unreachable ();
+}
+
+static inline void __attribute__((noreturn))
+exit_with_status_and_signal (int val, int signal)
+{
+ if (signal == 0)
+ val = val & 0xff;
+ else
+ {
+ val = (128 + signal) & 0xff;
+ signal = signal & 0xff;
+ }
+
+ exit_with_int ((0xffff << 16) | (signal << 8) | val);
+}
+
+#endif
diff --git a/newlib/libc/machine/amdgcn/exit.c b/newlib/libc/machine/amdgcn/exit.c
new file mode 100644
index 0000000..bdd532e
--- /dev/null
+++ b/newlib/libc/machine/amdgcn/exit.c
@@ -0,0 +1,23 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2014-2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <stdlib.h>
+#include "exit-value.h"
+
+void __attribute__((noreturn))
+exit (int val)
+{
+ exit_with_status_and_signal (val, 0);
+}
diff --git a/newlib/libc/machine/amdgcn/getreent.c b/newlib/libc/machine/amdgcn/getreent.c
new file mode 100644
index 0000000..acf10a9
--- /dev/null
+++ b/newlib/libc/machine/amdgcn/getreent.c
@@ -0,0 +1,79 @@
+/* get thread-specific reentrant pointer */
+
+#include <reent.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+/* Copied from the HSA documentation. */
+typedef struct hsa_signal_s {
+ uint64_t handle;
+} hsa_signal_t;
+typedef struct hsa_kernel_dispatch_packet_s {
+ uint16_t header ;
+ uint16_t setup;
+ uint16_t workgroup_size_x ;
+ uint16_t workgroup_size_y ;
+ uint16_t workgroup_size_z;
+ uint16_t reserved0;
+ uint32_t grid_size_x ;
+ uint32_t grid_size_y ;
+ uint32_t grid_size_z;
+ uint32_t private_segment_size;
+ uint32_t group_segment_size;
+ uint64_t kernel_object;
+ uint64_t reserved2;
+ hsa_signal_t completion_signal;
+} hsa_kernel_dispatch_packet_t;
+
+struct _reent *
+__getreent (void)
+{
+ /* Place the reent data at the top of the stack allocation.
+ s[0:1] contains a 48-bit private segment base address.
+ s11 contains the offset to the base of the stack.
+ s[4:5] contains the dispatch pointer.
+
+ WARNING: this code will break if s[0:3] is ever used for anything! */
+ const register long buffer_descriptor asm("s0");
+ long private_segment = buffer_descriptor & 0x0000ffffffffffff;
+ const register int stack_offset asm("s11");
+ const register hsa_kernel_dispatch_packet_t *dispatch_ptr asm("s4");
+
+ struct data {
+ int marker;
+ struct _reent reent;
+ } *data;
+
+ long stack_base = private_segment + stack_offset;
+ long stack_end = stack_base + dispatch_ptr->private_segment_size * 64;
+ long addr = (stack_end - sizeof(struct data)) & ~7;
+ data = (struct data *)addr;
+
+ register long sp asm("s16");
+ if (sp >= addr)
+ goto stackoverflow;
+
+ /* Place a marker in s3 to indicate that the reent data is initialized.
+ The register is known to hold part of an unused buffer descriptor
+ when the kernel is launched. This may not be unused forever, but
+ we already used s0 and s1 above, so this doesn't do extra harm. */
+ register int s3 asm("s3");
+ if (s3 != 123456)
+ {
+ asm("s_mov_b32 s3, 123456");
+ data->marker = 123456;
+
+ __builtin_memset (&data->reent, 0, sizeof(struct _reent));
+ _REENT_INIT_PTR_ZEROED (&data->reent);
+ }
+ else if (data->marker != 123456)
+ goto stackoverflow;
+
+
+ return &data->reent;
+
+stackoverflow:
+ write (2, "GCN Stack Overflow!\n", 20);
+ abort ();
+}
+
diff --git a/newlib/libc/machine/amdgcn/malloc_support.c b/newlib/libc/machine/amdgcn/malloc_support.c
new file mode 100644
index 0000000..4848c97
--- /dev/null
+++ b/newlib/libc/machine/amdgcn/malloc_support.c
@@ -0,0 +1,111 @@
+/*
+ * Support file for AMDGCN in newlib.
+ * Copyright (c) 2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <stdlib.h>
+#include <stdint.h>
+#include <reent.h>
+
+/* _sbrk_r expects us to use the real errno, not the reentrant one. */
+#include <errno.h>
+#undef errno
+extern int errno;
+
+/* The runtime passes in heap space like this. */
+struct heap {
+ int64_t size;
+ char data[0];
+};
+
+static char *__heap_ptr = (char*)-1;
+static char *__heap_end = (char*)-1;
+static int __heap_lock = 0;
+static void *__heap_lock_id = NULL;
+static int __heap_lock_cnt = 0;
+
+void *
+sbrk (ptrdiff_t nbytes)
+{
+ if (__heap_ptr == (char *)-1)
+ {
+ /* Find the heap from kernargs.
+ The kernargs pointer is in s[8:9].
+ This will break if the enable_sgpr_* flags are ever changed. */
+ char *kernargs;
+ asm ("s_mov_b64 %0, s[8:9]" : "=Sg"(kernargs));
+
+ /* The heap data is at kernargs[3]. */
+ struct heap *heap = *(struct heap **)(kernargs + 24);
+
+ __heap_ptr = heap->data;
+ __heap_end = __heap_ptr + heap->size;
+ }
+
+ if ((__heap_ptr + nbytes) >= __heap_end)
+ {
+ errno = ENOMEM;
+ return (void*)-1;
+ }
+
+ char *base = __heap_ptr;
+ __heap_ptr += nbytes;
+
+ return base;
+}
+
+void
+__malloc_lock (struct _reent *reent)
+{
+ void *id = reent;
+
+ if (id == __heap_lock_id)
+ {
+ if (__heap_lock_cnt < 1)
+ abort ();
+ ++__heap_lock_cnt;
+ return;
+ }
+
+ while (__sync_lock_test_and_set (&__heap_lock, 1))
+ /* A sleep seems like it should allow the wavefront to yeild (maybe?)
+ Use the shortest possible sleep time of 1*64 cycles. */
+ asm volatile ("s_sleep\t1" ::: "memory");
+
+ if (__heap_lock_id != NULL)
+ abort ();
+ if (__heap_lock_cnt != 0)
+ abort ();
+
+ __heap_lock_cnt = 1;
+ __heap_lock_id = id;
+}
+
+void
+__malloc_unlock (struct _reent *reent)
+{
+ void *id = reent;
+
+ if (id != __heap_lock_id)
+ abort ();
+ if (__heap_lock_cnt < 1)
+ abort ();
+
+ --__heap_lock_cnt;
+
+ if (__heap_lock_cnt > 0)
+ return;
+
+ __heap_lock_id = NULL;
+ __sync_lock_release (&__heap_lock);
+}
diff --git a/newlib/libc/machine/configure.in b/newlib/libc/machine/configure.in
index 8ebe68b..0d4068c 100644
--- a/newlib/libc/machine/configure.in
+++ b/newlib/libc/machine/configure.in
@@ -25,6 +25,7 @@ if test -n "${machine_dir}"; then
case ${machine_dir} in
a29k) AC_CONFIG_SUBDIRS(a29k) ;;
aarch64) AC_CONFIG_SUBDIRS(aarch64) ;;
+ amdgcn) AC_CONFIG_SUBDIRS(amdgcn) ;;
arc) AC_CONFIG_SUBDIRS(arc) ;;
arm) AC_CONFIG_SUBDIRS(arm) ;;
bfin) AC_CONFIG_SUBDIRS(bfin) ;;
diff --git a/newlib/libc/ssp/stack_protector.c b/newlib/libc/ssp/stack_protector.c
index ee014b6..cd51543 100644
--- a/newlib/libc/ssp/stack_protector.c
+++ b/newlib/libc/ssp/stack_protector.c
@@ -5,6 +5,11 @@
#include <string.h>
#include <unistd.h>
+#if defined(__AMDGCN__)
+/* GCN does not support constructors, yet. */
+uintptr_t __stack_chk_guard = 0x00000aff; /* 0, 0, '\n', 255 */
+
+#else
uintptr_t __stack_chk_guard = 0;
void
@@ -24,6 +29,7 @@ __stack_chk_init (void)
((unsigned char *)&__stack_chk_guard)[3] = 255;
#endif
}
+#endif
void
__attribute__((__noreturn__))
diff --git a/newlib/libc/sys/amdgcn/Makefile.am b/newlib/libc/sys/amdgcn/Makefile.am
new file mode 100644
index 0000000..1716776
--- /dev/null
+++ b/newlib/libc/sys/amdgcn/Makefile.am
@@ -0,0 +1,16 @@
+## Process this file with automake to generate Makefile.in
+
+AUTOMAKE_OPTIONS = cygnus
+
+INCLUDES = $(NEWLIB_CFLAGS) $(CROSS_CFLAGS) $(TARGET_CFLAGS)
+
+AM_CCASFLAGS = $(INCLUDES) $(CFLAGS)
+
+noinst_LIBRARIES = lib.a
+
+lib_a_SOURCES = close.c fstat.c isatty.c lseek.c read.c write.c
+lib_a_CCASFLAGS = $(AM_CCASFLAGS)
+lib_a_CFLAGS = $(AM_CFLAGS)
+
+ACLOCAL_AMFLAGS = -I ../../.. -I ../../../..
+CONFIG_STATUS_DEPENDENCIES = $(newlib_basedir)/configure.host
diff --git a/newlib/libc/sys/amdgcn/close.c b/newlib/libc/sys/amdgcn/close.c
new file mode 100644
index 0000000..5bce557
--- /dev/null
+++ b/newlib/libc/sys/amdgcn/close.c
@@ -0,0 +1,24 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <unistd.h>
+#include <errno.h>
+
+int close(int fildes)
+{
+ errno = EIO;
+ return -1;
+}
+
diff --git a/newlib/libc/sys/amdgcn/configure.in b/newlib/libc/sys/amdgcn/configure.in
new file mode 100644
index 0000000..74edb0a
--- /dev/null
+++ b/newlib/libc/sys/amdgcn/configure.in
@@ -0,0 +1,14 @@
+dnl This is the newlib/libc/sys/amdgcn configure.in file.
+dnl Process this file with autoconf to produce a configure script.
+
+AC_PREREQ(2.59)
+AC_INIT([newlib],[NEWLIB_VERSION])
+AC_CONFIG_SRCDIR([close.c])
+
+dnl Can't be done in NEWLIB_CONFIGURE because that confuses automake.
+AC_CONFIG_AUX_DIR(../../../..)
+
+NEWLIB_CONFIGURE(../../..)
+
+AC_CONFIG_FILES([Makefile])
+AC_OUTPUT
diff --git a/newlib/libc/sys/amdgcn/fstat.c b/newlib/libc/sys/amdgcn/fstat.c
new file mode 100644
index 0000000..b787158
--- /dev/null
+++ b/newlib/libc/sys/amdgcn/fstat.c
@@ -0,0 +1,23 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <unistd.h>
+#include <errno.h>
+
+int fstat(int fildes, struct stat *buf)
+{
+ errno = EIO;
+ return -1;
+}
diff --git a/newlib/libc/sys/amdgcn/isatty.c b/newlib/libc/sys/amdgcn/isatty.c
new file mode 100644
index 0000000..4268f2c
--- /dev/null
+++ b/newlib/libc/sys/amdgcn/isatty.c
@@ -0,0 +1,23 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <unistd.h>
+#include <errno.h>
+
+int isatty(int fd)
+{
+ errno = EINVAL;
+ return 0;
+}
diff --git a/newlib/libc/sys/amdgcn/lseek.c b/newlib/libc/sys/amdgcn/lseek.c
new file mode 100644
index 0000000..be3220b
--- /dev/null
+++ b/newlib/libc/sys/amdgcn/lseek.c
@@ -0,0 +1,24 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <unistd.h>
+#include <errno.h>
+
+off_t lseek(int fildes, off_t offset, int whence)
+{
+ errno = ESPIPE;
+ return -1;
+}
+
diff --git a/newlib/libc/sys/amdgcn/read.c b/newlib/libc/sys/amdgcn/read.c
new file mode 100644
index 0000000..97385e9
--- /dev/null
+++ b/newlib/libc/sys/amdgcn/read.c
@@ -0,0 +1,21 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <stdio.h>
+
+_READ_WRITE_RETURN_TYPE read (int fildes, void *buf, size_t nbyte)
+{
+ return 0;
+}
diff --git a/newlib/libc/sys/amdgcn/write.c b/newlib/libc/sys/amdgcn/write.c
new file mode 100644
index 0000000..ce5bd36
--- /dev/null
+++ b/newlib/libc/sys/amdgcn/write.c
@@ -0,0 +1,88 @@
+/*
+ * Support file for amdgcn in newlib.
+ * Copyright (c) 2014, 2017 Mentor Graphics.
+ *
+ * The authors hereby grant permission to use, copy, modify, distribute,
+ * and license this software and its documentation for any purpose, provided
+ * that existing copyright notices are retained in all copies and that this
+ * notice is included verbatim in any distributions. No written agreement,
+ * license, or royalty fee is required for any of the authorized uses.
+ * Modifications to this software may be copyrighted by their authors
+ * and need not follow the licensing terms described here, provided that
+ * the new terms are clearly indicated on the first page of each file where
+ * they apply.
+ */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <unistd.h>
+#include <errno.h>
+#include <string.h>
+
+/* This struct must match the one used by gcn-run and libgomp.
+ It holds all the data output from a kernel (besides mapping data).
+
+ The base address pointer can be found at kernargs+16.
+
+ The next_output counter must be atomically incremented for each
+ print output. Only when the print data is fully written can the
+ "written" flag be set. */
+struct output {
+ int return_value;
+ int next_output;
+ struct printf_data {
+ int written;
+ char msg[128];
+ int type;
+ union {
+ int64_t ivalue;
+ double dvalue;
+ char text[128];
+ };
+ } queue[1000];
+};
+
+_READ_WRITE_RETURN_TYPE write (int fd, const void *buf, size_t count)
+{
+ if (fd != 1 && fd != 2)
+ {
+ errno = EBADF;
+ return -1;
+ }
+
+ /* The output data is at ((void*)kernargs)[2]. */
+ register void **kernargs asm("s8");
+ struct output *data = (struct output *)kernargs[2];
+
+ /* Each output slot allows 256 bytes, so reserve as many as we need. */
+ int slot_count = ((count+1)/256)+1;
+ int index = __atomic_fetch_add (&data->next_output, slot_count,
+ __ATOMIC_ACQUIRE);
+ for (int c = count;
+ c >= 0 && index < 1000;
+ buf += 256, c -= 256, index++)
+ {
+ if (c < 128)
+ {
+ memcpy (data->queue[index].msg, buf, c);
+ data->queue[index].msg[c] = '\0';
+ data->queue[index].text[0] = '\0';
+ }
+ else if (c < 256)
+ {
+ memcpy (data->queue[index].msg, buf, 128);
+ memcpy (data->queue[index].text, buf+128, c-128);
+ data->queue[index].text[c-128] = '\0';
+ }
+ else
+ {
+ memcpy (data->queue[index].msg, buf, 128);
+ memcpy (data->queue[index].text, buf+128, 128);
+ }
+
+ data->queue[index].type = 3; /* Raw. */
+ __atomic_store_n (&data->queue[index].written, 1, __ATOMIC_RELEASE);
+ }
+
+ return count;
+}
diff --git a/newlib/libc/sys/configure.in b/newlib/libc/sys/configure.in
index bc6cb88..a65d1e7 100644
--- a/newlib/libc/sys/configure.in
+++ b/newlib/libc/sys/configure.in
@@ -23,6 +23,7 @@ fi
if test -n "${sys_dir}"; then
case ${sys_dir} in
a29khif) AC_CONFIG_SUBDIRS(a29khif) ;;
+ amdgcn) AC_CONFIG_SUBDIRS(amdgcn) ;;
arm) AC_CONFIG_SUBDIRS(arm) ;;
d10v) AC_CONFIG_SUBDIRS(d10v) ;;
decstation) AC_CONFIG_SUBDIRS(decstation) ;;