Skip to content

Commit ea0147c

Browse files
AlexMacleanAlexMaclean
committed
rebase + address comments
1 parent 1a4100e commit ea0147c

File tree

6 files changed

+39
-45
lines changed

6 files changed

+39
-45
lines changed

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -409,11 +409,11 @@ static void adjustByValArgAlignment(Argument *Arg, Value *ArgInParamAS,
409409
Worklist.pop();
410410

411411
for (User *CurUser : Ctx.InitialVal->users()) {
412-
if (auto *I = dyn_cast<LoadInst>(CurUser)) {
412+
if (auto *I = dyn_cast<LoadInst>(CurUser))
413413
Loads.push_back({I, Ctx.Offset});
414-
} else if (isa<BitCastInst>(CurUser) || isa<AddrSpaceCastInst>(CurUser)) {
414+
else if (isa<BitCastInst>(CurUser) || isa<AddrSpaceCastInst>(CurUser))
415415
Worklist.push({cast<Instruction>(CurUser), Ctx.Offset});
416-
} else if (auto *I = dyn_cast<GetElementPtrInst>(CurUser)) {
416+
else if (auto *I = dyn_cast<GetElementPtrInst>(CurUser)) {
417417
APInt OffsetAccumulated =
418418
APInt::getZero(DL.getIndexSizeInBits(ADDRESS_SPACE_PARAM));
419419

‎llvm/test/CodeGen/NVPTX/forward-ld-param.ll‎

Lines changed: 8 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@ define i32 @test_ld_param_const(ptr byval(i32) %a) {
77
; CHECK-LABEL: test_ld_param_const(
88
; CHECK: {
99
; CHECK-NEXT: .reg .b32 %r<2>;
10-
; CHECK-NEXT: .reg .b64 %rd<4>;
10+
; CHECK-NEXT: .reg .b64 %rd<2>;
1111
; CHECK-EMPTY:
1212
; CHECK-NEXT: // %bb.0:
1313
; CHECK-NEXT: ld.param.u32 %r1, [test_ld_param_const_param_0+4];
@@ -22,15 +22,13 @@ define i32 @test_ld_param_non_const(ptr byval([10 x i32]) %a, i32 %b) {
2222
; CHECK-LABEL: test_ld_param_non_const(
2323
; CHECK: {
2424
; CHECK-NEXT: .reg .b32 %r<2>;
25-
; CHECK-NEXT: .reg .b64 %rd<6>;
25+
; CHECK-NEXT: .reg .b64 %rd<4>;
2626
; CHECK-EMPTY:
2727
; CHECK-NEXT: // %bb.0:
2828
; CHECK-NEXT: mov.b64 %rd1, test_ld_param_non_const_param_0;
29-
; CHECK-NEXT: cvta.local.u64 %rd2, %rd1;
30-
; CHECK-NEXT: cvta.to.local.u64 %rd3, %rd2;
31-
; CHECK-NEXT: ld.param.s32 %rd4, [test_ld_param_non_const_param_1];
32-
; CHECK-NEXT: add.s64 %rd5, %rd3, %rd4;
33-
; CHECK-NEXT: ld.local.u32 %r1, [%rd5];
29+
; CHECK-NEXT: ld.param.s32 %rd2, [test_ld_param_non_const_param_1];
30+
; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;
31+
; CHECK-NEXT: ld.local.u32 %r1, [%rd3];
3432
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
3533
; CHECK-NEXT: ret;
3634
%p2 = getelementptr i8, ptr %a, i32 %b
@@ -89,15 +87,13 @@ define i32 @test_modify_param(ptr byval([10 x i32]) %a, i32 %b, i32 %c ) {
8987
; CHECK-LABEL: test_modify_param(
9088
; CHECK: {
9189
; CHECK-NEXT: .reg .b32 %r<3>;
92-
; CHECK-NEXT: .reg .b64 %rd<4>;
90+
; CHECK-NEXT: .reg .b64 %rd<2>;
9391
; CHECK-EMPTY:
9492
; CHECK-NEXT: // %bb.0:
9593
; CHECK-NEXT: mov.b64 %rd1, test_modify_param_param_0;
96-
; CHECK-NEXT: cvta.local.u64 %rd2, %rd1;
97-
; CHECK-NEXT: cvta.to.local.u64 %rd3, %rd2;
9894
; CHECK-NEXT: ld.param.u32 %r1, [test_modify_param_param_1];
9995
; CHECK-NEXT: ld.param.u32 %r2, [test_modify_param_param_2];
100-
; CHECK-NEXT: st.local.u32 [%rd3+2], %r1;
96+
; CHECK-NEXT: st.local.u32 [%rd1+2], %r1;
10197
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
10298
; CHECK-NEXT: ret;
10399
%p2 = getelementptr i8, ptr %a, i32 2
@@ -111,7 +107,7 @@ define i32 @test_multi_block(ptr byval([10 x i32]) %a, i1 %p) {
111107
; CHECK-NEXT: .reg .pred %p<3>;
112108
; CHECK-NEXT: .reg .b16 %rs<3>;
113109
; CHECK-NEXT: .reg .b32 %r<5>;
114-
; CHECK-NEXT: .reg .b64 %rd<4>;
110+
; CHECK-NEXT: .reg .b64 %rd<2>;
115111
; CHECK-EMPTY:
116112
; CHECK-NEXT: // %bb.0:
117113
; CHECK-NEXT: ld.param.u8 %rs1, [test_multi_block_param_1];

‎llvm/test/CodeGen/NVPTX/i128-array.ll‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -27,15 +27,15 @@ define [2 x i128] @foo(i64 %a, i32 %b) {
2727
define [2 x i128] @foo2(ptr byval([2 x i128]) %a) {
2828
; CHECK-LABEL: foo2(
2929
; CHECK: {
30-
; CHECK-NEXT: .reg .b64 %rd<9>;
30+
; CHECK-NEXT: .reg .b64 %rd<7>;
3131
; CHECK-EMPTY:
3232
; CHECK-NEXT: // %bb.0:
33-
; CHECK-NEXT: ld.param.u64 %rd5, [foo2_param_0+8];
34-
; CHECK-NEXT: ld.param.u64 %rd6, [foo2_param_0];
35-
; CHECK-NEXT: ld.param.u64 %rd7, [foo2_param_0+24];
36-
; CHECK-NEXT: ld.param.u64 %rd8, [foo2_param_0+16];
37-
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd5};
38-
; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd8, %rd7};
33+
; CHECK-NEXT: ld.param.u64 %rd3, [foo2_param_0+8];
34+
; CHECK-NEXT: ld.param.u64 %rd4, [foo2_param_0];
35+
; CHECK-NEXT: ld.param.u64 %rd5, [foo2_param_0+24];
36+
; CHECK-NEXT: ld.param.u64 %rd6, [foo2_param_0+16];
37+
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd3};
38+
; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd6, %rd5};
3939
; CHECK-NEXT: ret;
4040
%ptr0 = getelementptr [2 x i128], ptr %a, i64 0, i32 0
4141
%1 = load i128, i128* %ptr0

‎llvm/test/CodeGen/NVPTX/lower-args.ll‎

Lines changed: 10 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -33,30 +33,28 @@ define void @load_alignment(ptr nocapture readonly byval(%class.outer) align 8 %
3333
; PTX-LABEL: load_alignment(
3434
; PTX: {
3535
; PTX-NEXT: .reg .b32 %r<4>;
36-
; PTX-NEXT: .reg .b64 %rd<10>;
36+
; PTX-NEXT: .reg .b64 %rd<8>;
3737
; PTX-EMPTY:
3838
; PTX-NEXT: // %bb.0: // %entry
3939
; PTX-NEXT: mov.b64 %rd1, load_alignment_param_0;
40-
; PTX-NEXT: cvta.local.u64 %rd2, %rd1;
41-
; PTX-NEXT: cvta.to.local.u64 %rd3, %rd2;
42-
; PTX-NEXT: ld.local.u64 %rd4, [%rd3];
43-
; PTX-NEXT: ld.local.u64 %rd5, [%rd3+8];
44-
; PTX-NEXT: add.s64 %rd6, %rd3, 16;
45-
; PTX-NEXT: cvta.local.u64 %rd7, %rd6;
46-
; PTX-NEXT: ld.local.u32 %r1, [%rd3+16];
47-
; PTX-NEXT: ld.u32 %r2, [%rd4];
40+
; PTX-NEXT: ld.local.u64 %rd2, [%rd1];
41+
; PTX-NEXT: ld.local.u64 %rd3, [%rd1+8];
42+
; PTX-NEXT: add.s64 %rd4, %rd1, 16;
43+
; PTX-NEXT: cvta.local.u64 %rd5, %rd4;
44+
; PTX-NEXT: ld.local.u32 %r1, [%rd1+16];
45+
; PTX-NEXT: ld.u32 %r2, [%rd2];
4846
; PTX-NEXT: add.s32 %r3, %r2, %r1;
49-
; PTX-NEXT: st.u32 [%rd5], %r3;
47+
; PTX-NEXT: st.u32 [%rd3], %r3;
5048
; PTX-NEXT: { // callseq 0, 0
5149
; PTX-NEXT: .param .b64 param0;
52-
; PTX-NEXT: st.param.b64 [param0], %rd7;
50+
; PTX-NEXT: st.param.b64 [param0], %rd5;
5351
; PTX-NEXT: .param .b64 retval0;
5452
; PTX-NEXT: call.uni (retval0),
5553
; PTX-NEXT: escape,
5654
; PTX-NEXT: (
5755
; PTX-NEXT: param0
5856
; PTX-NEXT: );
59-
; PTX-NEXT: ld.param.b64 %rd8, [retval0];
57+
; PTX-NEXT: ld.param.b64 %rd6, [retval0];
6058
; PTX-NEXT: } // callseq 0
6159
; PTX-NEXT: ret;
6260
entry:

‎llvm/test/CodeGen/NVPTX/variadics-backend.ll‎

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -338,18 +338,18 @@ define dso_local i32 @variadics4(ptr noundef byval(%struct.S2) align 8 %first, .
338338
; CHECK-PTX-LABEL: variadics4(
339339
; CHECK-PTX: {
340340
; CHECK-PTX-NEXT: .reg .b32 %r<2>;
341-
; CHECK-PTX-NEXT: .reg .b64 %rd<12>;
341+
; CHECK-PTX-NEXT: .reg .b64 %rd<10>;
342342
; CHECK-PTX-EMPTY:
343343
; CHECK-PTX-NEXT: // %bb.0: // %entry
344-
; CHECK-PTX-NEXT: ld.param.u64 %rd4, [variadics4_param_1];
345-
; CHECK-PTX-NEXT: add.s64 %rd5, %rd4, 7;
346-
; CHECK-PTX-NEXT: and.b64 %rd6, %rd5, -8;
347-
; CHECK-PTX-NEXT: ld.u64 %rd7, [%rd6];
348-
; CHECK-PTX-NEXT: ld.param.u64 %rd8, [variadics4_param_0];
349-
; CHECK-PTX-NEXT: ld.param.u64 %rd9, [variadics4_param_0+8];
350-
; CHECK-PTX-NEXT: add.s64 %rd10, %rd8, %rd9;
351-
; CHECK-PTX-NEXT: add.s64 %rd11, %rd10, %rd7;
352-
; CHECK-PTX-NEXT: cvt.u32.u64 %r1, %rd11;
344+
; CHECK-PTX-NEXT: ld.param.u64 %rd2, [variadics4_param_1];
345+
; CHECK-PTX-NEXT: add.s64 %rd3, %rd2, 7;
346+
; CHECK-PTX-NEXT: and.b64 %rd4, %rd3, -8;
347+
; CHECK-PTX-NEXT: ld.u64 %rd5, [%rd4];
348+
; CHECK-PTX-NEXT: ld.param.u64 %rd6, [variadics4_param_0];
349+
; CHECK-PTX-NEXT: ld.param.u64 %rd7, [variadics4_param_0+8];
350+
; CHECK-PTX-NEXT: add.s64 %rd8, %rd6, %rd7;
351+
; CHECK-PTX-NEXT: add.s64 %rd9, %rd8, %rd5;
352+
; CHECK-PTX-NEXT: cvt.u32.u64 %r1, %rd9;
353353
; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %r1;
354354
; CHECK-PTX-NEXT: ret;
355355
entry:

‎llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -52,7 +52,7 @@ define internal fastcc [4 x i64] @callee_St8x4(ptr nocapture noundef readonly by
5252
; CHECK-LABEL: callee_St8x4(
5353
; CHECK: // @callee_St8x4
5454
; CHECK-NEXT: {
55-
; CHECK-NEXT: .reg .b32 %r<4>;
55+
; CHECK-NEXT: .reg .b32 %r<2>;
5656
; CHECK-NEXT: .reg .b64 %rd<5>;
5757
; CHECK-EMPTY:
5858
; CHECK-NEXT: // %bb.0:

0 commit comments

Comments
 (0)