.version 6.5
.target sm_30
.address_size 64

.visible .entry vector_extract(
    .param .u64 input_p,
    .param .u64 output_p
)
{
    .reg .u64           in_addr;
    .reg .u64           out_addr;
    .reg .u16           temp1;
    .reg .u16           temp2;
    .reg .u16           temp3;
    .reg .u16           temp4;
    .reg .v4.u16        foo;

	ld.param.u64    in_addr, [input_p];
    ld.param.u64    out_addr, [output_p];

    ld.global.v4.u8     {temp1, temp2, temp3, temp4}, [in_addr];
    mov.v4.u16          foo, {temp2, temp3, temp4, temp1};
    mov.v4.u16          {temp3, temp4, temp1, temp2}, foo;
    mov.v4.u16          {temp4, temp1, temp2, temp3}, {temp3, temp4, temp1, temp2};
    st.global.v4.u8     [out_addr], {temp1, temp2, temp3, temp4};
    ret;
}