.version 6.5
.target sm_30
.address_size 64

.const .align 8 .b16 constparams[4] = { 10, 20, 30, 40 };

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

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

    ld.const.b16    temp1, [constparams];
    ld.const.b16    temp2, [constparams+2];
    ld.const.b16    temp3, [constparams+4];
    ld.const.b16    temp4, [constparams+6];
    st.u16          [out_addr], temp1;
    st.u16          [out_addr+2], temp2;
    st.u16          [out_addr+4], temp3;
    st.u16          [out_addr+6], temp4;
	ret;
}