Skip to content

Commit d13098b

Browse files
AlexMacleanAlexMaclean
committed
pre-commit tests (use update_llc_test_checks)
1 parent 16e712e commit d13098b

File tree

3 files changed

+701
-364
lines changed

3 files changed

+701
-364
lines changed

‎llvm/test/CodeGen/NVPTX/convert-int-sm20.ll‎

Lines changed: 49 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -1,66 +1,91 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
12
; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s
23
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
34
; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 | %ptxas-verify %}
45
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
56

6-
77
;; Integer conversions happen inplicitly by loading/storing the proper types
88

9-
109
; i16
1110

1211
define i16 @cvt_i16_i32(i32 %x) {
13-
; CHECK: ld.param.b16 %r[[R0:[0-9]+]], [cvt_i16_i32_param_{{[0-9]+}}]
14-
; CHECK: st.param.b32 [func_retval{{[0-9]+}}], %r[[R0]]
15-
; CHECK: ret
12+
; CHECK-LABEL: cvt_i16_i32(
13+
; CHECK: {
14+
; CHECK-NEXT: .reg .b32 %r<2>;
15+
; CHECK-EMPTY:
16+
; CHECK-NEXT: // %bb.0:
17+
; CHECK-NEXT: ld.param.b16 %r1, [cvt_i16_i32_param_0];
18+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
19+
; CHECK-NEXT: ret;
1620
%a = trunc i32 %x to i16
1721
ret i16 %a
1822
}
1923

2024
define i16 @cvt_i16_i64(i64 %x) {
21-
; CHECK: ld.param.b16 %r[[R0:[0-9]+]], [cvt_i16_i64_param_{{[0-9]+}}]
22-
; CHECK: st.param.b32 [func_retval{{[0-9]+}}], %r[[R0]]
23-
; CHECK: ret
25+
; CHECK-LABEL: cvt_i16_i64(
26+
; CHECK: {
27+
; CHECK-NEXT: .reg .b32 %r<2>;
28+
; CHECK-EMPTY:
29+
; CHECK-NEXT: // %bb.0:
30+
; CHECK-NEXT: ld.param.b16 %r1, [cvt_i16_i64_param_0];
31+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
32+
; CHECK-NEXT: ret;
2433
%a = trunc i64 %x to i16
2534
ret i16 %a
2635
}
2736

28-
29-
3037
; i32
3138

3239
define i32 @cvt_i32_i16(i16 %x) {
33-
; CHECK: ld.param.b16 %r[[R0:[0-9]+]], [cvt_i32_i16_param_{{[0-9]+}}]
34-
; CHECK: st.param.b32 [func_retval{{[0-9]+}}], %r[[R0]]
35-
; CHECK: ret
40+
; CHECK-LABEL: cvt_i32_i16(
41+
; CHECK: {
42+
; CHECK-NEXT: .reg .b32 %r<2>;
43+
; CHECK-EMPTY:
44+
; CHECK-NEXT: // %bb.0:
45+
; CHECK-NEXT: ld.param.b16 %r1, [cvt_i32_i16_param_0];
46+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
47+
; CHECK-NEXT: ret;
3648
%a = zext i16 %x to i32
3749
ret i32 %a
3850
}
3951

4052
define i32 @cvt_i32_i64(i64 %x) {
41-
; CHECK: ld.param.b32 %r[[R0:[0-9]+]], [cvt_i32_i64_param_{{[0-9]+}}]
42-
; CHECK: st.param.b32 [func_retval{{[0-9]+}}], %r[[R0]]
43-
; CHECK: ret
53+
; CHECK-LABEL: cvt_i32_i64(
54+
; CHECK: {
55+
; CHECK-NEXT: .reg .b32 %r<2>;
56+
; CHECK-EMPTY:
57+
; CHECK-NEXT: // %bb.0:
58+
; CHECK-NEXT: ld.param.b32 %r1, [cvt_i32_i64_param_0];
59+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
60+
; CHECK-NEXT: ret;
4461
%a = trunc i64 %x to i32
4562
ret i32 %a
4663
}
4764

48-
49-
5065
; i64
5166

5267
define i64 @cvt_i64_i16(i16 %x) {
53-
; CHECK: ld.param.b16 %rd[[R0:[0-9]+]], [cvt_i64_i16_param_{{[0-9]+}}]
54-
; CHECK: st.param.b64 [func_retval{{[0-9]+}}], %rd[[R0]]
55-
; CHECK: ret
68+
; CHECK-LABEL: cvt_i64_i16(
69+
; CHECK: {
70+
; CHECK-NEXT: .reg .b64 %rd<2>;
71+
; CHECK-EMPTY:
72+
; CHECK-NEXT: // %bb.0:
73+
; CHECK-NEXT: ld.param.b16 %rd1, [cvt_i64_i16_param_0];
74+
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
75+
; CHECK-NEXT: ret;
5676
%a = zext i16 %x to i64
5777
ret i64 %a
5878
}
5979

6080
define i64 @cvt_i64_i32(i32 %x) {
61-
; CHECK: ld.param.b32 %rd[[R0:[0-9]+]], [cvt_i64_i32_param_{{[0-9]+}}]
62-
; CHECK: st.param.b64 [func_retval{{[0-9]+}}], %rd[[R0]]
63-
; CHECK: ret
81+
; CHECK-LABEL: cvt_i64_i32(
82+
; CHECK: {
83+
; CHECK-NEXT: .reg .b64 %rd<2>;
84+
; CHECK-EMPTY:
85+
; CHECK-NEXT: // %bb.0:
86+
; CHECK-NEXT: ld.param.b32 %rd1, [cvt_i64_i32_param_0];
87+
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
88+
; CHECK-NEXT: ret;
6489
%a = zext i32 %x to i64
6590
ret i64 %a
6691
}

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

Lines changed: 114 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,4 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
12
; Check that various LLVM idioms get lowered to NVPTX as expected.
23

34
; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s
@@ -8,114 +9,196 @@
89
%struct.S16 = type { i16, i16 }
910
%struct.S32 = type { i32, i32 }
1011

11-
; CHECK-LABEL: abs_i16(
1212
define i16 @abs_i16(i16 %a) {
13-
; CHECK: abs.s16
13+
; CHECK-LABEL: abs_i16(
14+
; CHECK: {
15+
; CHECK-NEXT: .reg .b16 %rs<3>;
16+
; CHECK-NEXT: .reg .b32 %r<2>;
17+
; CHECK-EMPTY:
18+
; CHECK-NEXT: // %bb.0:
19+
; CHECK-NEXT: ld.param.b16 %rs1, [abs_i16_param_0];
20+
; CHECK-NEXT: abs.s16 %rs2, %rs1;
21+
; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
22+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
23+
; CHECK-NEXT: ret;
1424
%neg = sub i16 0, %a
1525
%abs.cond = icmp sge i16 %a, 0
1626
%abs = select i1 %abs.cond, i16 %a, i16 %neg
1727
ret i16 %abs
1828
}
1929

20-
; CHECK-LABEL: abs_i32(
2130
define i32 @abs_i32(i32 %a) {
22-
; CHECK: abs.s32
31+
; CHECK-LABEL: abs_i32(
32+
; CHECK: {
33+
; CHECK-NEXT: .reg .b32 %r<3>;
34+
; CHECK-EMPTY:
35+
; CHECK-NEXT: // %bb.0:
36+
; CHECK-NEXT: ld.param.b32 %r1, [abs_i32_param_0];
37+
; CHECK-NEXT: abs.s32 %r2, %r1;
38+
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
39+
; CHECK-NEXT: ret;
2340
%neg = sub i32 0, %a
2441
%abs.cond = icmp sge i32 %a, 0
2542
%abs = select i1 %abs.cond, i32 %a, i32 %neg
2643
ret i32 %abs
2744
}
2845

29-
; CHECK-LABEL: abs_i64(
3046
define i64 @abs_i64(i64 %a) {
31-
; CHECK: abs.s64
47+
; CHECK-LABEL: abs_i64(
48+
; CHECK: {
49+
; CHECK-NEXT: .reg .b64 %rd<3>;
50+
; CHECK-EMPTY:
51+
; CHECK-NEXT: // %bb.0:
52+
; CHECK-NEXT: ld.param.b64 %rd1, [abs_i64_param_0];
53+
; CHECK-NEXT: abs.s64 %rd2, %rd1;
54+
; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
55+
; CHECK-NEXT: ret;
3256
%neg = sub i64 0, %a
3357
%abs.cond = icmp sge i64 %a, 0
3458
%abs = select i1 %abs.cond, i64 %a, i64 %neg
3559
ret i64 %abs
3660
}
3761

38-
; CHECK-LABEL: i32_to_2xi16(
3962
define %struct.S16 @i32_to_2xi16(i32 noundef %in) {
63+
; CHECK-LABEL: i32_to_2xi16(
64+
; CHECK: {
65+
; CHECK-NEXT: .reg .b16 %rs<3>;
66+
; CHECK-NEXT: .reg .b32 %r<2>;
67+
; CHECK-EMPTY:
68+
; CHECK-NEXT: // %bb.0:
69+
; CHECK-NEXT: ld.param.b32 %r1, [i32_to_2xi16_param_0];
70+
; CHECK-NEXT: cvt.u16.u32 %rs1, %r1;
71+
; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs2}, %r1; }
72+
; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
73+
; CHECK-NEXT: st.param.b16 [func_retval0+2], %rs2;
74+
; CHECK-NEXT: ret;
4075
%low = trunc i32 %in to i16
4176
%high32 = lshr i32 %in, 16
4277
%high = trunc i32 %high32 to i16
43-
; CHECK: ld.param.b32 %[[R32:r[0-9]+]], [i32_to_2xi16_param_0];
44-
; CHECK-DAG: cvt.u16.u32 %rs{{[0-9+]}}, %[[R32]];
45-
; CHECK-DAG mov.b32 {tmp, %rs{{[0-9+]}}}, %[[R32]];
4678
%s1 = insertvalue %struct.S16 poison, i16 %low, 0
4779
%s = insertvalue %struct.S16 %s1, i16 %high, 1
4880
ret %struct.S16 %s
4981
}
5082

51-
; CHECK-LABEL: i32_to_2xi16_lh(
5283
; Same as above, but with rearranged order of low/high parts.
5384
define %struct.S16 @i32_to_2xi16_lh(i32 noundef %in) {
85+
; CHECK-LABEL: i32_to_2xi16_lh(
86+
; CHECK: {
87+
; CHECK-NEXT: .reg .b16 %rs<3>;
88+
; CHECK-NEXT: .reg .b32 %r<2>;
89+
; CHECK-EMPTY:
90+
; CHECK-NEXT: // %bb.0:
91+
; CHECK-NEXT: ld.param.b32 %r1, [i32_to_2xi16_lh_param_0];
92+
; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r1; }
93+
; CHECK-NEXT: cvt.u16.u32 %rs2, %r1;
94+
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
95+
; CHECK-NEXT: st.param.b16 [func_retval0+2], %rs1;
96+
; CHECK-NEXT: ret;
5497
%high32 = lshr i32 %in, 16
5598
%high = trunc i32 %high32 to i16
5699
%low = trunc i32 %in to i16
57-
; CHECK: ld.param.b32 %[[R32:r[0-9]+]], [i32_to_2xi16_lh_param_0];
58-
; CHECK-DAG: cvt.u16.u32 %rs{{[0-9+]}}, %[[R32]];
59-
; CHECK-DAG mov.b32 {tmp, %rs{{[0-9+]}}}, %[[R32]];
60100
%s1 = insertvalue %struct.S16 poison, i16 %low, 0
61101
%s = insertvalue %struct.S16 %s1, i16 %high, 1
62102
ret %struct.S16 %s
63103
}
64104

65105

66-
; CHECK-LABEL: i32_to_2xi16_not(
67106
define %struct.S16 @i32_to_2xi16_not(i32 noundef %in) {
107+
; CHECK-LABEL: i32_to_2xi16_not(
108+
; CHECK: {
109+
; CHECK-NEXT: .reg .b16 %rs<3>;
110+
; CHECK-NEXT: .reg .b32 %r<3>;
111+
; CHECK-EMPTY:
112+
; CHECK-NEXT: // %bb.0:
113+
; CHECK-NEXT: ld.param.b32 %r1, [i32_to_2xi16_not_param_0];
114+
; CHECK-NEXT: cvt.u16.u32 %rs1, %r1;
115+
; CHECK-NEXT: shr.u32 %r2, %r1, 15;
116+
; CHECK-NEXT: cvt.u16.u32 %rs2, %r2;
117+
; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
118+
; CHECK-NEXT: st.param.b16 [func_retval0+2], %rs2;
119+
; CHECK-NEXT: ret;
68120
%low = trunc i32 %in to i16
69121
; Shift by any value other than 16 blocks the conversiopn to mov.
70122
%high32 = lshr i32 %in, 15
71123
%high = trunc i32 %high32 to i16
72-
; CHECK: cvt.u16.u32
73-
; CHECK: shr.u32
74-
; CHECK: cvt.u16.u32
75124
%s1 = insertvalue %struct.S16 poison, i16 %low, 0
76125
%s = insertvalue %struct.S16 %s1, i16 %high, 1
77126
ret %struct.S16 %s
78127
}
79128

80-
; CHECK-LABEL: i64_to_2xi32(
81129
define %struct.S32 @i64_to_2xi32(i64 noundef %in) {
130+
; CHECK-LABEL: i64_to_2xi32(
131+
; CHECK: {
132+
; CHECK-NEXT: .reg .b32 %r<3>;
133+
; CHECK-NEXT: .reg .b64 %rd<2>;
134+
; CHECK-EMPTY:
135+
; CHECK-NEXT: // %bb.0:
136+
; CHECK-NEXT: ld.param.b64 %rd1, [i64_to_2xi32_param_0];
137+
; CHECK-NEXT: cvt.u32.u64 %r1, %rd1;
138+
; CHECK-NEXT: { .reg .b32 tmp; mov.b64 {tmp, %r2}, %rd1; }
139+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
140+
; CHECK-NEXT: st.param.b32 [func_retval0+4], %r2;
141+
; CHECK-NEXT: ret;
82142
%low = trunc i64 %in to i32
83143
%high64 = lshr i64 %in, 32
84144
%high = trunc i64 %high64 to i32
85-
; CHECK: ld.param.b64 %[[R64:rd[0-9]+]], [i64_to_2xi32_param_0];
86-
; CHECK-DAG: cvt.u32.u64 %r{{[0-9+]}}, %[[R64]];
87-
; CHECK-DAG mov.b64 {tmp, %r{{[0-9+]}}}, %[[R64]];
88145
%s1 = insertvalue %struct.S32 poison, i32 %low, 0
89146
%s = insertvalue %struct.S32 %s1, i32 %high, 1
90147
ret %struct.S32 %s
91148
}
92149

93-
; CHECK-LABEL: i64_to_2xi32_not(
94150
define %struct.S32 @i64_to_2xi32_not(i64 noundef %in) {
151+
; CHECK-LABEL: i64_to_2xi32_not(
152+
; CHECK: {
153+
; CHECK-NEXT: .reg .b32 %r<3>;
154+
; CHECK-NEXT: .reg .b64 %rd<3>;
155+
; CHECK-EMPTY:
156+
; CHECK-NEXT: // %bb.0:
157+
; CHECK-NEXT: ld.param.b64 %rd1, [i64_to_2xi32_not_param_0];
158+
; CHECK-NEXT: cvt.u32.u64 %r1, %rd1;
159+
; CHECK-NEXT: shr.u64 %rd2, %rd1, 31;
160+
; CHECK-NEXT: cvt.u32.u64 %r2, %rd2;
161+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
162+
; CHECK-NEXT: st.param.b32 [func_retval0+4], %r2;
163+
; CHECK-NEXT: ret;
95164
%low = trunc i64 %in to i32
96165
; Shift by any value other than 32 blocks the conversiopn to mov.
97166
%high64 = lshr i64 %in, 31
98167
%high = trunc i64 %high64 to i32
99-
; CHECK: cvt.u32.u64
100-
; CHECK: shr.u64
101-
; CHECK: cvt.u32.u64
102168
%s1 = insertvalue %struct.S32 poison, i32 %low, 0
103169
%s = insertvalue %struct.S32 %s1, i32 %high, 1
104170
ret %struct.S32 %s
105171
}
106172

107-
; CHECK-LABEL: i32_to_2xi16_shr(
108173
; Make sure we do not get confused when our input itself is [al]shr.
109174
define %struct.S16 @i32_to_2xi16_shr(i32 noundef %i){
175+
; CHECK-LABEL: i32_to_2xi16_shr(
176+
; CHECK: {
177+
; CHECK-NEXT: .reg .b16 %rs<3>;
178+
; CHECK-NEXT: .reg .b32 %r<3>;
179+
; CHECK-EMPTY:
180+
; CHECK-NEXT: // %bb.0:
181+
; CHECK-NEXT: ld.param.b32 %r1, [i32_to_2xi16_shr_param_0];
182+
; CHECK-NEXT: { // callseq 0, 0
183+
; CHECK-NEXT: .param .b32 param0;
184+
; CHECK-NEXT: st.param.b32 [param0], %r1;
185+
; CHECK-NEXT: call.uni
186+
; CHECK-NEXT: escape_int,
187+
; CHECK-NEXT: (
188+
; CHECK-NEXT: param0
189+
; CHECK-NEXT: );
190+
; CHECK-NEXT: } // callseq 0
191+
; CHECK-NEXT: shr.s32 %r2, %r1, 16;
192+
; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r1; }
193+
; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs2}, %r2; }
194+
; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
195+
; CHECK-NEXT: st.param.b16 [func_retval0+2], %rs2;
196+
; CHECK-NEXT: ret;
110197
call void @escape_int(i32 %i); // Force %i to be loaded completely.
111198
%i1 = ashr i32 %i, 16
112199
%l = trunc i32 %i1 to i16
113200
%h32 = ashr i32 %i1, 16
114201
%h = trunc i32 %h32 to i16
115-
; CHECK: ld.param.b32 %[[R32:r[0-9]+]], [i32_to_2xi16_shr_param_0];
116-
; CHECK: shr.s32 %[[R32H:r[0-9]+]], %[[R32]], 16;
117-
; CHECK-DAG mov.b32 {tmp, %rs{{[0-9+]}}}, %[[R32]];
118-
; CHECK-DAG mov.b32 {tmp, %rs{{[0-9+]}}}, %[[R32H]];
119202
%s0 = insertvalue %struct.S16 poison, i16 %l, 0
120203
%s1 = insertvalue %struct.S16 %s0, i16 %h, 1
121204
ret %struct.S16 %s1

0 commit comments

Comments
 (0)