.version 6.5
.target sm_30
.address_size 64

.visible .entry reg_local(
	.param .u64 input,
	.param .u64 output
)
{
    .local .align 8 .b8 local_x[8];
	.reg .u64 	    in_addr;
    .reg .u64 	    out_addr;
    .reg .b64 	    temp;

	ld.param.u64 	in_addr, [input];
    ld.param.u64 	out_addr, [output];

    ld.global.u64   temp, [in_addr];
    st.u64          [local_x], temp + 1;
    ld.u64          temp, [local_x+0];
    st.global.u64   [out_addr+0], temp;
	ret;
}