.version 6.5
.target sm_30
.address_size 64

.visible .entry stateful_ld_st_ntid_sub(
    .param .u64 input,
    .param .u64 output
)
{
    .reg .b64           in_addr1;
    .reg .b64           in_addr2;
    .reg .b64           in_addr3;
    .reg .b64           out_addr1;
    .reg .b64           out_addr2;
    .reg .b64           out_addr3;
    .reg .u32           tid_32;
    .reg .u64           tid_64;
    .reg .u64           temp;

    ld.param.u64        in_addr1, [input];
    ld.param.u64        out_addr1, [output];

    cvta.to.global.u64  in_addr2, in_addr1;
    cvta.to.global.u64  out_addr2, out_addr1;

    mov.u32             tid_32, %tid.x;
    cvt.u64.u32         tid_64, tid_32;

    sub.s64             in_addr3, in_addr2, tid_64;
    sub.s64             out_addr3, out_addr2, tid_64;

    ld.global.u64       temp, [in_addr3+-0];
    st.global.u64       [out_addr3+-0], temp;
    ret;
}