[llvm] [NVPTX] Add reverseBranchCondition and CBranchOther (PR #191887)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 13 14:15:36 PDT 2026
https://github.com/modiking created https://github.com/llvm/llvm-project/pull/191887
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.
>From d3114dce53a041eb65b3910ad84eb93daab0d9c9 Mon Sep 17 00:00:00 2001
From: root <mmo at nvidia.com>
Date: Mon, 13 Apr 2026 14:03:18 -0700
Subject: [PATCH] [NVPTX] Add reverseBranchCondition and CBranchOther
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.
Made-with: Cursor
---
llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp | 33 ++++++++++----
llvm/lib/Target/NVPTX/NVPTXInstrInfo.h | 2 +
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 2 +
llvm/test/CodeGen/NVPTX/i128.ll | 36 +++++++---------
llvm/test/CodeGen/NVPTX/jump-table.ll | 55 +++++++++++-------------
5 files changed, 69 insertions(+), 59 deletions(-)
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
More information about the llvm-commits
mailing list