Skip to content

Commit e7c680c

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

File tree

5 files changed

+14
-48
lines changed

5 files changed

+14
-48
lines changed

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

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3513,9 +3513,7 @@ NVPTXTargetLowering::LowerReturn(SDValue Chain, CallingConv::ID CallConv,
35133513
Val = DAG.getBuildVector(VT, dl, StoreVals);
35143514
}
35153515

3516-
SDValue RetSymbol =
3517-
DAG.getNode(NVPTXISD::Wrapper, dl, MVT::i32,
3518-
DAG.getTargetExternalSymbol("func_retval0", MVT::i32));
3516+
SDValue RetSymbol = DAG.getExternalSymbol("func_retval0", MVT::i32);
35193517
SDValue Ptr =
35203518
DAG.getObjectPtrOffset(dl, RetSymbol, TypeSize::getFixed(Offsets[I]));
35213519

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

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -169,11 +169,7 @@ define %struct.S16 @i32_to_2xi16_shr(i32 noundef %i){
169169
; CHECK-NEXT: { // callseq 0, 0
170170
; CHECK-NEXT: .param .b32 param0;
171171
; CHECK-NEXT: st.param.b32 [param0], %r1;
172-
; CHECK-NEXT: call.uni
173-
; CHECK-NEXT: escape_int,
174-
; CHECK-NEXT: (
175-
; CHECK-NEXT: param0
176-
; CHECK-NEXT: );
172+
; CHECK-NEXT: call.uni escape_int, (param0);
177173
; CHECK-NEXT: } // callseq 0
178174
; CHECK-NEXT: shr.s32 %r2, %r1, 16;
179175
; CHECK-NEXT: shr.u32 %r3, %r2, 16;

‎llvm/test/CodeGen/NVPTX/param-add.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ define i32 @test(%struct.1float alignstack(32) %data) {
2929
; CHECK-NEXT: st.param.b8 [param0+3], %r4;
3030
; CHECK-NEXT: .param .b32 retval0;
3131
; CHECK-NEXT: call.uni (retval0), callee, (param0);
32-
; CHECK-NEXT: ld.param.b32 %r14, [retval0];
32+
; CHECK-NEXT: ld.param.b32 %r5, [retval0];
3333
; CHECK-NEXT: } // callseq 0
3434
; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
3535
; CHECK-NEXT: ret;

‎llvm/test/CodeGen/NVPTX/param-overalign.ll‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -106,10 +106,10 @@ define alignstack(8) %struct.float2 @aligned_return(%struct.float2 %a ) {
106106
; CHECK-NEXT: .reg .b32 %r<3>;
107107
; CHECK-EMPTY:
108108
; CHECK-NEXT: // %bb.0:
109-
; CHECK-NEXT: ld.param.b32 %r1, [aligned_return_param_0+4];
110-
; CHECK-NEXT: ld.param.b32 %r2, [aligned_return_param_0];
111-
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
112-
; CHECK-NEXT: st.param.b32 [func_retval0+4], %r1;
109+
; CHECK-NEXT: ld.param.b32 %r1, [aligned_return_param_0];
110+
; CHECK-NEXT: ld.param.b32 %r2, [aligned_return_param_0+4];
111+
; CHECK-NEXT: st.param.b32 [func_retval0+4], %r2;
112+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
113113
; CHECK-NEXT: ret;
114114
ret %struct.float2 %a
115115
}

‎llvm/test/CodeGen/NVPTX/unaligned-param-load-store.ll‎

Lines changed: 7 additions & 35 deletions
Original file line numberDiff line numberDiff line change
@@ -43,11 +43,7 @@ define %s_i8i16p @test_s_i8i16p(%s_i8i16p %a) {
4343
; CHECK-NEXT: st.param.b8 [param0+4], %rs4;
4444
; CHECK-NEXT: st.param.b64 [param0+8], %rd1;
4545
; CHECK-NEXT: .param .align 8 .b8 retval0[16];
46-
; CHECK-NEXT: call.uni (retval0),
47-
; CHECK-NEXT: test_s_i8i16p,
48-
; CHECK-NEXT: (
49-
; CHECK-NEXT: param0
50-
; CHECK-NEXT: );
46+
; CHECK-NEXT: call.uni (retval0), test_s_i8i16p, (param0);
5147
; CHECK-NEXT: ld.param.b16 %rs7, [retval0];
5248
; CHECK-NEXT: ld.param.b8 %rs8, [retval0+2];
5349
; CHECK-NEXT: ld.param.b8 %rs9, [retval0+3];
@@ -98,11 +94,7 @@ define %s_i8i32p @test_s_i8i32p(%s_i8i32p %a) {
9894
; CHECK-NEXT: st.param.b8 [param0+8], %r9;
9995
; CHECK-NEXT: st.param.b64 [param0+16], %rd1;
10096
; CHECK-NEXT: .param .align 8 .b8 retval0[24];
101-
; CHECK-NEXT: call.uni (retval0),
102-
; CHECK-NEXT: test_s_i8i32p,
103-
; CHECK-NEXT: (
104-
; CHECK-NEXT: param0
105-
; CHECK-NEXT: );
97+
; CHECK-NEXT: call.uni (retval0), test_s_i8i32p, (param0);
10698
; CHECK-NEXT: ld.param.b32 %r14, [retval0];
10799
; CHECK-NEXT: ld.param.b8 %rs2, [retval0+4];
108100
; CHECK-NEXT: ld.param.b8 %rs3, [retval0+5];
@@ -180,11 +172,7 @@ define %s_i8i64p @test_s_i8i64p(%s_i8i64p %a) {
180172
; CHECK-NEXT: st.param.b8 [param0+16], %rd30;
181173
; CHECK-NEXT: st.param.b64 [param0+24], %rd3;
182174
; CHECK-NEXT: .param .align 8 .b8 retval0[32];
183-
; CHECK-NEXT: call.uni (retval0),
184-
; CHECK-NEXT: test_s_i8i64p,
185-
; CHECK-NEXT: (
186-
; CHECK-NEXT: param0
187-
; CHECK-NEXT: );
175+
; CHECK-NEXT: call.uni (retval0), test_s_i8i64p, (param0);
188176
; CHECK-NEXT: ld.param.b64 %rd31, [retval0];
189177
; CHECK-NEXT: ld.param.b8 %rs2, [retval0+8];
190178
; CHECK-NEXT: ld.param.b8 %rs3, [retval0+9];
@@ -269,11 +257,7 @@ define %s_i8f16p @test_s_i8f16p(%s_i8f16p %a) {
269257
; CHECK-NEXT: st.param.b8 [param0+4], %rs4;
270258
; CHECK-NEXT: st.param.b64 [param0+8], %rd1;
271259
; CHECK-NEXT: .param .align 8 .b8 retval0[16];
272-
; CHECK-NEXT: call.uni (retval0),
273-
; CHECK-NEXT: test_s_i8f16p,
274-
; CHECK-NEXT: (
275-
; CHECK-NEXT: param0
276-
; CHECK-NEXT: );
260+
; CHECK-NEXT: call.uni (retval0), test_s_i8f16p, (param0);
277261
; CHECK-NEXT: ld.param.b16 %rs7, [retval0];
278262
; CHECK-NEXT: ld.param.b8 %rs8, [retval0+2];
279263
; CHECK-NEXT: ld.param.b8 %rs9, [retval0+3];
@@ -324,11 +308,7 @@ define %s_i8f16x2p @test_s_i8f16x2p(%s_i8f16x2p %a) {
324308
; CHECK-NEXT: st.param.b8 [param0+8], %r9;
325309
; CHECK-NEXT: st.param.b64 [param0+16], %rd1;
326310
; CHECK-NEXT: .param .align 8 .b8 retval0[24];
327-
; CHECK-NEXT: call.uni (retval0),
328-
; CHECK-NEXT: test_s_i8f16x2p,
329-
; CHECK-NEXT: (
330-
; CHECK-NEXT: param0
331-
; CHECK-NEXT: );
311+
; CHECK-NEXT: call.uni (retval0), test_s_i8f16x2p, (param0);
332312
; CHECK-NEXT: ld.param.b32 %r14, [retval0];
333313
; CHECK-NEXT: ld.param.b8 %rs2, [retval0+4];
334314
; CHECK-NEXT: ld.param.b8 %rs3, [retval0+5];
@@ -387,11 +367,7 @@ define %s_i8f32p @test_s_i8f32p(%s_i8f32p %a) {
387367
; CHECK-NEXT: st.param.b8 [param0+8], %r9;
388368
; CHECK-NEXT: st.param.b64 [param0+16], %rd1;
389369
; CHECK-NEXT: .param .align 8 .b8 retval0[24];
390-
; CHECK-NEXT: call.uni (retval0),
391-
; CHECK-NEXT: test_s_i8f32p,
392-
; CHECK-NEXT: (
393-
; CHECK-NEXT: param0
394-
; CHECK-NEXT: );
370+
; CHECK-NEXT: call.uni (retval0), test_s_i8f32p, (param0);
395371
; CHECK-NEXT: ld.param.b32 %r14, [retval0];
396372
; CHECK-NEXT: ld.param.b8 %rs2, [retval0+4];
397373
; CHECK-NEXT: ld.param.b8 %rs3, [retval0+5];
@@ -469,11 +445,7 @@ define %s_i8f64p @test_s_i8f64p(%s_i8f64p %a) {
469445
; CHECK-NEXT: st.param.b8 [param0+16], %rd30;
470446
; CHECK-NEXT: st.param.b64 [param0+24], %rd3;
471447
; CHECK-NEXT: .param .align 8 .b8 retval0[32];
472-
; CHECK-NEXT: call.uni (retval0),
473-
; CHECK-NEXT: test_s_i8f64p,
474-
; CHECK-NEXT: (
475-
; CHECK-NEXT: param0
476-
; CHECK-NEXT: );
448+
; CHECK-NEXT: call.uni (retval0), test_s_i8f64p, (param0);
477449
; CHECK-NEXT: ld.param.b64 %rd31, [retval0];
478450
; CHECK-NEXT: ld.param.b8 %rs2, [retval0+8];
479451
; CHECK-NEXT: ld.param.b8 %rs3, [retval0+9];

0 commit comments

Comments
 (0)