Skip to content

Commit 49dc163

Browse files
RKSimonRKSimon
committed
[NVPTX] mulwide.ll - regenerate test checks
1 parent c438603 commit 49dc163

File tree

1 file changed

+254
-40
lines changed

1 file changed

+254
-40
lines changed

‎llvm/test/CodeGen/NVPTX/mulwide.ll

Lines changed: 254 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -1,111 +1,325 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
12
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O3 | FileCheck %s --check-prefix=OPT
23
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | FileCheck %s --check-prefix=NOOPT
34
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O3 | %ptxas-verify %}
45
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | %ptxas-verify %}
56

6-
; OPT-LABEL: @mulwide16
7-
; NOOPT-LABEL: @mulwide16
87
define i32 @mulwide16(i16 %a, i16 %b) {
9-
; OPT: mul.wide.s16
10-
; NOOPT: mul.lo.s32
8+
; OPT-LABEL: mulwide16(
9+
; OPT: {
10+
; OPT-NEXT: .reg .b16 %rs<3>;
11+
; OPT-NEXT: .reg .b32 %r<2>;
12+
; OPT-EMPTY:
13+
; OPT-NEXT: // %bb.0:
14+
; OPT-NEXT: ld.param.b16 %rs1, [mulwide16_param_0];
15+
; OPT-NEXT: ld.param.b16 %rs2, [mulwide16_param_1];
16+
; OPT-NEXT: mul.wide.s16 %r1, %rs1, %rs2;
17+
; OPT-NEXT: st.param.b32 [func_retval0], %r1;
18+
; OPT-NEXT: ret;
19+
;
20+
; NOOPT-LABEL: mulwide16(
21+
; NOOPT: {
22+
; NOOPT-NEXT: .reg .b16 %rs<3>;
23+
; NOOPT-NEXT: .reg .b32 %r<4>;
24+
; NOOPT-EMPTY:
25+
; NOOPT-NEXT: // %bb.0:
26+
; NOOPT-NEXT: ld.param.b16 %rs2, [mulwide16_param_1];
27+
; NOOPT-NEXT: ld.param.b16 %rs1, [mulwide16_param_0];
28+
; NOOPT-NEXT: cvt.s32.s16 %r1, %rs1;
29+
; NOOPT-NEXT: cvt.s32.s16 %r2, %rs2;
30+
; NOOPT-NEXT: mul.lo.s32 %r3, %r1, %r2;
31+
; NOOPT-NEXT: st.param.b32 [func_retval0], %r3;
32+
; NOOPT-NEXT: ret;
1133
%val0 = sext i16 %a to i32
1234
%val1 = sext i16 %b to i32
1335
%val2 = mul i32 %val0, %val1
1436
ret i32 %val2
1537
}
1638

17-
; OPT-LABEL: @mulwideu16
18-
; NOOPT-LABEL: @mulwideu16
1939
define i32 @mulwideu16(i16 %a, i16 %b) {
20-
; OPT: mul.wide.u16
21-
; NOOPT: mul.lo.s32
40+
; OPT-LABEL: mulwideu16(
41+
; OPT: {
42+
; OPT-NEXT: .reg .b16 %rs<3>;
43+
; OPT-NEXT: .reg .b32 %r<2>;
44+
; OPT-EMPTY:
45+
; OPT-NEXT: // %bb.0:
46+
; OPT-NEXT: ld.param.b16 %rs1, [mulwideu16_param_0];
47+
; OPT-NEXT: ld.param.b16 %rs2, [mulwideu16_param_1];
48+
; OPT-NEXT: mul.wide.u16 %r1, %rs1, %rs2;
49+
; OPT-NEXT: st.param.b32 [func_retval0], %r1;
50+
; OPT-NEXT: ret;
51+
;
52+
; NOOPT-LABEL: mulwideu16(
53+
; NOOPT: {
54+
; NOOPT-NEXT: .reg .b16 %rs<3>;
55+
; NOOPT-NEXT: .reg .b32 %r<4>;
56+
; NOOPT-EMPTY:
57+
; NOOPT-NEXT: // %bb.0:
58+
; NOOPT-NEXT: ld.param.b16 %rs2, [mulwideu16_param_1];
59+
; NOOPT-NEXT: ld.param.b16 %rs1, [mulwideu16_param_0];
60+
; NOOPT-NEXT: cvt.u32.u16 %r1, %rs1;
61+
; NOOPT-NEXT: cvt.u32.u16 %r2, %rs2;
62+
; NOOPT-NEXT: mul.lo.s32 %r3, %r1, %r2;
63+
; NOOPT-NEXT: st.param.b32 [func_retval0], %r3;
64+
; NOOPT-NEXT: ret;
2265
%val0 = zext i16 %a to i32
2366
%val1 = zext i16 %b to i32
2467
%val2 = mul i32 %val0, %val1
2568
ret i32 %val2
2669
}
2770

28-
; OPT-LABEL: @mulwide8
29-
; NOOPT-LABEL: @mulwide8
3071
define i32 @mulwide8(i8 %a, i8 %b) {
31-
; OPT: mul.wide.s16
32-
; NOOPT: mul.lo.s32
72+
; OPT-LABEL: mulwide8(
73+
; OPT: {
74+
; OPT-NEXT: .reg .b16 %rs<3>;
75+
; OPT-NEXT: .reg .b32 %r<2>;
76+
; OPT-EMPTY:
77+
; OPT-NEXT: // %bb.0:
78+
; OPT-NEXT: ld.param.s8 %rs1, [mulwide8_param_0];
79+
; OPT-NEXT: ld.param.s8 %rs2, [mulwide8_param_1];
80+
; OPT-NEXT: mul.wide.s16 %r1, %rs1, %rs2;
81+
; OPT-NEXT: st.param.b32 [func_retval0], %r1;
82+
; OPT-NEXT: ret;
83+
;
84+
; NOOPT-LABEL: mulwide8(
85+
; NOOPT: {
86+
; NOOPT-NEXT: .reg .b16 %rs<3>;
87+
; NOOPT-NEXT: .reg .b32 %r<6>;
88+
; NOOPT-EMPTY:
89+
; NOOPT-NEXT: // %bb.0:
90+
; NOOPT-NEXT: ld.param.b8 %rs2, [mulwide8_param_1];
91+
; NOOPT-NEXT: ld.param.b8 %rs1, [mulwide8_param_0];
92+
; NOOPT-NEXT: cvt.u32.u16 %r1, %rs1;
93+
; NOOPT-NEXT: cvt.s32.s8 %r2, %r1;
94+
; NOOPT-NEXT: cvt.u32.u16 %r3, %rs2;
95+
; NOOPT-NEXT: cvt.s32.s8 %r4, %r3;
96+
; NOOPT-NEXT: mul.lo.s32 %r5, %r2, %r4;
97+
; NOOPT-NEXT: st.param.b32 [func_retval0], %r5;
98+
; NOOPT-NEXT: ret;
3399
%val0 = sext i8 %a to i32
34100
%val1 = sext i8 %b to i32
35101
%val2 = mul i32 %val0, %val1
36102
ret i32 %val2
37103
}
38104

39-
; OPT-LABEL: @mulwideu8
40-
; NOOPT-LABEL: @mulwideu8
41105
define i32 @mulwideu8(i8 %a, i8 %b) {
42-
; OPT: mul.wide.u16
43-
; NOOPT: mul.lo.s32
106+
; OPT-LABEL: mulwideu8(
107+
; OPT: {
108+
; OPT-NEXT: .reg .b16 %rs<3>;
109+
; OPT-NEXT: .reg .b32 %r<2>;
110+
; OPT-EMPTY:
111+
; OPT-NEXT: // %bb.0:
112+
; OPT-NEXT: ld.param.b8 %rs1, [mulwideu8_param_0];
113+
; OPT-NEXT: ld.param.b8 %rs2, [mulwideu8_param_1];
114+
; OPT-NEXT: mul.wide.u16 %r1, %rs1, %rs2;
115+
; OPT-NEXT: st.param.b32 [func_retval0], %r1;
116+
; OPT-NEXT: ret;
117+
;
118+
; NOOPT-LABEL: mulwideu8(
119+
; NOOPT: {
120+
; NOOPT-NEXT: .reg .b16 %rs<3>;
121+
; NOOPT-NEXT: .reg .b32 %r<6>;
122+
; NOOPT-EMPTY:
123+
; NOOPT-NEXT: // %bb.0:
124+
; NOOPT-NEXT: ld.param.b8 %rs2, [mulwideu8_param_1];
125+
; NOOPT-NEXT: ld.param.b8 %rs1, [mulwideu8_param_0];
126+
; NOOPT-NEXT: cvt.u32.u16 %r1, %rs1;
127+
; NOOPT-NEXT: and.b32 %r2, %r1, 255;
128+
; NOOPT-NEXT: cvt.u32.u16 %r3, %rs2;
129+
; NOOPT-NEXT: and.b32 %r4, %r3, 255;
130+
; NOOPT-NEXT: mul.lo.s32 %r5, %r2, %r4;
131+
; NOOPT-NEXT: st.param.b32 [func_retval0], %r5;
132+
; NOOPT-NEXT: ret;
44133
%val0 = zext i8 %a to i32
45134
%val1 = zext i8 %b to i32
46135
%val2 = mul i32 %val0, %val1
47136
ret i32 %val2
48137
}
49138

50-
; OPT-LABEL: @mulwide32
51-
; NOOPT-LABEL: @mulwide32
52139
define i64 @mulwide32(i32 %a, i32 %b) {
53-
; OPT: mul.wide.s32
54-
; NOOPT: mul.lo.s64
140+
; OPT-LABEL: mulwide32(
141+
; OPT: {
142+
; OPT-NEXT: .reg .b32 %r<3>;
143+
; OPT-NEXT: .reg .b64 %rd<2>;
144+
; OPT-EMPTY:
145+
; OPT-NEXT: // %bb.0:
146+
; OPT-NEXT: ld.param.b32 %r1, [mulwide32_param_0];
147+
; OPT-NEXT: ld.param.b32 %r2, [mulwide32_param_1];
148+
; OPT-NEXT: mul.wide.s32 %rd1, %r1, %r2;
149+
; OPT-NEXT: st.param.b64 [func_retval0], %rd1;
150+
; OPT-NEXT: ret;
151+
;
152+
; NOOPT-LABEL: mulwide32(
153+
; NOOPT: {
154+
; NOOPT-NEXT: .reg .b32 %r<3>;
155+
; NOOPT-NEXT: .reg .b64 %rd<4>;
156+
; NOOPT-EMPTY:
157+
; NOOPT-NEXT: // %bb.0:
158+
; NOOPT-NEXT: ld.param.b32 %r2, [mulwide32_param_1];
159+
; NOOPT-NEXT: ld.param.b32 %r1, [mulwide32_param_0];
160+
; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;
161+
; NOOPT-NEXT: cvt.s64.s32 %rd2, %r2;
162+
; NOOPT-NEXT: mul.lo.s64 %rd3, %rd1, %rd2;
163+
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd3;
164+
; NOOPT-NEXT: ret;
55165
%val0 = sext i32 %a to i64
56166
%val1 = sext i32 %b to i64
57167
%val2 = mul i64 %val0, %val1
58168
ret i64 %val2
59169
}
60170

61-
; OPT-LABEL: @mulwideu32
62-
; NOOPT-LABEL: @mulwideu32
63171
define i64 @mulwideu32(i32 %a, i32 %b) {
64-
; OPT: mul.wide.u32
65-
; NOOPT: mul.lo.s64
172+
; OPT-LABEL: mulwideu32(
173+
; OPT: {
174+
; OPT-NEXT: .reg .b32 %r<3>;
175+
; OPT-NEXT: .reg .b64 %rd<2>;
176+
; OPT-EMPTY:
177+
; OPT-NEXT: // %bb.0:
178+
; OPT-NEXT: ld.param.b32 %r1, [mulwideu32_param_0];
179+
; OPT-NEXT: ld.param.b32 %r2, [mulwideu32_param_1];
180+
; OPT-NEXT: mul.wide.u32 %rd1, %r1, %r2;
181+
; OPT-NEXT: st.param.b64 [func_retval0], %rd1;
182+
; OPT-NEXT: ret;
183+
;
184+
; NOOPT-LABEL: mulwideu32(
185+
; NOOPT: {
186+
; NOOPT-NEXT: .reg .b32 %r<3>;
187+
; NOOPT-NEXT: .reg .b64 %rd<4>;
188+
; NOOPT-EMPTY:
189+
; NOOPT-NEXT: // %bb.0:
190+
; NOOPT-NEXT: ld.param.b32 %r2, [mulwideu32_param_1];
191+
; NOOPT-NEXT: ld.param.b32 %r1, [mulwideu32_param_0];
192+
; NOOPT-NEXT: cvt.u64.u32 %rd1, %r1;
193+
; NOOPT-NEXT: cvt.u64.u32 %rd2, %r2;
194+
; NOOPT-NEXT: mul.lo.s64 %rd3, %rd1, %rd2;
195+
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd3;
196+
; NOOPT-NEXT: ret;
66197
%val0 = zext i32 %a to i64
67198
%val1 = zext i32 %b to i64
68199
%val2 = mul i64 %val0, %val1
69200
ret i64 %val2
70201
}
71202

72-
; OPT-LABEL: @mulwideu7
73-
; NOOPT-LABEL: @mulwideu7
74203
define i64 @mulwideu7(i7 %a, i7 %b) {
75-
; OPT: mul.wide.u32
76-
; NOOPT: mul.lo.s64
204+
; OPT-LABEL: mulwideu7(
205+
; OPT: {
206+
; OPT-NEXT: .reg .b32 %r<3>;
207+
; OPT-NEXT: .reg .b64 %rd<2>;
208+
; OPT-EMPTY:
209+
; OPT-NEXT: // %bb.0:
210+
; OPT-NEXT: ld.param.b8 %r1, [mulwideu7_param_0];
211+
; OPT-NEXT: ld.param.b8 %r2, [mulwideu7_param_1];
212+
; OPT-NEXT: mul.wide.u32 %rd1, %r1, %r2;
213+
; OPT-NEXT: st.param.b64 [func_retval0], %rd1;
214+
; OPT-NEXT: ret;
215+
;
216+
; NOOPT-LABEL: mulwideu7(
217+
; NOOPT: {
218+
; NOOPT-NEXT: .reg .b16 %rs<3>;
219+
; NOOPT-NEXT: .reg .b64 %rd<6>;
220+
; NOOPT-EMPTY:
221+
; NOOPT-NEXT: // %bb.0:
222+
; NOOPT-NEXT: ld.param.b8 %rs2, [mulwideu7_param_1];
223+
; NOOPT-NEXT: ld.param.b8 %rs1, [mulwideu7_param_0];
224+
; NOOPT-NEXT: cvt.u64.u16 %rd1, %rs1;
225+
; NOOPT-NEXT: and.b64 %rd2, %rd1, 127;
226+
; NOOPT-NEXT: cvt.u64.u16 %rd3, %rs2;
227+
; NOOPT-NEXT: and.b64 %rd4, %rd3, 127;
228+
; NOOPT-NEXT: mul.lo.s64 %rd5, %rd2, %rd4;
229+
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd5;
230+
; NOOPT-NEXT: ret;
77231
%val0 = zext i7 %a to i64
78232
%val1 = zext i7 %b to i64
79233
%val2 = mul i64 %val0, %val1
80234
ret i64 %val2
81235
}
82236

83-
; OPT-LABEL: @mulwides7
84-
; NOOPT-LABEL: @mulwides7
85237
define i64 @mulwides7(i7 %a, i7 %b) {
86-
; OPT: mul.wide.s32
87-
; NOOPT: mul.lo.s64
238+
; OPT-LABEL: mulwides7(
239+
; OPT: {
240+
; OPT-NEXT: .reg .b32 %r<5>;
241+
; OPT-NEXT: .reg .b64 %rd<2>;
242+
; OPT-EMPTY:
243+
; OPT-NEXT: // %bb.0:
244+
; OPT-NEXT: ld.param.b8 %r1, [mulwides7_param_0];
245+
; OPT-NEXT: bfe.s32 %r2, %r1, 0, 7;
246+
; OPT-NEXT: ld.param.b8 %r3, [mulwides7_param_1];
247+
; OPT-NEXT: bfe.s32 %r4, %r3, 0, 7;
248+
; OPT-NEXT: mul.wide.s32 %rd1, %r2, %r4;
249+
; OPT-NEXT: st.param.b64 [func_retval0], %rd1;
250+
; OPT-NEXT: ret;
251+
;
252+
; NOOPT-LABEL: mulwides7(
253+
; NOOPT: {
254+
; NOOPT-NEXT: .reg .b16 %rs<3>;
255+
; NOOPT-NEXT: .reg .b64 %rd<6>;
256+
; NOOPT-EMPTY:
257+
; NOOPT-NEXT: // %bb.0:
258+
; NOOPT-NEXT: ld.param.b8 %rs2, [mulwides7_param_1];
259+
; NOOPT-NEXT: ld.param.b8 %rs1, [mulwides7_param_0];
260+
; NOOPT-NEXT: cvt.u64.u16 %rd1, %rs1;
261+
; NOOPT-NEXT: bfe.s64 %rd2, %rd1, 0, 7;
262+
; NOOPT-NEXT: cvt.u64.u16 %rd3, %rs2;
263+
; NOOPT-NEXT: bfe.s64 %rd4, %rd3, 0, 7;
264+
; NOOPT-NEXT: mul.lo.s64 %rd5, %rd2, %rd4;
265+
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd5;
266+
; NOOPT-NEXT: ret;
88267
%val0 = sext i7 %a to i64
89268
%val1 = sext i7 %b to i64
90269
%val2 = mul i64 %val0, %val1
91270
ret i64 %val2
92271
}
93272

94-
; OPT-LABEL: @shl30
95-
; NOOPT-LABEL: @shl30
96273
define i64 @shl30(i32 %a) {
97-
; OPT: mul.wide
98-
; NOOPT: shl.b64
274+
; OPT-LABEL: shl30(
275+
; OPT: {
276+
; OPT-NEXT: .reg .b32 %r<2>;
277+
; OPT-NEXT: .reg .b64 %rd<2>;
278+
; OPT-EMPTY:
279+
; OPT-NEXT: // %bb.0:
280+
; OPT-NEXT: ld.param.b32 %r1, [shl30_param_0];
281+
; OPT-NEXT: mul.wide.s32 %rd1, %r1, 1073741824;
282+
; OPT-NEXT: st.param.b64 [func_retval0], %rd1;
283+
; OPT-NEXT: ret;
284+
;
285+
; NOOPT-LABEL: shl30(
286+
; NOOPT: {
287+
; NOOPT-NEXT: .reg .b32 %r<2>;
288+
; NOOPT-NEXT: .reg .b64 %rd<3>;
289+
; NOOPT-EMPTY:
290+
; NOOPT-NEXT: // %bb.0:
291+
; NOOPT-NEXT: ld.param.b32 %r1, [shl30_param_0];
292+
; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;
293+
; NOOPT-NEXT: shl.b64 %rd2, %rd1, 30;
294+
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd2;
295+
; NOOPT-NEXT: ret;
99296
%conv = sext i32 %a to i64
100297
%shl = shl i64 %conv, 30
101298
ret i64 %shl
102299
}
103300

104-
; OPT-LABEL: @shl31
105-
; NOOPT-LABEL: @shl31
106301
define i64 @shl31(i32 %a) {
107-
; OPT-NOT: mul.wide
108-
; NOOPT-NOT: mul.wide
302+
; OPT-LABEL: shl31(
303+
; OPT: {
304+
; OPT-NEXT: .reg .b64 %rd<3>;
305+
; OPT-EMPTY:
306+
; OPT-NEXT: // %bb.0:
307+
; OPT-NEXT: ld.param.s32 %rd1, [shl31_param_0];
308+
; OPT-NEXT: shl.b64 %rd2, %rd1, 31;
309+
; OPT-NEXT: st.param.b64 [func_retval0], %rd2;
310+
; OPT-NEXT: ret;
311+
;
312+
; NOOPT-LABEL: shl31(
313+
; NOOPT: {
314+
; NOOPT-NEXT: .reg .b32 %r<2>;
315+
; NOOPT-NEXT: .reg .b64 %rd<3>;
316+
; NOOPT-EMPTY:
317+
; NOOPT-NEXT: // %bb.0:
318+
; NOOPT-NEXT: ld.param.b32 %r1, [shl31_param_0];
319+
; NOOPT-NEXT: cvt.s64.s32 %rd1, %r1;
320+
; NOOPT-NEXT: shl.b64 %rd2, %rd1, 31;
321+
; NOOPT-NEXT: st.param.b64 [func_retval0], %rd2;
322+
; NOOPT-NEXT: ret;
109323
%conv = sext i32 %a to i64
110324
%shl = shl i64 %conv, 31
111325
ret i64 %shl

0 commit comments

Comments
 (0)