newlib/libc/machine/amdgcn/getreent.c

   1 /* get thread-specific reentrant pointer */
   2
   3 #include <reent.h>
   4 #include <stdint.h>
   5 #include <stdlib.h>
   6 #include <unistd.h>
   7
   8 /* Copied from the HSA documentation.  */
   9 typedef struct hsa_signal_s {
  10   uint64_t handle;
  11 } hsa_signal_t;
  12 typedef struct hsa_kernel_dispatch_packet_s {
  13   uint16_t header ;
  14   uint16_t setup;
  15   uint16_t workgroup_size_x ;
  16   uint16_t workgroup_size_y ;
  17   uint16_t workgroup_size_z;
  18   uint16_t reserved0;
  19   uint32_t grid_size_x ;
  20   uint32_t grid_size_y ;
  21   uint32_t grid_size_z;
  22   uint32_t private_segment_size;
  23   uint32_t group_segment_size;
  24   uint64_t kernel_object;
  25   uint64_t reserved2;
  26   hsa_signal_t completion_signal;
  27 } hsa_kernel_dispatch_packet_t;
  28
  29 struct _reent *
  30 __getreent (void)
  31 {
  32   /* Place the reent data at the top of the stack allocation.  */
  33   struct data {
  34     int marker;
  35     struct _reent reent;
  36   } *data;
  37
  38 #if defined(__has_builtin) \
  39     && __has_builtin(__builtin_gcn_get_stack_limit) \
  40     && __has_builtin(__builtin_gcn_first_call_this_thread_p)
  41   unsigned long addr = (((unsigned long) __builtin_gcn_get_stack_limit()
  42                          - sizeof(struct data)) & ~7);
  43   data = (struct data *)addr;
  44
  45   register long sp asm("s16");
  46
  47   if (sp >= addr)
  48     goto stackoverflow;
  49   if (__builtin_gcn_first_call_this_thread_p())
  50     {
  51       data->marker = 12345;
  52       __builtin_memset (&data->reent, 0, sizeof(struct _reent));
  53       _REENT_INIT_PTR_ZEROED (&data->reent);
  54     }
  55   else if (data->marker != 12345)
  56     goto stackoverflow;
  57 #else
  58   /* s[0:1] contains a 48-bit private segment base address.
  59      s11 contains the offset to the base of the stack.
  60      s[4:5] contains the dispatch pointer.
  61
  62      WARNING: this code will break if s[0:1] is ever used for anything!  */
  63   const register unsigned long buffer_descriptor asm("s0");
  64   unsigned long private_segment = buffer_descriptor & 0x0000ffffffffffff;
  65   const register unsigned int stack_offset asm("s11");
  66   const register hsa_kernel_dispatch_packet_t *dispatch_ptr asm("s4");
  67
  68   unsigned long stack_base = private_segment + stack_offset;
  69   unsigned long stack_end = stack_base + dispatch_ptr->private_segment_size * 64;
  70   unsigned long addr = (stack_end - sizeof(struct data)) & ~7;
  71   data = (struct data *)addr;
  72
  73   register long sp asm("s16");
  74   if (sp >= addr)
  75     goto stackoverflow;
  76
  77   /* Stash a marker in the unused upper 16 bits of s[0:1] to indicate that
  78      the reent data is initialized.  */
  79   const register unsigned int s1 asm("s1");
  80   unsigned int marker = s1 >> 16;
  81   if (marker != 12345)
  82     {
  83       asm("s_and_b32\ts1, s1, 0xffff");
  84       asm("s_or_b32\ts1, s1, (12345 << 16)");
  85       data->marker = 12345;
  86
  87       __builtin_memset (&data->reent, 0, sizeof(struct _reent));
  88       _REENT_INIT_PTR_ZEROED (&data->reent);
  89     }
  90   else if (data->marker != 12345)
  91     goto stackoverflow;
  92 #endif
  93
  94   return &data->reent;
  95
  96 stackoverflow:
  97     write (2, "GCN Stack Overflow!\n", 20);
  98     abort ();
  99 }
 100