contrib/gcc-5.0/libgcc/config/nvptx/crt0.s

   1         .version 3.1
   2         .target sm_30
   3         .address_size 64
   4
   5 .global .u64 %__exitval;
   6 // BEGIN GLOBAL FUNCTION DEF: abort
   7 .visible .func abort
   8 {
   9         .reg .u64 %rd1;
  10         ld.global.u64   %rd1,[%__exitval];
  11         st.u32   [%rd1], 255;
  12         exit;
  13 }
  14 // BEGIN GLOBAL FUNCTION DEF: exit
  15 .visible .func exit (.param .u32 %arg)
  16 {
  17         .reg .u64 %rd1;
  18         .reg .u32 %val;
  19         ld.param.u32 %val,[%arg];
  20         ld.global.u64   %rd1,[%__exitval];
  21         st.u32   [%rd1], %val;
  22         exit;
  23 }
  24
  25 .extern .func (.param.u32 retval) main (.param.u32 argc, .param.u64 argv);
  26
  27 .visible .entry __main (.param .u64 __retval, .param.u32 __argc, .param.u64 __argv)
  28 {
  29         .reg .u32 %r<3>;
  30         .reg .u64 %rd<3>;
  31         .param.u32 %argc;
  32         .param.u64 %argp;
  33         .param.u32 %mainret;
  34         ld.param.u64    %rd0, [__retval];
  35         st.global.u64   [%__exitval], %rd0;
  36
  37         ld.param.u32    %r1, [__argc];
  38         ld.param.u64    %rd1, [__argv];
  39         st.param.u32    [%argc], %r1;
  40         st.param.u64    [%argp], %rd1;
  41         call.uni        (%mainret), main, (%argc, %argp);
  42         ld.param.u32    %r1,[%mainret];
  43         st.s32   [%rd0], %r1;
  44         exit;
  45 }