.version 6.5
.target sm_30
.address_size 64

.visible .entry vshr(
	.param .u64 input,
	.param .u64 output
)
{
	.reg .u64 	    in_addr;
    .reg .u64 	    out_addr;
    .reg .u32 	    temp1;
    .reg .u32 	    temp2;
    .reg .u32 	    temp3;
    .reg .u32 	    temp4;

    ld.param.u64 	            in_addr, [input];
    ld.param.u64 	            out_addr, [output];
    ld.b32                      temp2, [in_addr];
    ld.b32                      temp3, [in_addr+4];
    ld.b32                      temp4, [in_addr+8];

    vshr.u32.u32.u32.clamp.add  temp1, temp2, temp3, temp4;

    st.u32                      [out_addr], temp1;
	ret;
}