.version 6.5
.target sm_30
.address_size 64

.visible .entry mad_s32(
	.param .u64 input,
	.param .u64 output
)
{
	.reg .u64 	    in_addr;
    .reg .u64 	    out_addr;
    .reg .s32 	    dst;
    .reg .s32 	    src1;
    .reg .s32 	    src2;
    .reg .s32 	    src3;

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

    ld.s32          src1, [in_addr];
    ld.s32          src2, [in_addr+4];
    ld.s32          src3, [in_addr+8];
	mad.lo.s32      dst, src1, src2, src3;
    st.s32          [out_addr], dst;
    st.s32          [out_addr+4], dst;
    st.s32          [out_addr+8], dst;
	ret;
}