gcc/libgcc/config/nvptx/crt0.s

	.version 3.1
	.target	sm_30
	.address_size 64

.global .u64 %__exitval;
// BEGIN GLOBAL FUNCTION DEF: abort
.visible .func abort
{
        .reg .u64 %rd1;
        ld.global.u64   %rd1,[%__exitval];
        st.u32   [%rd1], 255;
        exit;
}
// BEGIN GLOBAL FUNCTION DEF: exit
.visible .func exit (.param .u32 %arg)
{
        .reg .u64 %rd1;
	.reg .u32 %val;
	ld.param.u32 %val,[%arg];
        ld.global.u64   %rd1,[%__exitval];
        st.u32   [%rd1], %val;
        exit;
}

.extern .func (.param.u32 retval) main (.param.u32 argc, .param.u64 argv);

.visible .entry __main (.param .u64 __retval, .param.u32 __argc, .param.u64 __argv)
{
        .reg .u32 %r<3>;
        .reg .u64 %rd<3>;
	.param.u32 %argc;
	.param.u64 %argp;
	.param.u32 %mainret;
        ld.param.u64    %rd0, [__retval];
        st.global.u64   [%__exitval], %rd0;

	ld.param.u32	%r1, [__argc];
	ld.param.u64	%rd1, [__argv];
	st.param.u32	[%argc], %r1;
	st.param.u64	[%argp], %rd1;
        call.uni        (%mainret), main, (%argc, %argp);
	ld.param.u32	%r1,[%mainret];
        st.s32   [%rd0], %r1;
        exit;
}