mirror of
git://sourceware.org/git/newlib-cygwin.git
synced 2025-01-22 06:57:36 +08:00
1787e9d033
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.
80 lines
2.2 KiB
C
80 lines
2.2 KiB
C
/* 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 ();
|
|
}
|
|
|