Skip to content

Commit 15c4720

Browse files
AlexMacleanAlexMaclean
committed
fixup tests after rebase
1 parent dc4b6f5 commit 15c4720

File tree

5 files changed

+101
-122
lines changed

5 files changed

+101
-122
lines changed

‎llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5732,7 +5732,6 @@ static SDValue combinePRMT(SDNode *N, TargetLowering::DAGCombinerInfo &DCI,
57325732
return SDValue();
57335733
}
57345734

5735-
57365735
// During call lowering we wrap the return values in a ProxyReg node which
57375736
// depend on the chain value produced by the completed call. This ensures that
57385737
// the full call is emitted in cases where libcalls are used to legalize

‎llvm/test/CodeGen/NVPTX/aggregate-return.ll‎

Lines changed: 24 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -16,8 +16,8 @@ define void @test_v2f32(<2 x float> %input, ptr %output) {
1616
; CHECK-NEXT: ld.param.b64 %rd1, [test_v2f32_param_0];
1717
; CHECK-NEXT: { // callseq 0, 0
1818
; CHECK-NEXT: .param .align 8 .b8 param0[8];
19-
; CHECK-NEXT: st.param.b64 [param0], %rd1;
2019
; CHECK-NEXT: .param .align 8 .b8 retval0[8];
20+
; CHECK-NEXT: st.param.b64 [param0], %rd1;
2121
; CHECK-NEXT: call.uni (retval0), barv, (param0);
2222
; CHECK-NEXT: ld.param.b64 %rd2, [retval0];
2323
; CHECK-NEXT: } // callseq 0
@@ -32,24 +32,24 @@ define void @test_v2f32(<2 x float> %input, ptr %output) {
3232
define void @test_v3f32(<3 x float> %input, ptr %output) {
3333
; CHECK-LABEL: test_v3f32(
3434
; CHECK: {
35-
; CHECK-NEXT: .reg .b32 %r<10>;
36-
; CHECK-NEXT: .reg .b64 %rd<2>;
35+
; CHECK-NEXT: .reg .b32 %r<4>;
36+
; CHECK-NEXT: .reg .b64 %rd<5>;
3737
; CHECK-EMPTY:
3838
; CHECK-NEXT: // %bb.0:
39-
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_v3f32_param_0];
40-
; CHECK-NEXT: ld.param.b32 %r3, [test_v3f32_param_0+8];
39+
; CHECK-NEXT: ld.param.b64 %rd1, [test_v3f32_param_0];
40+
; CHECK-NEXT: ld.param.b32 %r1, [test_v3f32_param_0+8];
4141
; CHECK-NEXT: { // callseq 1, 0
4242
; CHECK-NEXT: .param .align 16 .b8 param0[16];
43-
; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, %r2};
44-
; CHECK-NEXT: st.param.b32 [param0+8], %r3;
4543
; CHECK-NEXT: .param .align 16 .b8 retval0[16];
44+
; CHECK-NEXT: st.param.b32 [param0+8], %r1;
45+
; CHECK-NEXT: st.param.b64 [param0], %rd1;
4646
; CHECK-NEXT: call.uni (retval0), barv3, (param0);
47-
; CHECK-NEXT: ld.param.v2.b32 {%r4, %r5}, [retval0];
48-
; CHECK-NEXT: ld.param.b32 %r6, [retval0+8];
47+
; CHECK-NEXT: ld.param.b32 %r2, [retval0+8];
48+
; CHECK-NEXT: ld.param.b64 %rd2, [retval0];
4949
; CHECK-NEXT: } // callseq 1
50-
; CHECK-NEXT: ld.param.b64 %rd1, [test_v3f32_param_1];
51-
; CHECK-NEXT: st.v2.b32 [%rd1], {%r4, %r5};
52-
; CHECK-NEXT: st.b32 [%rd1+8], %r6;
50+
; CHECK-NEXT: ld.param.b64 %rd4, [test_v3f32_param_1];
51+
; CHECK-NEXT: st.b32 [%rd4+8], %r2;
52+
; CHECK-NEXT: st.b64 [%rd4], %rd2;
5353
; CHECK-NEXT: ret;
5454
%call = tail call <3 x float> @barv3(<3 x float> %input)
5555
; Make sure we don't load more values than than we need to.
@@ -68,16 +68,16 @@ define void @test_a2f32([2 x float] %input, ptr %output) {
6868
; CHECK-NEXT: ld.param.b32 %r2, [test_a2f32_param_0+4];
6969
; CHECK-NEXT: { // callseq 2, 0
7070
; CHECK-NEXT: .param .align 4 .b8 param0[8];
71-
; CHECK-NEXT: st.param.b32 [param0], %r1;
72-
; CHECK-NEXT: st.param.b32 [param0+4], %r2;
7371
; CHECK-NEXT: .param .align 4 .b8 retval0[8];
72+
; CHECK-NEXT: st.param.b32 [param0+4], %r2;
73+
; CHECK-NEXT: st.param.b32 [param0], %r1;
7474
; CHECK-NEXT: call.uni (retval0), bara, (param0);
75-
; CHECK-NEXT: ld.param.b32 %r3, [retval0];
76-
; CHECK-NEXT: ld.param.b32 %r4, [retval0+4];
75+
; CHECK-NEXT: ld.param.b32 %r3, [retval0+4];
76+
; CHECK-NEXT: ld.param.b32 %r4, [retval0];
7777
; CHECK-NEXT: } // callseq 2
7878
; CHECK-NEXT: ld.param.b64 %rd1, [test_a2f32_param_1];
79-
; CHECK-NEXT: st.b32 [%rd1+4], %r4;
80-
; CHECK-NEXT: st.b32 [%rd1], %r3;
79+
; CHECK-NEXT: st.b32 [%rd1+4], %r3;
80+
; CHECK-NEXT: st.b32 [%rd1], %r4;
8181
; CHECK-NEXT: ret;
8282
%call = tail call [2 x float] @bara([2 x float] %input)
8383
store [2 x float] %call, ptr %output, align 4
@@ -95,16 +95,16 @@ define void @test_s2f32({float, float} %input, ptr %output) {
9595
; CHECK-NEXT: ld.param.b32 %r2, [test_s2f32_param_0+4];
9696
; CHECK-NEXT: { // callseq 3, 0
9797
; CHECK-NEXT: .param .align 4 .b8 param0[8];
98-
; CHECK-NEXT: st.param.b32 [param0], %r1;
99-
; CHECK-NEXT: st.param.b32 [param0+4], %r2;
10098
; CHECK-NEXT: .param .align 4 .b8 retval0[8];
99+
; CHECK-NEXT: st.param.b32 [param0+4], %r2;
100+
; CHECK-NEXT: st.param.b32 [param0], %r1;
101101
; CHECK-NEXT: call.uni (retval0), bars, (param0);
102-
; CHECK-NEXT: ld.param.b32 %r3, [retval0];
103-
; CHECK-NEXT: ld.param.b32 %r4, [retval0+4];
102+
; CHECK-NEXT: ld.param.b32 %r3, [retval0+4];
103+
; CHECK-NEXT: ld.param.b32 %r4, [retval0];
104104
; CHECK-NEXT: } // callseq 3
105105
; CHECK-NEXT: ld.param.b64 %rd1, [test_s2f32_param_1];
106-
; CHECK-NEXT: st.b32 [%rd1+4], %r4;
107-
; CHECK-NEXT: st.b32 [%rd1], %r3;
106+
; CHECK-NEXT: st.b32 [%rd1+4], %r3;
107+
; CHECK-NEXT: st.b32 [%rd1], %r4;
108108
; CHECK-NEXT: ret;
109109
%call = tail call {float, float} @bars({float, float} %input)
110110
store {float, float} %call, ptr %output, align 4

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

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -859,10 +859,10 @@ define <2 x float> @test_call(<2 x float> %a, <2 x float> %b) #0 {
859859
; CHECK-NEXT: ld.param.b64 %rd1, [test_call_param_0];
860860
; CHECK-NEXT: { // callseq 0, 0
861861
; CHECK-NEXT: .param .align 8 .b8 param0[8];
862-
; CHECK-NEXT: st.param.b64 [param0], %rd1;
863862
; CHECK-NEXT: .param .align 8 .b8 param1[8];
864-
; CHECK-NEXT: st.param.b64 [param1], %rd2;
865863
; CHECK-NEXT: .param .align 8 .b8 retval0[8];
864+
; CHECK-NEXT: st.param.b64 [param1], %rd2;
865+
; CHECK-NEXT: st.param.b64 [param0], %rd1;
866866
; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);
867867
; CHECK-NEXT: ld.param.b64 %rd3, [retval0];
868868
; CHECK-NEXT: } // callseq 0
@@ -882,10 +882,10 @@ define <2 x float> @test_call_flipped(<2 x float> %a, <2 x float> %b) #0 {
882882
; CHECK-NEXT: ld.param.b64 %rd1, [test_call_flipped_param_0];
883883
; CHECK-NEXT: { // callseq 1, 0
884884
; CHECK-NEXT: .param .align 8 .b8 param0[8];
885-
; CHECK-NEXT: st.param.b64 [param0], %rd2;
886885
; CHECK-NEXT: .param .align 8 .b8 param1[8];
887-
; CHECK-NEXT: st.param.b64 [param1], %rd1;
888886
; CHECK-NEXT: .param .align 8 .b8 retval0[8];
887+
; CHECK-NEXT: st.param.b64 [param1], %rd1;
888+
; CHECK-NEXT: st.param.b64 [param0], %rd2;
889889
; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);
890890
; CHECK-NEXT: ld.param.b64 %rd3, [retval0];
891891
; CHECK-NEXT: } // callseq 1
@@ -905,10 +905,10 @@ define <2 x float> @test_tailcall_flipped(<2 x float> %a, <2 x float> %b) #0 {
905905
; CHECK-NEXT: ld.param.b64 %rd1, [test_tailcall_flipped_param_0];
906906
; CHECK-NEXT: { // callseq 2, 0
907907
; CHECK-NEXT: .param .align 8 .b8 param0[8];
908-
; CHECK-NEXT: st.param.b64 [param0], %rd2;
909908
; CHECK-NEXT: .param .align 8 .b8 param1[8];
910-
; CHECK-NEXT: st.param.b64 [param1], %rd1;
911909
; CHECK-NEXT: .param .align 8 .b8 retval0[8];
910+
; CHECK-NEXT: st.param.b64 [param1], %rd1;
911+
; CHECK-NEXT: st.param.b64 [param0], %rd2;
912912
; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);
913913
; CHECK-NEXT: ld.param.b64 %rd3, [retval0];
914914
; CHECK-NEXT: } // callseq 2

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

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -1273,10 +1273,10 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 {
12731273
; O0-NEXT: ld.param.b32 %r1, [test_call_param_0];
12741274
; O0-NEXT: { // callseq 0, 0
12751275
; O0-NEXT: .param .align 4 .b8 param0[4];
1276-
; O0-NEXT: st.param.b32 [param0], %r1;
12771276
; O0-NEXT: .param .align 4 .b8 param1[4];
1278-
; O0-NEXT: st.param.b32 [param1], %r2;
12791277
; O0-NEXT: .param .align 4 .b8 retval0[4];
1278+
; O0-NEXT: st.param.b32 [param1], %r2;
1279+
; O0-NEXT: st.param.b32 [param0], %r1;
12801280
; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
12811281
; O0-NEXT: ld.param.b32 %r3, [retval0];
12821282
; O0-NEXT: } // callseq 0
@@ -1289,13 +1289,13 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 {
12891289
; O3-EMPTY:
12901290
; O3-NEXT: // %bb.0:
12911291
; O3-NEXT: ld.param.b32 %r1, [test_call_param_0];
1292-
; O3-NEXT: ld.param.b32 %r2, [test_call_param_1];
12931292
; O3-NEXT: { // callseq 0, 0
12941293
; O3-NEXT: .param .align 4 .b8 param0[4];
1295-
; O3-NEXT: st.param.b32 [param0], %r1;
12961294
; O3-NEXT: .param .align 4 .b8 param1[4];
1297-
; O3-NEXT: st.param.b32 [param1], %r2;
12981295
; O3-NEXT: .param .align 4 .b8 retval0[4];
1296+
; O3-NEXT: ld.param.b32 %r2, [test_call_param_1];
1297+
; O3-NEXT: st.param.b32 [param1], %r2;
1298+
; O3-NEXT: st.param.b32 [param0], %r1;
12991299
; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
13001300
; O3-NEXT: ld.param.b32 %r3, [retval0];
13011301
; O3-NEXT: } // callseq 0
@@ -1315,10 +1315,10 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
13151315
; O0-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];
13161316
; O0-NEXT: { // callseq 1, 0
13171317
; O0-NEXT: .param .align 4 .b8 param0[4];
1318-
; O0-NEXT: st.param.b32 [param0], %r2;
13191318
; O0-NEXT: .param .align 4 .b8 param1[4];
1320-
; O0-NEXT: st.param.b32 [param1], %r1;
13211319
; O0-NEXT: .param .align 4 .b8 retval0[4];
1320+
; O0-NEXT: st.param.b32 [param1], %r1;
1321+
; O0-NEXT: st.param.b32 [param0], %r2;
13221322
; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
13231323
; O0-NEXT: ld.param.b32 %r3, [retval0];
13241324
; O0-NEXT: } // callseq 1
@@ -1331,13 +1331,13 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
13311331
; O3-EMPTY:
13321332
; O3-NEXT: // %bb.0:
13331333
; O3-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];
1334-
; O3-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];
13351334
; O3-NEXT: { // callseq 1, 0
13361335
; O3-NEXT: .param .align 4 .b8 param0[4];
1337-
; O3-NEXT: st.param.b32 [param0], %r2;
13381336
; O3-NEXT: .param .align 4 .b8 param1[4];
1339-
; O3-NEXT: st.param.b32 [param1], %r1;
13401337
; O3-NEXT: .param .align 4 .b8 retval0[4];
1338+
; O3-NEXT: st.param.b32 [param1], %r1;
1339+
; O3-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];
1340+
; O3-NEXT: st.param.b32 [param0], %r2;
13411341
; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
13421342
; O3-NEXT: ld.param.b32 %r3, [retval0];
13431343
; O3-NEXT: } // callseq 1
@@ -1357,10 +1357,10 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
13571357
; O0-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];
13581358
; O0-NEXT: { // callseq 2, 0
13591359
; O0-NEXT: .param .align 4 .b8 param0[4];
1360-
; O0-NEXT: st.param.b32 [param0], %r2;
13611360
; O0-NEXT: .param .align 4 .b8 param1[4];
1362-
; O0-NEXT: st.param.b32 [param1], %r1;
13631361
; O0-NEXT: .param .align 4 .b8 retval0[4];
1362+
; O0-NEXT: st.param.b32 [param1], %r1;
1363+
; O0-NEXT: st.param.b32 [param0], %r2;
13641364
; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
13651365
; O0-NEXT: ld.param.b32 %r3, [retval0];
13661366
; O0-NEXT: } // callseq 2
@@ -1373,13 +1373,13 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
13731373
; O3-EMPTY:
13741374
; O3-NEXT: // %bb.0:
13751375
; O3-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];
1376-
; O3-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];
13771376
; O3-NEXT: { // callseq 2, 0
13781377
; O3-NEXT: .param .align 4 .b8 param0[4];
1379-
; O3-NEXT: st.param.b32 [param0], %r2;
13801378
; O3-NEXT: .param .align 4 .b8 param1[4];
1381-
; O3-NEXT: st.param.b32 [param1], %r1;
13821379
; O3-NEXT: .param .align 4 .b8 retval0[4];
1380+
; O3-NEXT: st.param.b32 [param1], %r1;
1381+
; O3-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];
1382+
; O3-NEXT: st.param.b32 [param0], %r2;
13831383
; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
13841384
; O3-NEXT: ld.param.b32 %r3, [retval0];
13851385
; O3-NEXT: } // callseq 2

0 commit comments

Comments
 (0)