Skip to content

Commit b966989

Browse files
AlexMacleanAlexMaclean
committed
fixup tests after rebase
1 parent 0f5249c commit b966989

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
@@ -1305,10 +1305,10 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 {
13051305
; O0-NEXT: ld.param.b32 %r1, [test_call_param_0];
13061306
; O0-NEXT: { // callseq 0, 0
13071307
; O0-NEXT: .param .align 4 .b8 param0[4];
1308-
; O0-NEXT: st.param.b32 [param0], %r1;
13091308
; O0-NEXT: .param .align 4 .b8 param1[4];
1310-
; O0-NEXT: st.param.b32 [param1], %r2;
13111309
; O0-NEXT: .param .align 4 .b8 retval0[4];
1310+
; O0-NEXT: st.param.b32 [param1], %r2;
1311+
; O0-NEXT: st.param.b32 [param0], %r1;
13121312
; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
13131313
; O0-NEXT: ld.param.b32 %r3, [retval0];
13141314
; O0-NEXT: } // callseq 0
@@ -1321,13 +1321,13 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 {
13211321
; O3-EMPTY:
13221322
; O3-NEXT: // %bb.0:
13231323
; O3-NEXT: ld.param.b32 %r1, [test_call_param_0];
1324-
; O3-NEXT: ld.param.b32 %r2, [test_call_param_1];
13251324
; O3-NEXT: { // callseq 0, 0
13261325
; O3-NEXT: .param .align 4 .b8 param0[4];
1327-
; O3-NEXT: st.param.b32 [param0], %r1;
13281326
; O3-NEXT: .param .align 4 .b8 param1[4];
1329-
; O3-NEXT: st.param.b32 [param1], %r2;
13301327
; O3-NEXT: .param .align 4 .b8 retval0[4];
1328+
; O3-NEXT: ld.param.b32 %r2, [test_call_param_1];
1329+
; O3-NEXT: st.param.b32 [param1], %r2;
1330+
; O3-NEXT: st.param.b32 [param0], %r1;
13311331
; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
13321332
; O3-NEXT: ld.param.b32 %r3, [retval0];
13331333
; O3-NEXT: } // callseq 0
@@ -1347,10 +1347,10 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
13471347
; O0-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];
13481348
; O0-NEXT: { // callseq 1, 0
13491349
; O0-NEXT: .param .align 4 .b8 param0[4];
1350-
; O0-NEXT: st.param.b32 [param0], %r2;
13511350
; O0-NEXT: .param .align 4 .b8 param1[4];
1352-
; O0-NEXT: st.param.b32 [param1], %r1;
13531351
; O0-NEXT: .param .align 4 .b8 retval0[4];
1352+
; O0-NEXT: st.param.b32 [param1], %r1;
1353+
; O0-NEXT: st.param.b32 [param0], %r2;
13541354
; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
13551355
; O0-NEXT: ld.param.b32 %r3, [retval0];
13561356
; O0-NEXT: } // callseq 1
@@ -1363,13 +1363,13 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
13631363
; O3-EMPTY:
13641364
; O3-NEXT: // %bb.0:
13651365
; O3-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];
1366-
; O3-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];
13671366
; O3-NEXT: { // callseq 1, 0
13681367
; O3-NEXT: .param .align 4 .b8 param0[4];
1369-
; O3-NEXT: st.param.b32 [param0], %r2;
13701368
; O3-NEXT: .param .align 4 .b8 param1[4];
1371-
; O3-NEXT: st.param.b32 [param1], %r1;
13721369
; O3-NEXT: .param .align 4 .b8 retval0[4];
1370+
; O3-NEXT: st.param.b32 [param1], %r1;
1371+
; O3-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];
1372+
; O3-NEXT: st.param.b32 [param0], %r2;
13731373
; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
13741374
; O3-NEXT: ld.param.b32 %r3, [retval0];
13751375
; O3-NEXT: } // callseq 1
@@ -1389,10 +1389,10 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
13891389
; O0-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];
13901390
; O0-NEXT: { // callseq 2, 0
13911391
; O0-NEXT: .param .align 4 .b8 param0[4];
1392-
; O0-NEXT: st.param.b32 [param0], %r2;
13931392
; O0-NEXT: .param .align 4 .b8 param1[4];
1394-
; O0-NEXT: st.param.b32 [param1], %r1;
13951393
; O0-NEXT: .param .align 4 .b8 retval0[4];
1394+
; O0-NEXT: st.param.b32 [param1], %r1;
1395+
; O0-NEXT: st.param.b32 [param0], %r2;
13961396
; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
13971397
; O0-NEXT: ld.param.b32 %r3, [retval0];
13981398
; O0-NEXT: } // callseq 2
@@ -1405,13 +1405,13 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
14051405
; O3-EMPTY:
14061406
; O3-NEXT: // %bb.0:
14071407
; O3-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];
1408-
; O3-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];
14091408
; O3-NEXT: { // callseq 2, 0
14101409
; O3-NEXT: .param .align 4 .b8 param0[4];
1411-
; O3-NEXT: st.param.b32 [param0], %r2;
14121410
; O3-NEXT: .param .align 4 .b8 param1[4];
1413-
; O3-NEXT: st.param.b32 [param1], %r1;
14141411
; O3-NEXT: .param .align 4 .b8 retval0[4];
1412+
; O3-NEXT: st.param.b32 [param1], %r1;
1413+
; O3-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];
1414+
; O3-NEXT: st.param.b32 [param0], %r2;
14151415
; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
14161416
; O3-NEXT: ld.param.b32 %r3, [retval0];
14171417
; O3-NEXT: } // callseq 2

0 commit comments

Comments
 (0)