libs/newlib-cygwin
libs
/
newlib-cygwin
mirror of git://sourceware.org/git/newlib-cygwin.git
4
0
Fork 0
Code Issues Projects Releases Wiki Activity
newlib-cygwin/newlib/libc/machine/amdgcn/getreent.c

101 lines
2.8 KiB
C
Raw Normal View History

AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
/* get thread-specific reentrant pointer */
#include <reent.h>
#include <stdint.h>
#include <stdlib.h>
Add missing includes. These missing includes were causing build warnings, but also a real bug in which the "size" parameter to "write" was being passed in 32-bit, whereas it ought to be 64-bit. This led to intermittent bad behaviour.
2019-03-25 23:33:53 +08:00
#include <unistd.h>
AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
/* 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)
{
amdgcn: Use __builtin_gcn_ in libc/machine/amdgcn/getreent.c Call __builtin_gcn_get_stack_limit and __builtin_gcn_first_call_this_thread_p to reduce dependency on some register/layout assumptions by using the new GCC mainline (GCC 13) builtins, if they are available. If not, the existing code is used.
2022-11-22 21:53:48 +08:00
/* Place the reent data at the top of the stack allocation. */
struct data {
int marker;
struct _reent reent;
} *data;
#if defined(__has_builtin) \
&& __has_builtin(__builtin_gcn_get_stack_limit) \
&& __has_builtin(__builtin_gcn_first_call_this_thread_p)
unsigned long addr = (((unsigned long) __builtin_gcn_get_stack_limit()
- sizeof(struct data)) & ~7);
data = (struct data *)addr;
register long sp asm("s16");
if (sp >= addr)
goto stackoverflow;
if (__builtin_gcn_first_call_this_thread_p())
{
data->marker = 12345;
__builtin_memset (&data->reent, 0, sizeof(struct _reent));
_REENT_INIT_PTR_ZEROED (&data->reent);
}
else if (data->marker != 12345)
goto stackoverflow;
#else
/* s[0:1] contains a 48-bit private segment base address.
AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
s11 contains the offset to the base of the stack.
s[4:5] contains the dispatch pointer.
amdgcn: Use __builtin_gcn_ in libc/machine/amdgcn/getreent.c Call __builtin_gcn_get_stack_limit and __builtin_gcn_first_call_this_thread_p to reduce dependency on some register/layout assumptions by using the new GCC mainline (GCC 13) builtins, if they are available. If not, the existing code is used.
2022-11-22 21:53:48 +08:00
Stash reent marker in upper bits of s1 on AMD GCN s[0:3] contain a descriptor used to set up the initial value of the stack, but only the lower 48 bits of s[0:1] are currently used. The reent marker is currently set in s3, but by stashing it in the upper 16 bits of s[0:1] instead, s3 can be freed up for other purposes.
2019-11-08 06:46:41 +08:00
WARNING: this code will break if s[0:1] is ever used for anything! */
Fix __getreent stack calculations for AMD GCN From: Andrew Stubbs <ams@codesourcery.com> Fix a bug in which the high-part of 64-bit values are being corrupted, leading to erroneous stack overflow errors. The problem was only that the mixed-size calculations are being treated as signed when they should be unsigned.
2019-06-08 01:57:45 +08:00
const register unsigned long buffer_descriptor asm("s0");
unsigned long private_segment = buffer_descriptor & 0x0000ffffffffffff;
const register unsigned int stack_offset asm("s11");
AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
const register hsa_kernel_dispatch_packet_t *dispatch_ptr asm("s4");
Fix __getreent stack calculations for AMD GCN From: Andrew Stubbs <ams@codesourcery.com> Fix a bug in which the high-part of 64-bit values are being corrupted, leading to erroneous stack overflow errors. The problem was only that the mixed-size calculations are being treated as signed when they should be unsigned.
2019-06-08 01:57:45 +08:00
unsigned long stack_base = private_segment + stack_offset;
unsigned long stack_end = stack_base + dispatch_ptr->private_segment_size * 64;
unsigned long addr = (stack_end - sizeof(struct data)) & ~7;
AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
data = (struct data *)addr;
register long sp asm("s16");
if (sp >= addr)
goto stackoverflow;
Stash reent marker in upper bits of s1 on AMD GCN s[0:3] contain a descriptor used to set up the initial value of the stack, but only the lower 48 bits of s[0:1] are currently used. The reent marker is currently set in s3, but by stashing it in the upper 16 bits of s[0:1] instead, s3 can be freed up for other purposes.
2019-11-08 06:46:41 +08:00
/* Stash a marker in the unused upper 16 bits of s[0:1] to indicate that
the reent data is initialized. */
const register unsigned int s1 asm("s1");
unsigned int marker = s1 >> 16;
if (marker != 12345)
AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
{
Stash reent marker in upper bits of s1 on AMD GCN s[0:3] contain a descriptor used to set up the initial value of the stack, but only the lower 48 bits of s[0:1] are currently used. The reent marker is currently set in s3, but by stashing it in the upper 16 bits of s[0:1] instead, s3 can be freed up for other purposes.
2019-11-08 06:46:41 +08:00
asm("s_and_b32\ts1, s1, 0xffff");
asm("s_or_b32\ts1, s1, (12345 << 16)");
data->marker = 12345;
AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
__builtin_memset (&data->reent, 0, sizeof(struct _reent));
_REENT_INIT_PTR_ZEROED (&data->reent);
}
Stash reent marker in upper bits of s1 on AMD GCN s[0:3] contain a descriptor used to set up the initial value of the stack, but only the lower 48 bits of s[0:1] are currently used. The reent marker is currently set in s3, but by stashing it in the upper 16 bits of s[0:1] instead, s3 can be freed up for other purposes.
2019-11-08 06:46:41 +08:00
else if (data->marker != 12345)
AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
goto stackoverflow;
amdgcn: Use __builtin_gcn_ in libc/machine/amdgcn/getreent.c Call __builtin_gcn_get_stack_limit and __builtin_gcn_first_call_this_thread_p to reduce dependency on some register/layout assumptions by using the new GCC mainline (GCC 13) builtins, if they are available. If not, the existing code is used.
2022-11-22 21:53:48 +08:00
#endif
AMD GCN Port contributed by Andrew Stubbs <ams@codesourcery.com> 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.
2019-01-15 23:36:50 +08:00
return &data->reent;
stackoverflow:
write (2, "GCN Stack Overflow!\n", 20);
abort ();
}