[llvm] [NVPTX] Add reverseBranchCondition and CBranchOther (PR #191887)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 13 14:16:11 PDT 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: modiking
<details>
<summary>Changes</summary>
Add `CBranchOther` instruction for inverted predicate branches (`@!p bra`) and implement `reverseBranchCondition` to support branch condition inversion. Update `analyzeBranch`, `insertBranch`, and `removeBranch` to handle both `CBranch` and `CBranchOther`.
This enables passes like branch folding to properly reverse branch conditions, and is a prerequisite for SETP predicate inversion CSE.
Split from #<!-- -->185755 per reviewer feedback.
---
Full diff: https://github.com/llvm/llvm-project/pull/191887.diff
5 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp (+25-8)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.h (+2)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+2)
- (modified) llvm/test/CodeGen/NVPTX/i128.ll (+16-20)
- (modified) llvm/test/CodeGen/NVPTX/jump-table.ll (+24-31)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
index db2d96f5ff532..3516106a918d8 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
@@ -99,10 +99,12 @@ bool NVPTXInstrInfo::analyzeBranch(MachineBasicBlock &MBB,
if (LastInst.getOpcode() == NVPTX::GOTO) {
TBB = LastInst.getOperand(0).getMBB();
return false;
- } else if (LastInst.getOpcode() == NVPTX::CBranch) {
+ } else if (LastInst.getOpcode() == NVPTX::CBranch ||
+ LastInst.getOpcode() == NVPTX::CBranchOther) {
// Block ends with fall-through condbranch.
TBB = LastInst.getOperand(1).getMBB();
Cond.push_back(LastInst.getOperand(0));
+ Cond.push_back(MachineOperand::CreateImm(LastInst.getOpcode()));
return false;
}
// Otherwise, don't know what this is.
@@ -116,11 +118,13 @@ bool NVPTXInstrInfo::analyzeBranch(MachineBasicBlock &MBB,
if (I != MBB.begin() && isUnpredicatedTerminator(*--I))
return true;
- // If the block ends with NVPTX::GOTO and NVPTX:CBranch, handle it.
- if (SecondLastInst.getOpcode() == NVPTX::CBranch &&
+ // If the block ends with a conditional branch and NVPTX::GOTO, handle it.
+ if ((SecondLastInst.getOpcode() == NVPTX::CBranch ||
+ SecondLastInst.getOpcode() == NVPTX::CBranchOther) &&
LastInst.getOpcode() == NVPTX::GOTO) {
TBB = SecondLastInst.getOperand(1).getMBB();
Cond.push_back(SecondLastInst.getOperand(0));
+ Cond.push_back(MachineOperand::CreateImm(SecondLastInst.getOpcode()));
FBB = LastInst.getOperand(0).getMBB();
return false;
}
@@ -147,7 +151,8 @@ unsigned NVPTXInstrInfo::removeBranch(MachineBasicBlock &MBB,
if (I == MBB.begin())
return 0;
--I;
- if (I->getOpcode() != NVPTX::GOTO && I->getOpcode() != NVPTX::CBranch)
+ if (I->getOpcode() != NVPTX::GOTO && I->getOpcode() != NVPTX::CBranch &&
+ I->getOpcode() != NVPTX::CBranchOther)
return 0;
// Remove the branch.
@@ -158,7 +163,7 @@ unsigned NVPTXInstrInfo::removeBranch(MachineBasicBlock &MBB,
if (I == MBB.begin())
return 1;
--I;
- if (I->getOpcode() != NVPTX::CBranch)
+ if (I->getOpcode() != NVPTX::CBranch && I->getOpcode() != NVPTX::CBranchOther)
return 1;
// Remove the branch.
@@ -176,7 +181,7 @@ unsigned NVPTXInstrInfo::insertBranch(MachineBasicBlock &MBB,
// Shouldn't be a fall through.
assert(TBB && "insertBranch must not be told to insert a fallthrough");
- assert((Cond.size() == 1 || Cond.size() == 0) &&
+ assert((Cond.size() == 2 || Cond.size() == 0) &&
"NVPTX branch conditions have two components!");
// One-way branch.
@@ -184,12 +189,24 @@ unsigned NVPTXInstrInfo::insertBranch(MachineBasicBlock &MBB,
if (Cond.empty()) // Unconditional branch
BuildMI(&MBB, DL, get(NVPTX::GOTO)).addMBB(TBB);
else // Conditional branch
- BuildMI(&MBB, DL, get(NVPTX::CBranch)).add(Cond[0]).addMBB(TBB);
+ BuildMI(&MBB, DL, get(Cond[1].getImm())).add(Cond[0]).addMBB(TBB);
return 1;
}
// Two-way Conditional Branch.
- BuildMI(&MBB, DL, get(NVPTX::CBranch)).add(Cond[0]).addMBB(TBB);
+ BuildMI(&MBB, DL, get(Cond[1].getImm())).add(Cond[0]).addMBB(TBB);
BuildMI(&MBB, DL, get(NVPTX::GOTO)).addMBB(FBB);
return 2;
}
+
+bool NVPTXInstrInfo::reverseBranchCondition(
+ SmallVectorImpl<MachineOperand> &Cond) const {
+ assert(Cond.size() == 2 && "Invalid NVPTX branch condition!");
+ if (Cond[1].getImm() == NVPTX::CBranch)
+ Cond[1].setImm(NVPTX::CBranchOther);
+ else if (Cond[1].getImm() == NVPTX::CBranchOther)
+ Cond[1].setImm(NVPTX::CBranch);
+ else
+ return true;
+ return false;
+}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
index 3c150b1e26785..c0b7e77cf63a5 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
@@ -67,6 +67,8 @@ class NVPTXInstrInfo : public NVPTXGenInstrInfo {
MachineBasicBlock *FBB, ArrayRef<MachineOperand> Cond,
const DebugLoc &DL,
int *BytesAdded = nullptr) const override;
+ bool
+ reverseBranchCondition(SmallVectorImpl<MachineOperand> &Cond) const override;
};
} // namespace llvm
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 858635bc99815..d8ed1d20101a1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -2422,6 +2422,8 @@ let isTerminator=1 in {
def CBranch : NVPTXInst<(outs), (ins B1:$a, brtarget:$target),
"@$a bra \t$target;",
[(brcond i1:$a, bb:$target)]>;
+ def CBranchOther : NVPTXInst<(outs), (ins B1:$a, brtarget:$target),
+ "@!$a bra \t$target;", []>;
let isBarrier=1 in
def GOTO : BasicNVPTXInst<(outs), (ins brtarget:$target),
diff --git a/llvm/test/CodeGen/NVPTX/i128.ll b/llvm/test/CodeGen/NVPTX/i128.ll
index 8a5e0a00a20eb..df1214c3f5e00 100644
--- a/llvm/test/CodeGen/NVPTX/i128.ll
+++ b/llvm/test/CodeGen/NVPTX/i128.ll
@@ -58,7 +58,7 @@ define i128 @srem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: selp.b64 %rd77, 0, %rd2, %p13;
; CHECK-NEXT: or.pred %p15, %p13, %p14;
; CHECK-NEXT: @%p15 bra $L__BB0_5;
-; CHECK-NEXT: // %bb.3: // %udiv-bb1
+; CHECK-NEXT: // %bb.1: // %udiv-bb1
; CHECK-NEXT: add.cc.s64 %rd71, %rd27, 1;
; CHECK-NEXT: addc.cc.s64 %rd72, %rd28, 0;
; CHECK-NEXT: or.b64 %rd31, %rd71, %rd72;
@@ -77,7 +77,7 @@ define i128 @srem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: mov.b64 %rd70, 0;
; CHECK-NEXT: mov.b64 %rd69, %rd70;
; CHECK-NEXT: @%p16 bra $L__BB0_4;
-; CHECK-NEXT: // %bb.1: // %udiv-preheader
+; CHECK-NEXT: // %bb.2: // %udiv-preheader
; CHECK-NEXT: cvt.u32.u64 %r9, %rd71;
; CHECK-NEXT: shr.u64 %rd36, %rd2, %r9;
; CHECK-NEXT: sub.s32 %r10, 64, %r9;
@@ -91,7 +91,7 @@ define i128 @srem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: add.cc.s64 %rd6, %rd4, -1;
; CHECK-NEXT: addc.cc.s64 %rd7, %rd5, -1;
; CHECK-NEXT: mov.b64 %rd69, %rd70;
-; CHECK-NEXT: $L__BB0_2: // %udiv-do-while
+; CHECK-NEXT: $L__BB0_3: // %udiv-do-while
; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
; CHECK-NEXT: shr.u64 %rd40, %rd73, 63;
; CHECK-NEXT: shl.b64 %rd41, %rd74, 1;
@@ -117,8 +117,7 @@ define i128 @srem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: addc.cc.s64 %rd72, %rd72, -1;
; CHECK-NEXT: or.b64 %rd55, %rd71, %rd72;
; CHECK-NEXT: setp.eq.b64 %p19, %rd55, 0;
-; CHECK-NEXT: @%p19 bra $L__BB0_4;
-; CHECK-NEXT: bra.uni $L__BB0_2;
+; CHECK-NEXT: @!%p19 bra $L__BB0_3;
; CHECK-NEXT: $L__BB0_4: // %udiv-loop-exit
; CHECK-NEXT: shr.u64 %rd56, %rd75, 63;
; CHECK-NEXT: shl.b64 %rd57, %rd76, 1;
@@ -187,7 +186,7 @@ define i128 @urem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: selp.b64 %rd64, 0, %rd5, %p11;
; CHECK-NEXT: or.pred %p13, %p11, %p12;
; CHECK-NEXT: @%p13 bra $L__BB1_5;
-; CHECK-NEXT: // %bb.3: // %udiv-bb1
+; CHECK-NEXT: // %bb.1: // %udiv-bb1
; CHECK-NEXT: add.cc.s64 %rd58, %rd18, 1;
; CHECK-NEXT: addc.cc.s64 %rd59, %rd19, 0;
; CHECK-NEXT: or.b64 %rd22, %rd58, %rd59;
@@ -206,7 +205,7 @@ define i128 @urem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: mov.b64 %rd57, 0;
; CHECK-NEXT: mov.b64 %rd56, %rd57;
; CHECK-NEXT: @%p14 bra $L__BB1_4;
-; CHECK-NEXT: // %bb.1: // %udiv-preheader
+; CHECK-NEXT: // %bb.2: // %udiv-preheader
; CHECK-NEXT: cvt.u32.u64 %r9, %rd58;
; CHECK-NEXT: shr.u64 %rd27, %rd5, %r9;
; CHECK-NEXT: sub.s32 %r10, 64, %r9;
@@ -220,7 +219,7 @@ define i128 @urem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: add.cc.s64 %rd3, %rd1, -1;
; CHECK-NEXT: addc.cc.s64 %rd4, %rd2, -1;
; CHECK-NEXT: mov.b64 %rd56, %rd57;
-; CHECK-NEXT: $L__BB1_2: // %udiv-do-while
+; CHECK-NEXT: $L__BB1_3: // %udiv-do-while
; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
; CHECK-NEXT: shr.u64 %rd31, %rd60, 63;
; CHECK-NEXT: shl.b64 %rd32, %rd61, 1;
@@ -246,8 +245,7 @@ define i128 @urem_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: addc.cc.s64 %rd59, %rd59, -1;
; CHECK-NEXT: or.b64 %rd46, %rd58, %rd59;
; CHECK-NEXT: setp.eq.b64 %p17, %rd46, 0;
-; CHECK-NEXT: @%p17 bra $L__BB1_4;
-; CHECK-NEXT: bra.uni $L__BB1_2;
+; CHECK-NEXT: @!%p17 bra $L__BB1_3;
; CHECK-NEXT: $L__BB1_4: // %udiv-loop-exit
; CHECK-NEXT: shr.u64 %rd47, %rd62, 63;
; CHECK-NEXT: shl.b64 %rd48, %rd63, 1;
@@ -358,7 +356,7 @@ define i128 @sdiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: selp.b64 %rd72, 0, %rd1, %p13;
; CHECK-NEXT: or.pred %p15, %p13, %p14;
; CHECK-NEXT: @%p15 bra $L__BB4_5;
-; CHECK-NEXT: // %bb.3: // %udiv-bb1
+; CHECK-NEXT: // %bb.1: // %udiv-bb1
; CHECK-NEXT: add.cc.s64 %rd66, %rd28, 1;
; CHECK-NEXT: addc.cc.s64 %rd67, %rd29, 0;
; CHECK-NEXT: or.b64 %rd32, %rd66, %rd67;
@@ -377,7 +375,7 @@ define i128 @sdiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: mov.b64 %rd65, 0;
; CHECK-NEXT: mov.b64 %rd64, %rd65;
; CHECK-NEXT: @%p16 bra $L__BB4_4;
-; CHECK-NEXT: // %bb.1: // %udiv-preheader
+; CHECK-NEXT: // %bb.2: // %udiv-preheader
; CHECK-NEXT: cvt.u32.u64 %r9, %rd66;
; CHECK-NEXT: shr.u64 %rd37, %rd1, %r9;
; CHECK-NEXT: sub.s32 %r10, 64, %r9;
@@ -391,7 +389,7 @@ define i128 @sdiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: add.cc.s64 %rd6, %rd3, -1;
; CHECK-NEXT: addc.cc.s64 %rd7, %rd4, -1;
; CHECK-NEXT: mov.b64 %rd64, %rd65;
-; CHECK-NEXT: $L__BB4_2: // %udiv-do-while
+; CHECK-NEXT: $L__BB4_3: // %udiv-do-while
; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
; CHECK-NEXT: shr.u64 %rd41, %rd68, 63;
; CHECK-NEXT: shl.b64 %rd42, %rd69, 1;
@@ -417,8 +415,7 @@ define i128 @sdiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: addc.cc.s64 %rd67, %rd67, -1;
; CHECK-NEXT: or.b64 %rd56, %rd66, %rd67;
; CHECK-NEXT: setp.eq.b64 %p19, %rd56, 0;
-; CHECK-NEXT: @%p19 bra $L__BB4_4;
-; CHECK-NEXT: bra.uni $L__BB4_2;
+; CHECK-NEXT: @!%p19 bra $L__BB4_3;
; CHECK-NEXT: $L__BB4_4: // %udiv-loop-exit
; CHECK-NEXT: shr.u64 %rd57, %rd70, 63;
; CHECK-NEXT: shl.b64 %rd58, %rd71, 1;
@@ -481,7 +478,7 @@ define i128 @udiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: selp.b64 %rd58, 0, %rd3, %p11;
; CHECK-NEXT: or.pred %p13, %p11, %p12;
; CHECK-NEXT: @%p13 bra $L__BB5_5;
-; CHECK-NEXT: // %bb.3: // %udiv-bb1
+; CHECK-NEXT: // %bb.1: // %udiv-bb1
; CHECK-NEXT: add.cc.s64 %rd52, %rd18, 1;
; CHECK-NEXT: addc.cc.s64 %rd53, %rd19, 0;
; CHECK-NEXT: or.b64 %rd22, %rd52, %rd53;
@@ -500,7 +497,7 @@ define i128 @udiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: mov.b64 %rd51, 0;
; CHECK-NEXT: mov.b64 %rd50, %rd51;
; CHECK-NEXT: @%p14 bra $L__BB5_4;
-; CHECK-NEXT: // %bb.1: // %udiv-preheader
+; CHECK-NEXT: // %bb.2: // %udiv-preheader
; CHECK-NEXT: cvt.u32.u64 %r9, %rd52;
; CHECK-NEXT: shr.u64 %rd27, %rd3, %r9;
; CHECK-NEXT: sub.s32 %r10, 64, %r9;
@@ -514,7 +511,7 @@ define i128 @udiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: add.cc.s64 %rd1, %rd5, -1;
; CHECK-NEXT: addc.cc.s64 %rd2, %rd6, -1;
; CHECK-NEXT: mov.b64 %rd50, %rd51;
-; CHECK-NEXT: $L__BB5_2: // %udiv-do-while
+; CHECK-NEXT: $L__BB5_3: // %udiv-do-while
; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
; CHECK-NEXT: shr.u64 %rd31, %rd54, 63;
; CHECK-NEXT: shl.b64 %rd32, %rd55, 1;
@@ -540,8 +537,7 @@ define i128 @udiv_i128(i128 %lhs, i128 %rhs) {
; CHECK-NEXT: addc.cc.s64 %rd53, %rd53, -1;
; CHECK-NEXT: or.b64 %rd46, %rd52, %rd53;
; CHECK-NEXT: setp.eq.b64 %p17, %rd46, 0;
-; CHECK-NEXT: @%p17 bra $L__BB5_4;
-; CHECK-NEXT: bra.uni $L__BB5_2;
+; CHECK-NEXT: @!%p17 bra $L__BB5_3;
; CHECK-NEXT: $L__BB5_4: // %udiv-loop-exit
; CHECK-NEXT: shr.u64 %rd47, %rd56, 63;
; CHECK-NEXT: shl.b64 %rd48, %rd57, 1;
diff --git a/llvm/test/CodeGen/NVPTX/jump-table.ll b/llvm/test/CodeGen/NVPTX/jump-table.ll
index 4d391f85e978a..8f253f5861aea 100644
--- a/llvm/test/CodeGen/NVPTX/jump-table.ll
+++ b/llvm/test/CodeGen/NVPTX/jump-table.ll
@@ -53,9 +53,8 @@ define void @foo(i32 %i) {
; PTX50-NEXT: @%p4 bra $L__BB0_7;
; PTX50-NEXT: // %bb.2: // %entry
; PTX50-NEXT: setp.eq.b32 %p5, %r1, 1;
-; PTX50-NEXT: @%p5 bra $L__BB0_3;
-; PTX50-NEXT: bra.uni $L__BB0_9;
-; PTX50-NEXT: $L__BB0_3: // %case1
+; PTX50-NEXT: @!%p5 bra $L__BB0_9;
+; PTX50-NEXT: // %bb.3: // %case1
; PTX50-NEXT: st.global.b32 [out], 1;
; PTX50-NEXT: bra.uni $L__BB0_9;
; PTX50-NEXT: $L__BB0_4: // %entry
@@ -63,9 +62,8 @@ define void @foo(i32 %i) {
; PTX50-NEXT: @%p2 bra $L__BB0_8;
; PTX50-NEXT: // %bb.5: // %entry
; PTX50-NEXT: setp.eq.b32 %p3, %r1, 3;
-; PTX50-NEXT: @%p3 bra $L__BB0_6;
-; PTX50-NEXT: bra.uni $L__BB0_9;
-; PTX50-NEXT: $L__BB0_6: // %case3
+; PTX50-NEXT: @!%p3 bra $L__BB0_9;
+; PTX50-NEXT: // %bb.6: // %case3
; PTX50-NEXT: st.global.b32 [out], 3;
; PTX50-NEXT: bra.uni $L__BB0_9;
; PTX50-NEXT: $L__BB0_7: // %case0
@@ -116,18 +114,18 @@ define i32 @test2(i32 %tmp158) {
; PTX60-NEXT: @%p1 bra $L__BB1_4;
; PTX60-NEXT: // %bb.1: // %entry
; PTX60-NEXT: setp.lt.u32 %p4, %r1, 6;
-; PTX60-NEXT: @%p4 bra $L__BB1_3;
+; PTX60-NEXT: @%p4 bra $L__BB1_6;
; PTX60-NEXT: // %bb.2: // %entry
; PTX60-NEXT: setp.lt.s32 %p5, %r1, -2147483645;
-; PTX60-NEXT: @%p5 bra $L__BB1_3;
-; PTX60-NEXT: bra.uni $L__BB1_6;
+; PTX60-NEXT: @%p5 bra $L__BB1_6;
+; PTX60-NEXT: bra.uni $L__BB1_3;
; PTX60-NEXT: $L__BB1_4: // %entry
; PTX60-NEXT: add.s32 %r2, %r1, -120;
; PTX60-NEXT: setp.gt.u32 %p2, %r2, 5;
; PTX60-NEXT: @%p2 bra $L__BB1_5;
; PTX60-NEXT: // %bb.12: // %entry
; PTX60-NEXT: $L_brx_0: .branchtargets
-; PTX60-NEXT: $L__BB1_3,
+; PTX60-NEXT: $L__BB1_6,
; PTX60-NEXT: $L__BB1_7,
; PTX60-NEXT: $L__BB1_8,
; PTX60-NEXT: $L__BB1_9,
@@ -139,15 +137,14 @@ define i32 @test2(i32 %tmp158) {
; PTX60-NEXT: ret;
; PTX60-NEXT: $L__BB1_5: // %entry
; PTX60-NEXT: setp.eq.b32 %p3, %r1, 1024;
-; PTX60-NEXT: @%p3 bra $L__BB1_3;
-; PTX60-NEXT: bra.uni $L__BB1_6;
-; PTX60-NEXT: $L__BB1_3: // %bb338
+; PTX60-NEXT: @!%p3 bra $L__BB1_3;
+; PTX60-NEXT: $L__BB1_6: // %bb338
; PTX60-NEXT: st.param.b32 [func_retval0], 11;
; PTX60-NEXT: ret;
; PTX60-NEXT: $L__BB1_10: // %bb342
; PTX60-NEXT: st.param.b32 [func_retval0], 15;
; PTX60-NEXT: ret;
-; PTX60-NEXT: $L__BB1_6: // %bb336
+; PTX60-NEXT: $L__BB1_3: // %bb336
; PTX60-NEXT: st.param.b32 [func_retval0], 10;
; PTX60-NEXT: ret;
; PTX60-NEXT: $L__BB1_8: // %bb340
@@ -171,16 +168,15 @@ define i32 @test2(i32 %tmp158) {
; PTX50-NEXT: @%p1 bra $L__BB1_4;
; PTX50-NEXT: // %bb.1: // %entry
; PTX50-NEXT: setp.lt.u32 %p11, %r1, 6;
-; PTX50-NEXT: @%p11 bra $L__BB1_3;
+; PTX50-NEXT: @%p11 bra $L__BB1_15;
; PTX50-NEXT: // %bb.2: // %entry
; PTX50-NEXT: setp.lt.s32 %p12, %r1, -2147483645;
-; PTX50-NEXT: @%p12 bra $L__BB1_3;
-; PTX50-NEXT: bra.uni $L__BB1_15;
+; PTX50-NEXT: @%p12 bra $L__BB1_15;
+; PTX50-NEXT: bra.uni $L__BB1_3;
; PTX50-NEXT: $L__BB1_4: // %entry
; PTX50-NEXT: setp.gt.s32 %p2, %r1, 122;
-; PTX50-NEXT: @%p2 bra $L__BB1_9;
-; PTX50-NEXT: bra.uni $L__BB1_5;
-; PTX50-NEXT: $L__BB1_9: // %entry
+; PTX50-NEXT: @!%p2 bra $L__BB1_5;
+; PTX50-NEXT: // %bb.9: // %entry
; PTX50-NEXT: setp.gt.s32 %p3, %r1, 124;
; PTX50-NEXT: @%p3 bra $L__BB1_13;
; PTX50-NEXT: // %bb.10: // %entry
@@ -188,22 +184,20 @@ define i32 @test2(i32 %tmp158) {
; PTX50-NEXT: @%p6 bra $L__BB1_17;
; PTX50-NEXT: // %bb.11: // %entry
; PTX50-NEXT: setp.eq.b32 %p7, %r1, 124;
-; PTX50-NEXT: @%p7 bra $L__BB1_12;
-; PTX50-NEXT: bra.uni $L__BB1_15;
-; PTX50-NEXT: $L__BB1_12: // %bb342
+; PTX50-NEXT: @!%p7 bra $L__BB1_3;
+; PTX50-NEXT: // %bb.12: // %bb342
; PTX50-NEXT: st.param.b32 [func_retval0], 15;
; PTX50-NEXT: ret;
; PTX50-NEXT: $L__BB1_5: // %entry
; PTX50-NEXT: setp.eq.b32 %p8, %r1, 120;
-; PTX50-NEXT: @%p8 bra $L__BB1_3;
+; PTX50-NEXT: @%p8 bra $L__BB1_15;
; PTX50-NEXT: // %bb.6: // %entry
; PTX50-NEXT: setp.eq.b32 %p9, %r1, 121;
; PTX50-NEXT: @%p9 bra $L__BB1_16;
; PTX50-NEXT: // %bb.7: // %entry
; PTX50-NEXT: setp.eq.b32 %p10, %r1, 122;
-; PTX50-NEXT: @%p10 bra $L__BB1_8;
-; PTX50-NEXT: bra.uni $L__BB1_15;
-; PTX50-NEXT: $L__BB1_8: // %bb340
+; PTX50-NEXT: @!%p10 bra $L__BB1_3;
+; PTX50-NEXT: // %bb.8: // %bb340
; PTX50-NEXT: st.param.b32 [func_retval0], 13;
; PTX50-NEXT: ret;
; PTX50-NEXT: $L__BB1_13: // %entry
@@ -211,9 +205,8 @@ define i32 @test2(i32 %tmp158) {
; PTX50-NEXT: @%p4 bra $L__BB1_18;
; PTX50-NEXT: // %bb.14: // %entry
; PTX50-NEXT: setp.eq.b32 %p5, %r1, 1024;
-; PTX50-NEXT: @%p5 bra $L__BB1_3;
-; PTX50-NEXT: bra.uni $L__BB1_15;
-; PTX50-NEXT: $L__BB1_3: // %bb338
+; PTX50-NEXT: @!%p5 bra $L__BB1_3;
+; PTX50-NEXT: $L__BB1_15: // %bb338
; PTX50-NEXT: st.param.b32 [func_retval0], 11;
; PTX50-NEXT: ret;
; PTX50-NEXT: $L__BB1_17: // %bb341
@@ -222,7 +215,7 @@ define i32 @test2(i32 %tmp158) {
; PTX50-NEXT: $L__BB1_18: // %bb343
; PTX50-NEXT: st.param.b32 [func_retval0], 18;
; PTX50-NEXT: ret;
-; PTX50-NEXT: $L__BB1_15: // %bb336
+; PTX50-NEXT: $L__BB1_3: // %bb336
; PTX50-NEXT: st.param.b32 [func_retval0], 10;
; PTX50-NEXT: ret;
; PTX50-NEXT: $L__BB1_16: // %bb339
``````````
</details>
https://github.com/llvm/llvm-project/pull/191887
More information about the llvm-commits
mailing list