.version 6.5
.target sm_30
.address_size 64

.visible .entry mul_wide(
	.param .u64 input,
	.param .u64 output
)
{
	.reg .u64 	    in_addr;
    .reg .u64 	    out_addr;
    .reg .s32 	    inp1;
    .reg .s32 	    inp2;
    .reg .s64 	    result;

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

    ld.global.s32          inp1, [in_addr];
    ld.global.s32          inp2, [in_addr+4];
	mul.wide.s32    result, inp1, inp2;
    st.u64          [out_addr], result;
	ret;
}