Skip to content

Commit ca2f50c

Browse files
AlexMacleanAlexMaclean
committed
update after rebase
1 parent e7c680c commit ca2f50c

15 files changed

+239
-387
lines changed

‎llvm/test/CodeGen/MIR/NVPTX/floating-point-immediate-operands.mir‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ body: |
5050
; CHECK: %6:b32 = FADD_rnf32ri %5, float 6.250000e+00
5151
%6 = FADD_rnf32ri %5, float 6.250000e+00
5252
%7 = FMUL_rnf32rr %6, %4
53-
ST_i32 %7, 0, 0, 101, 3, 32, &func_retval0, 0 :: (store (s32), addrspace 101)
53+
ST_i32 %7, 0, 0, 101, 32, &func_retval0, 0 :: (store (s32), addrspace 101)
5454
Return
5555
...
5656
---
@@ -76,6 +76,6 @@ body: |
7676
; CHECK: %6:b32 = FADD_rnf32ri %5, float 0x7FF8000000000000
7777
%6 = FADD_rnf32ri %5, float 0x7FF8000000000000
7878
%7 = FMUL_rnf32rr %6, %4
79-
ST_i32 %7, 0, 0, 101, 3, 32, &func_retval0, 0 :: (store (s32), addrspace 101)
79+
ST_i32 %7, 0, 0, 101, 32, &func_retval0, 0 :: (store (s32), addrspace 101)
8080
Return
8181
...

‎llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -9,11 +9,10 @@ target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
99
define <2 x bfloat> @test_ret_const() #0 {
1010
; CHECK-LABEL: test_ret_const(
1111
; CHECK: {
12-
; CHECK-NEXT: .reg .b32 %r<2>;
12+
; CHECK-EMPTY:
1313
; CHECK-EMPTY:
1414
; CHECK-NEXT: // %bb.0:
15-
; CHECK-NEXT: mov.b32 %r1, 1073758080;
16-
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
15+
; CHECK-NEXT: st.param.b32 [func_retval0], 1073758080;
1716
; CHECK-NEXT: ret;
1817
ret <2 x bfloat> <bfloat 1.0, bfloat 2.0>
1918
}

‎llvm/test/CodeGen/NVPTX/compute-ptx-value-vts.ll‎

Lines changed: 14 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -6,57 +6,49 @@ target triple = "nvptx-nvidia-cuda"
66
define <6 x half> @half6() {
77
; CHECK-LABEL: half6(
88
; CHECK: {
9-
; CHECK-NEXT: .reg .b32 %r<2>;
9+
; CHECK-EMPTY:
1010
; CHECK-EMPTY:
1111
; CHECK-NEXT: // %bb.0:
12-
; CHECK-NEXT: mov.b32 %r1, 0;
13-
; CHECK-NEXT: st.param.b32 [func_retval0+8], %r1;
14-
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r1};
12+
; CHECK-NEXT: st.param.b32 [func_retval0+8], 0;
13+
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {0, 0};
1514
; CHECK-NEXT: ret;
1615
ret <6 x half> zeroinitializer
1716
}
1817

1918
define <10 x half> @half10() {
2019
; CHECK-LABEL: half10(
2120
; CHECK: {
22-
; CHECK-NEXT: .reg .b32 %r<2>;
21+
; CHECK-EMPTY:
2322
; CHECK-EMPTY:
2423
; CHECK-NEXT: // %bb.0:
25-
; CHECK-NEXT: mov.b32 %r1, 0;
26-
; CHECK-NEXT: st.param.b32 [func_retval0+16], %r1;
27-
; CHECK-NEXT: st.param.v2.b32 [func_retval0+8], {%r1, %r1};
28-
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r1};
24+
; CHECK-NEXT: st.param.b32 [func_retval0+16], 0;
25+
; CHECK-NEXT: st.param.v2.b32 [func_retval0+8], {0, 0};
26+
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {0, 0};
2927
; CHECK-NEXT: ret;
3028
ret <10 x half> zeroinitializer
3129
}
3230

3331
define <12 x i8> @byte12() {
3432
; CHECK-LABEL: byte12(
3533
; CHECK: {
36-
; CHECK-NEXT: .reg .b32 %r<2>;
37-
; CHECK-NEXT: .reg .b64 %rd<2>;
34+
; CHECK-EMPTY:
3835
; CHECK-EMPTY:
3936
; CHECK-NEXT: // %bb.0:
40-
; CHECK-NEXT: mov.b32 %r1, 0;
41-
; CHECK-NEXT: st.param.b32 [func_retval0+8], %r1;
42-
; CHECK-NEXT: mov.b64 %rd1, 0;
43-
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
37+
; CHECK-NEXT: st.param.b32 [func_retval0+8], 0;
38+
; CHECK-NEXT: st.param.b64 [func_retval0], 0;
4439
; CHECK-NEXT: ret;
4540
ret <12 x i8> zeroinitializer
4641
}
4742

4843
define <20 x i8> @byte20() {
4944
; CHECK-LABEL: byte20(
5045
; CHECK: {
51-
; CHECK-NEXT: .reg .b32 %r<2>;
52-
; CHECK-NEXT: .reg .b64 %rd<2>;
46+
; CHECK-EMPTY:
5347
; CHECK-EMPTY:
5448
; CHECK-NEXT: // %bb.0:
55-
; CHECK-NEXT: mov.b32 %r1, 0;
56-
; CHECK-NEXT: st.param.b32 [func_retval0+16], %r1;
57-
; CHECK-NEXT: mov.b64 %rd1, 0;
58-
; CHECK-NEXT: st.param.b64 [func_retval0+8], %rd1;
59-
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
49+
; CHECK-NEXT: st.param.b32 [func_retval0+16], 0;
50+
; CHECK-NEXT: st.param.b64 [func_retval0+8], 0;
51+
; CHECK-NEXT: st.param.b64 [func_retval0], 0;
6052
; CHECK-NEXT: ret;
6153
ret <20 x i8> zeroinitializer
6254
}

‎llvm/test/CodeGen/NVPTX/f16-instructions.ll‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -43,8 +43,7 @@
4343
target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
4444

4545
; CHECK-LABEL: test_ret_const(
46-
; CHECK: mov.b16 [[R:%rs[0-9]+]], 0x3C00;
47-
; CHECK-NEXT: st.param.b16 [func_retval0], [[R]];
46+
; CHECK: st.param.b16 [func_retval0], 0x3C00;
4847
; CHECK-NEXT: ret;
4948
define half @test_ret_const() #0 {
5049
ret half 1.0

‎llvm/test/CodeGen/NVPTX/f16x2-instructions.ll‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -34,11 +34,10 @@ target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
3434
define <2 x half> @test_ret_const() #0 {
3535
; CHECK-LABEL: test_ret_const(
3636
; CHECK: {
37-
; CHECK-NEXT: .reg .b32 %r<2>;
37+
; CHECK-EMPTY:
3838
; CHECK-EMPTY:
3939
; CHECK-NEXT: // %bb.0:
40-
; CHECK-NEXT: mov.b32 %r1, 1073757184;
41-
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
40+
; CHECK-NEXT: st.param.b32 [func_retval0], 1073757184;
4241
; CHECK-NEXT: ret;
4342
ret <2 x half> <half 1.0, half 2.0>
4443
}

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

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -116,11 +116,10 @@ define ptx_device float @f32_iir(float %x) {
116116
define ptx_device float @f32_iii(float %x) {
117117
; CHECK-LABEL: f32_iii(
118118
; CHECK: {
119-
; CHECK-NEXT: .reg .b32 %r<2>;
119+
; CHECK-EMPTY:
120120
; CHECK-EMPTY:
121121
; CHECK-NEXT: // %bb.0:
122-
; CHECK-NEXT: mov.b32 %r1, 1092616192;
123-
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
122+
; CHECK-NEXT: st.param.b32 [func_retval0], 1092616192;
124123
; CHECK-NEXT: ret;
125124
%r = call float @llvm.fma.f32(float 2.0, float 3.0, float 4.0)
126125
ret float %r

‎llvm/test/CodeGen/NVPTX/i1-icmp.ll‎

Lines changed: 30 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@ define i32 @icmp_i1_eq(i32 %a, i32 %b) {
88
; CHECK-LABEL: icmp_i1_eq(
99
; CHECK: {
1010
; CHECK-NEXT: .reg .pred %p<4>;
11-
; CHECK-NEXT: .reg .b32 %r<5>;
11+
; CHECK-NEXT: .reg .b32 %r<3>;
1212
; CHECK-EMPTY:
1313
; CHECK-NEXT: // %bb.0:
1414
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_eq_param_0];
@@ -18,12 +18,10 @@ define i32 @icmp_i1_eq(i32 %a, i32 %b) {
1818
; CHECK-NEXT: xor.pred %p3, %p1, %p2;
1919
; CHECK-NEXT: @%p3 bra $L__BB0_2;
2020
; CHECK-NEXT: // %bb.1: // %bb1
21-
; CHECK-NEXT: mov.b32 %r4, 1;
22-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
21+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
2322
; CHECK-NEXT: ret;
2423
; CHECK-NEXT: $L__BB0_2: // %bb2
25-
; CHECK-NEXT: mov.b32 %r3, 127;
26-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
24+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
2725
; CHECK-NEXT: ret;
2826
%p1 = icmp sgt i32 %a, 1
2927
%p2 = icmp sgt i32 %b, 1
@@ -39,7 +37,7 @@ define i32 @icmp_i1_ne(i32 %a, i32 %b) {
3937
; CHECK-LABEL: icmp_i1_ne(
4038
; CHECK: {
4139
; CHECK-NEXT: .reg .pred %p<5>;
42-
; CHECK-NEXT: .reg .b32 %r<5>;
40+
; CHECK-NEXT: .reg .b32 %r<3>;
4341
; CHECK-EMPTY:
4442
; CHECK-NEXT: // %bb.0:
4543
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_ne_param_0];
@@ -50,12 +48,10 @@ define i32 @icmp_i1_ne(i32 %a, i32 %b) {
5048
; CHECK-NEXT: not.pred %p4, %p3;
5149
; CHECK-NEXT: @%p4 bra $L__BB1_2;
5250
; CHECK-NEXT: // %bb.1: // %bb1
53-
; CHECK-NEXT: mov.b32 %r4, 1;
54-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
51+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
5552
; CHECK-NEXT: ret;
5653
; CHECK-NEXT: $L__BB1_2: // %bb2
57-
; CHECK-NEXT: mov.b32 %r3, 127;
58-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
54+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
5955
; CHECK-NEXT: ret;
6056
%p1 = icmp sgt i32 %a, 1
6157
%p2 = icmp sgt i32 %b, 1
@@ -71,7 +67,7 @@ define i32 @icmp_i1_sgt(i32 %a, i32 %b) {
7167
; CHECK-LABEL: icmp_i1_sgt(
7268
; CHECK: {
7369
; CHECK-NEXT: .reg .pred %p<4>;
74-
; CHECK-NEXT: .reg .b32 %r<5>;
70+
; CHECK-NEXT: .reg .b32 %r<3>;
7571
; CHECK-EMPTY:
7672
; CHECK-NEXT: // %bb.0:
7773
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_sgt_param_0];
@@ -81,12 +77,10 @@ define i32 @icmp_i1_sgt(i32 %a, i32 %b) {
8177
; CHECK-NEXT: or.pred %p3, %p1, %p2;
8278
; CHECK-NEXT: @%p3 bra $L__BB2_2;
8379
; CHECK-NEXT: // %bb.1: // %bb1
84-
; CHECK-NEXT: mov.b32 %r4, 1;
85-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
80+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
8681
; CHECK-NEXT: ret;
8782
; CHECK-NEXT: $L__BB2_2: // %bb2
88-
; CHECK-NEXT: mov.b32 %r3, 127;
89-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
83+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
9084
; CHECK-NEXT: ret;
9185
%p1 = icmp sgt i32 %a, 1
9286
%p2 = icmp sgt i32 %b, 1
@@ -102,7 +96,7 @@ define i32 @icmp_i1_slt(i32 %a, i32 %b) {
10296
; CHECK-LABEL: icmp_i1_slt(
10397
; CHECK: {
10498
; CHECK-NEXT: .reg .pred %p<4>;
105-
; CHECK-NEXT: .reg .b32 %r<5>;
99+
; CHECK-NEXT: .reg .b32 %r<3>;
106100
; CHECK-EMPTY:
107101
; CHECK-NEXT: // %bb.0:
108102
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_slt_param_0];
@@ -112,12 +106,10 @@ define i32 @icmp_i1_slt(i32 %a, i32 %b) {
112106
; CHECK-NEXT: or.pred %p3, %p2, %p1;
113107
; CHECK-NEXT: @%p3 bra $L__BB3_2;
114108
; CHECK-NEXT: // %bb.1: // %bb1
115-
; CHECK-NEXT: mov.b32 %r4, 1;
116-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
109+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
117110
; CHECK-NEXT: ret;
118111
; CHECK-NEXT: $L__BB3_2: // %bb2
119-
; CHECK-NEXT: mov.b32 %r3, 127;
120-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
112+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
121113
; CHECK-NEXT: ret;
122114
%p1 = icmp sgt i32 %a, 1
123115
%p2 = icmp sgt i32 %b, 1
@@ -133,7 +125,7 @@ define i32 @icmp_i1_sge(i32 %a, i32 %b) {
133125
; CHECK-LABEL: icmp_i1_sge(
134126
; CHECK: {
135127
; CHECK-NEXT: .reg .pred %p<4>;
136-
; CHECK-NEXT: .reg .b32 %r<5>;
128+
; CHECK-NEXT: .reg .b32 %r<3>;
137129
; CHECK-EMPTY:
138130
; CHECK-NEXT: // %bb.0:
139131
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_sge_param_0];
@@ -143,12 +135,10 @@ define i32 @icmp_i1_sge(i32 %a, i32 %b) {
143135
; CHECK-NEXT: and.pred %p3, %p1, %p2;
144136
; CHECK-NEXT: @%p3 bra $L__BB4_2;
145137
; CHECK-NEXT: // %bb.1: // %bb1
146-
; CHECK-NEXT: mov.b32 %r4, 1;
147-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
138+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
148139
; CHECK-NEXT: ret;
149140
; CHECK-NEXT: $L__BB4_2: // %bb2
150-
; CHECK-NEXT: mov.b32 %r3, 127;
151-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
141+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
152142
; CHECK-NEXT: ret;
153143
%p1 = icmp sgt i32 %a, 1
154144
%p2 = icmp sgt i32 %b, 1
@@ -164,7 +154,7 @@ define i32 @icmp_i1_sle(i32 %a, i32 %b) {
164154
; CHECK-LABEL: icmp_i1_sle(
165155
; CHECK: {
166156
; CHECK-NEXT: .reg .pred %p<4>;
167-
; CHECK-NEXT: .reg .b32 %r<5>;
157+
; CHECK-NEXT: .reg .b32 %r<3>;
168158
; CHECK-EMPTY:
169159
; CHECK-NEXT: // %bb.0:
170160
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_sle_param_0];
@@ -174,12 +164,10 @@ define i32 @icmp_i1_sle(i32 %a, i32 %b) {
174164
; CHECK-NEXT: and.pred %p3, %p2, %p1;
175165
; CHECK-NEXT: @%p3 bra $L__BB5_2;
176166
; CHECK-NEXT: // %bb.1: // %bb1
177-
; CHECK-NEXT: mov.b32 %r4, 1;
178-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
167+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
179168
; CHECK-NEXT: ret;
180169
; CHECK-NEXT: $L__BB5_2: // %bb2
181-
; CHECK-NEXT: mov.b32 %r3, 127;
182-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
170+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
183171
; CHECK-NEXT: ret;
184172
%p1 = icmp sgt i32 %a, 1
185173
%p2 = icmp sgt i32 %b, 1
@@ -195,7 +183,7 @@ define i32 @icmp_i1_uge(i32 %a, i32 %b) {
195183
; CHECK-LABEL: icmp_i1_uge(
196184
; CHECK: {
197185
; CHECK-NEXT: .reg .pred %p<4>;
198-
; CHECK-NEXT: .reg .b32 %r<5>;
186+
; CHECK-NEXT: .reg .b32 %r<3>;
199187
; CHECK-EMPTY:
200188
; CHECK-NEXT: // %bb.0:
201189
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_uge_param_0];
@@ -205,12 +193,10 @@ define i32 @icmp_i1_uge(i32 %a, i32 %b) {
205193
; CHECK-NEXT: and.pred %p3, %p2, %p1;
206194
; CHECK-NEXT: @%p3 bra $L__BB6_2;
207195
; CHECK-NEXT: // %bb.1: // %bb1
208-
; CHECK-NEXT: mov.b32 %r4, 1;
209-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
196+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
210197
; CHECK-NEXT: ret;
211198
; CHECK-NEXT: $L__BB6_2: // %bb2
212-
; CHECK-NEXT: mov.b32 %r3, 127;
213-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
199+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
214200
; CHECK-NEXT: ret;
215201
%p1 = icmp sgt i32 %a, 1
216202
%p2 = icmp sgt i32 %b, 1
@@ -226,7 +212,7 @@ define i32 @icmp_i1_ugt(i32 %a, i32 %b) {
226212
; CHECK-LABEL: icmp_i1_ugt(
227213
; CHECK: {
228214
; CHECK-NEXT: .reg .pred %p<4>;
229-
; CHECK-NEXT: .reg .b32 %r<5>;
215+
; CHECK-NEXT: .reg .b32 %r<3>;
230216
; CHECK-EMPTY:
231217
; CHECK-NEXT: // %bb.0:
232218
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_ugt_param_0];
@@ -236,12 +222,10 @@ define i32 @icmp_i1_ugt(i32 %a, i32 %b) {
236222
; CHECK-NEXT: or.pred %p3, %p2, %p1;
237223
; CHECK-NEXT: @%p3 bra $L__BB7_2;
238224
; CHECK-NEXT: // %bb.1: // %bb1
239-
; CHECK-NEXT: mov.b32 %r4, 1;
240-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
225+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
241226
; CHECK-NEXT: ret;
242227
; CHECK-NEXT: $L__BB7_2: // %bb2
243-
; CHECK-NEXT: mov.b32 %r3, 127;
244-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
228+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
245229
; CHECK-NEXT: ret;
246230
%p1 = icmp sgt i32 %a, 1
247231
%p2 = icmp sgt i32 %b, 1
@@ -257,7 +241,7 @@ define i32 @icmp_i1_ule(i32 %a, i32 %b) {
257241
; CHECK-LABEL: icmp_i1_ule(
258242
; CHECK: {
259243
; CHECK-NEXT: .reg .pred %p<4>;
260-
; CHECK-NEXT: .reg .b32 %r<5>;
244+
; CHECK-NEXT: .reg .b32 %r<3>;
261245
; CHECK-EMPTY:
262246
; CHECK-NEXT: // %bb.0:
263247
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_ule_param_0];
@@ -267,12 +251,10 @@ define i32 @icmp_i1_ule(i32 %a, i32 %b) {
267251
; CHECK-NEXT: and.pred %p3, %p1, %p2;
268252
; CHECK-NEXT: @%p3 bra $L__BB8_2;
269253
; CHECK-NEXT: // %bb.1: // %bb1
270-
; CHECK-NEXT: mov.b32 %r4, 1;
271-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
254+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
272255
; CHECK-NEXT: ret;
273256
; CHECK-NEXT: $L__BB8_2: // %bb2
274-
; CHECK-NEXT: mov.b32 %r3, 127;
275-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
257+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
276258
; CHECK-NEXT: ret;
277259
%p1 = icmp sgt i32 %a, 1
278260
%p2 = icmp sgt i32 %b, 1
@@ -288,7 +270,7 @@ define i32 @icmp_i1_ult(i32 %a, i32 %b) {
288270
; CHECK-LABEL: icmp_i1_ult(
289271
; CHECK: {
290272
; CHECK-NEXT: .reg .pred %p<4>;
291-
; CHECK-NEXT: .reg .b32 %r<5>;
273+
; CHECK-NEXT: .reg .b32 %r<3>;
292274
; CHECK-EMPTY:
293275
; CHECK-NEXT: // %bb.0:
294276
; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_ult_param_0];
@@ -298,12 +280,10 @@ define i32 @icmp_i1_ult(i32 %a, i32 %b) {
298280
; CHECK-NEXT: or.pred %p3, %p1, %p2;
299281
; CHECK-NEXT: @%p3 bra $L__BB9_2;
300282
; CHECK-NEXT: // %bb.1: // %bb1
301-
; CHECK-NEXT: mov.b32 %r4, 1;
302-
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
283+
; CHECK-NEXT: st.param.b32 [func_retval0], 1;
303284
; CHECK-NEXT: ret;
304285
; CHECK-NEXT: $L__BB9_2: // %bb2
305-
; CHECK-NEXT: mov.b32 %r3, 127;
306-
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
286+
; CHECK-NEXT: st.param.b32 [func_retval0], 127;
307287
; CHECK-NEXT: ret;
308288
%p1 = icmp sgt i32 %a, 1
309289
%p2 = icmp sgt i32 %b, 1

‎llvm/test/CodeGen/NVPTX/i128-ld-st.ll‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -7,15 +7,14 @@ target triple = "nvptx64-nvidia-cuda"
77
define i128 @foo(ptr %p, ptr %o) {
88
; CHECK-LABEL: foo(
99
; CHECK: {
10-
; CHECK-NEXT: .reg .b64 %rd<5>;
10+
; CHECK-NEXT: .reg .b64 %rd<4>;
1111
; CHECK-EMPTY:
1212
; CHECK-NEXT: // %bb.0:
1313
; CHECK-NEXT: ld.param.b64 %rd2, [foo_param_1];
1414
; CHECK-NEXT: ld.param.b64 %rd1, [foo_param_0];
1515
; CHECK-NEXT: ld.b8 %rd3, [%rd1];
1616
; CHECK-NEXT: st.v2.b64 [%rd2], {%rd3, 0};
17-
; CHECK-NEXT: mov.b64 %rd4, 0;
18-
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd4};
17+
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd3, 0};
1918
; CHECK-NEXT: ret;
2019
%c = load i8, ptr %p, align 1
2120
%i = zext i8 %c to i128

0 commit comments

Comments
 (0)