Skip to content

Commit cb43464

Browse files
ChengjunpChengjunp
committed
Rename variables in tests & Use dag helper to build nodes
1 parent c4bc027 commit cb43464

File tree

4 files changed

+41
-57
lines changed

4 files changed

+41
-57
lines changed

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3111,14 +3111,14 @@ SDValue NVPTXTargetLowering::LowerCopyToReg_128(SDValue Op,
31113111
SDNode *Node = Op.getNode();
31123112
SDLoc DL(Node);
31133113

3114-
SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::v2i64, Op->getOperand(2));
3114+
SDValue Cast = DAG.getBitcast(MVT::v2i64, Op->getOperand(2));
31153115
SDValue Lo = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
31163116
DAG.getIntPtrConstant(0, DL));
31173117
SDValue Hi = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
31183118
DAG.getIntPtrConstant(1, DL));
31193119

31203120
SmallVector<SDValue, 5> NewOps(Op->getNumOperands() + 1);
3121-
SmallVector<EVT, 3> ResultsType(Node->value_begin(), Node->value_end());
3121+
SmallVector<EVT, 3> ResultsType(Node->values());
31223122

31233123
NewOps[0] = Op->getOperand(0); // Chain
31243124
NewOps[1] = Op->getOperand(1); // Dst Reg

‎llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll‎

Lines changed: 4 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -83,10 +83,10 @@ define void @test_b128_input_from_select(ptr nocapture readonly %flag) {
8383
; CHECK-NEXT: ret;
8484

8585
%1 = addrspacecast ptr %flag to ptr addrspace(1)
86-
%tmp1 = load i8, ptr addrspace(1) %1, align 1
87-
%tobool.not = icmp eq i8 %tmp1, 0
88-
%. = select i1 %tobool.not, i128 24, i128 42
89-
tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %.)
86+
%2 = load i8, ptr addrspace(1) %1, align 1
87+
%3 = icmp eq i8 %2, 0
88+
%4 = select i1 %3, i128 24, i128 42
89+
tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %4)
9090
ret void
9191
}
9292

@@ -146,8 +146,3 @@ define void @test_use_of_b128_output(ptr nocapture readonly %data) {
146146
store <2 x i64> %5, ptr addrspace(1) @value, align 16
147147
ret void
148148
}
149-
150-
!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
151-
152-
!0 = !{i32 2, i32 0, i32 3, i32 1}
153-
!1 = !{i32 2, i32 0}

‎llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll‎

Lines changed: 15 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -103,26 +103,20 @@ define void @test_corner_values() {
103103
; CHECK-NEXT: // end inline asm
104104
; CHECK-NEXT: ret;
105105

106-
%tmp = load ptr, ptr addrspace(1) @v64, align 8
107-
%add.ptr2 = getelementptr inbounds i64, ptr %tmp, i64 1
108-
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -1, ptr %tmp, ptr nonnull %add.ptr2, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_max to ptr))
109-
%tmp3 = load ptr, ptr addrspace(1) @v64, align 8
110-
%add.ptr4 = getelementptr inbounds i64, ptr %tmp3, i64 2
111-
%add.ptr6 = getelementptr inbounds i64, ptr %tmp3, i64 3
112-
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 170141183460469231731687303715884105727, ptr nonnull %add.ptr4, ptr nonnull %add.ptr6, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_max to ptr))
113-
%tmp7 = load ptr, ptr addrspace(1) @v64, align 8
114-
%add.ptr8 = getelementptr inbounds i64, ptr %tmp7, i64 4
115-
%add.ptr10 = getelementptr inbounds i64, ptr %tmp7, i64 5
116-
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -170141183460469231731687303715884105728, ptr nonnull %add.ptr8, ptr nonnull %add.ptr10, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_min to ptr))
117-
%tmp11 = load ptr, ptr addrspace(1) @v64, align 8
118-
%add.ptr12 = getelementptr inbounds i64, ptr %tmp11, i64 6
119-
%add.ptr14 = getelementptr inbounds i64, ptr %tmp11, i64 7
120-
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 0, ptr nonnull %add.ptr12, ptr nonnull %add.ptr14, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_zero to ptr))
106+
%1 = load ptr, ptr addrspace(1) @v64, align 8
107+
%2 = getelementptr inbounds i64, ptr %1, i64 1
108+
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -1, ptr %1, ptr nonnull %2, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_max to ptr))
109+
%3 = load ptr, ptr addrspace(1) @v64, align 8
110+
%4 = getelementptr inbounds i64, ptr %3, i64 2
111+
%5 = getelementptr inbounds i64, ptr %3, i64 3
112+
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 170141183460469231731687303715884105727, ptr nonnull %4, ptr nonnull %5, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_max to ptr))
113+
%6 = load ptr, ptr addrspace(1) @v64, align 8
114+
%7 = getelementptr inbounds i64, ptr %6, i64 4
115+
%8 = getelementptr inbounds i64, ptr %6, i64 5
116+
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -170141183460469231731687303715884105728, ptr nonnull %7, ptr nonnull %8, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_min to ptr))
117+
%9 = load ptr, ptr addrspace(1) @v64, align 8
118+
%10 = getelementptr inbounds i64, ptr %9, i64 6
119+
%11 = getelementptr inbounds i64, ptr %9, i64 7
120+
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 0, ptr nonnull %10, ptr nonnull %11, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_zero to ptr))
121121
ret void
122122
}
123-
124-
125-
!nvvmir.version = !{!2, !3, !2, !3, !3, !2, !2, !2, !3}
126-
127-
!2 = !{i32 2, i32 0, i32 3, i32 1}
128-
!3 = !{i32 2, i32 0}

‎llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll‎

Lines changed: 20 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -18,11 +18,11 @@ define void @test_b128_in_loop() {
1818
; CHECK-NEXT: ld.global.s32 %rd1, [size];
1919
; CHECK-NEXT: setp.eq.s64 %p1, %rd1, 0;
2020
; CHECK-NEXT: @%p1 bra $L__BB0_3;
21-
; CHECK-NEXT: // %bb.1: // %.lr.ph.preheader
21+
; CHECK-NEXT: // %bb.1: // %BB1
2222
; CHECK-NEXT: ld.global.u64 %rd13, [x+8];
2323
; CHECK-NEXT: ld.global.u64 %rd12, [x];
2424
; CHECK-NEXT: mov.u64 %rd14, 0;
25-
; CHECK-NEXT: $L__BB0_2: // %.lr.ph
25+
; CHECK-NEXT: $L__BB0_2: // %BB2
2626
; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
2727
; CHECK-NEXT: mov.b128 %rq1, {%rd12, %rd13};
2828
; CHECK-NEXT: // begin inline asm
@@ -40,33 +40,28 @@ define void @test_b128_in_loop() {
4040
; CHECK-NEXT: add.s64 %rd14, %rd14, 1;
4141
; CHECK-NEXT: setp.ne.s64 %p2, %rd1, %rd14;
4242
; CHECK-NEXT: @%p2 bra $L__BB0_2;
43-
; CHECK-NEXT: $L__BB0_3: // %._crit_edge
43+
; CHECK-NEXT: $L__BB0_3: // %BB3
4444
; CHECK-NEXT: ret;
4545

46-
%tmp11 = load i32, ptr addrspace(1) @size, align 4
47-
%cmp3.not = icmp eq i32 %tmp11, 0
48-
br i1 %cmp3.not, label %._crit_edge, label %.lr.ph.preheader
46+
%1 = load i32, ptr addrspace(1) @size, align 4
47+
%2 = icmp eq i32 %1, 0
48+
br i1 %2, label %BB3, label %BB1
4949

50-
.lr.ph.preheader: ; preds = %0
51-
%x.promoted5 = load i128, ptr addrspace(1) @x, align 16
52-
%umax = sext i32 %tmp11 to i64
53-
br label %.lr.ph
50+
BB1: ; preds = %0
51+
%3 = load i128, ptr addrspace(1) @x, align 16
52+
%4 = sext i32 %1 to i64
53+
br label %BB2
5454

55-
.lr.ph: ; preds = %.lr.ph, %.lr.ph.preheader
56-
%1 = phi i128 [ %2, %.lr.ph ], [ %x.promoted5, %.lr.ph.preheader ]
57-
%i.04 = phi i64 [ %inc, %.lr.ph ], [ 0, %.lr.ph.preheader ]
58-
%2 = tail call i128 asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, $1;\0A\09mov.b128 $0, {lo, hi};\0A\09}", "=q,l,0"(i64 %i.04, i128 %1)
59-
%3 = bitcast i128 %2 to <2 x i64>
60-
store <2 x i64> %3, ptr addrspace(1) @x, align 16
61-
%inc = add nuw i64 %i.04, 1
62-
%exitcond.not = icmp eq i64 %inc, %umax
63-
br i1 %exitcond.not, label %._crit_edge, label %.lr.ph
55+
BB2: ; preds = %BB2, %BB1
56+
%5 = phi i128 [ %7, %BB2 ], [ %3, %BB1 ]
57+
%6 = phi i64 [ %9, %BB2 ], [ 0, %BB1 ]
58+
%7 = tail call i128 asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, $1;\0A\09mov.b128 $0, {lo, hi};\0A\09}", "=q,l,0"(i64 %6, i128 %5)
59+
%8 = bitcast i128 %7 to <2 x i64>
60+
store <2 x i64> %8, ptr addrspace(1) @x, align 16
61+
%9 = add nuw i64 %6, 1
62+
%10 = icmp eq i64 %9, %4
63+
br i1 %10, label %BB3, label %BB2
6464

65-
._crit_edge: ; preds = %.lr.ph, %0
65+
BB3: ; preds = %BB2, %0
6666
ret void
6767
}
68-
69-
!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
70-
71-
!0 = !{i32 2, i32 0, i32 3, i32 1}
72-
!1 = !{i32 2, i32 0}

0 commit comments

Comments
 (0)