1 /*
2 Copyright (c) 2014-2017 Mentor Graphics.
3
4 The authors hereby grant permission to use, copy, modify, distribute,
5 and license this software and its documentation for any purpose, provided
6 that existing copyright notices are retained in all copies and that this
7 notice is included verbatim in any distributions. No written agreement,
8 license, or royalty fee is required for any of the authorized uses.
9 Modifications to this software may be copyrighted by their authors
10 and need not follow the licensing terms described here, provided that
11 the new terms are clearly indicated on the first page of each file where
12 they apply.
13 */
14 /* get thread-specific reentrant pointer */
15
16 #include <reent.h>
17 #include <stdint.h>
18 #include <stdlib.h>
19 #include <unistd.h>
20
21 /* Copied from the HSA documentation. */
22 typedef struct hsa_signal_s {
23 uint64_t handle;
24 } hsa_signal_t;
25 typedef struct hsa_kernel_dispatch_packet_s {
26 uint16_t header ;
27 uint16_t setup;
28 uint16_t workgroup_size_x ;
29 uint16_t workgroup_size_y ;
30 uint16_t workgroup_size_z;
31 uint16_t reserved0;
32 uint32_t grid_size_x ;
33 uint32_t grid_size_y ;
34 uint32_t grid_size_z;
35 uint32_t private_segment_size;
36 uint32_t group_segment_size;
37 uint64_t kernel_object;
38 uint64_t reserved2;
39 hsa_signal_t completion_signal;
40 } hsa_kernel_dispatch_packet_t;
41
42 struct _reent *
__getreent(void)43 __getreent (void)
44 {
45 /* Place the reent data at the top of the stack allocation. */
46 struct data {
47 int marker;
48 struct _reent reent;
49 } *data;
50
51 #if defined(__has_builtin) \
52 && __has_builtin(__builtin_gcn_get_stack_limit) \
53 && __has_builtin(__builtin_gcn_first_call_this_thread_p)
54 unsigned long addr = (((unsigned long) __builtin_gcn_get_stack_limit()
55 - sizeof(struct data)) & ~7);
56 data = (struct data *)addr;
57
58 register long sp asm("s16");
59
60 if (sp >= addr)
61 goto stackoverflow;
62 if (__builtin_gcn_first_call_this_thread_p())
63 {
64 data->marker = 12345;
65 __builtin_memset (&data->reent, 0, sizeof(struct _reent));
66 _REENT_INIT_PTR_ZEROED (&data->reent);
67 }
68 else if (data->marker != 12345)
69 goto stackoverflow;
70 #else
71 /* s[0:1] contains a 48-bit private segment base address.
72 s11 contains the offset to the base of the stack.
73 s[4:5] contains the dispatch pointer.
74
75 WARNING: this code will break if s[0:1] is ever used for anything! */
76 const register unsigned long buffer_descriptor __asm__("s0");
77 unsigned long private_segment = buffer_descriptor & 0x0000ffffffffffff;
78 const register unsigned int stack_offset __asm__("s11");
79 const register hsa_kernel_dispatch_packet_t *dispatch_ptr __asm__("s4");
80
81 unsigned long stack_base = private_segment + stack_offset;
82 unsigned long stack_end = stack_base + dispatch_ptr->private_segment_size * 64;
83 unsigned long addr = (stack_end - sizeof(struct data)) & ~7;
84 data = (struct data *)addr;
85
86 register long sp __asm__("s16");
87 if (sp >= addr)
88 goto stackoverflow;
89
90 /* Stash a marker in the unused upper 16 bits of s[0:1] to indicate that
91 the reent data is initialized. */
92 const register unsigned int s1 __asm__("s1");
93 unsigned int marker = s1 >> 16;
94 if (marker != 12345)
95 {
96 __asm__("s_and_b32\ts1, s1, 0xffff");
97 __asm__("s_or_b32\ts1, s1, (12345 << 16)");
98 data->marker = 12345;
99
100 __builtin_memset (&data->reent, 0, sizeof(struct _reent));
101 _REENT_INIT_PTR_ZEROED (&data->reent);
102 }
103 else if (data->marker != 12345)
104 goto stackoverflow;
105 #endif
106
107 return &data->reent;
108
109 stackoverflow:
110 write (2, "GCN Stack Overflow!\n", 20);
111 abort ();
112 }
113
114
