blob: 2d6f8a52ba0ceac1c8d0a07c2c2d031a6a801042 (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
|
.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;
}
|