[llvm] [NVPTX] Add commutativity to SETP instructions to enable MachineCSE of inverted predicates (PR #191888)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 13 14:16:39 PDT 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: modiking
<details>
<summary>Changes</summary>
Inverted predicates can be used freely in PTX. If we can invert a predicate and CSE the generating instruction we can save calculating the inverse.
Teach the NVPTX `commuteInstructionImpl` that SETP instructions can be inverted to allow CSEing with previous SETP that match the inverted form. This also inverts the branch users of the predicate to maintain correctness.
Currently only allow the SETP inversion if all users are branches. Future work can extend this to `sel` and `not` instructions.
**Depends on #<!-- -->191887** (the first commit in this PR is from that stacked PR and should be reviewed there).
Split from #<!-- -->185755 per reviewer feedback.
---
Patch is 194.23 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/191888.diff
21 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp (+205-8)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.h (+10)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+4-2)
- (modified) llvm/test/CodeGen/NVPTX/i128.ll (+16-20)
- (modified) llvm/test/CodeGen/NVPTX/jump-table.ll (+24-31)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-bfloat16.ll (+695)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-float16.ll (+695)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-float32.ll (+663)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-float64.ll (+679)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-int16.ll (+437)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-int32.ll (+427)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-int64.ll (+437)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-multiple-users.ll (+50)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion-vector-float.ll (+99)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-no-inversion-bfloat16.ll (+224)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-no-inversion-float16.ll (+224)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-no-inversion-float32.ll (+214)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-no-inversion-float64.ll (+219)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-no-inversion-int16.ll (+224)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-no-inversion-int32.ll (+219)
- (added) llvm/test/CodeGen/NVPTX/machine-cse-predicate-no-inversion-int64.ll (+224)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
index db2d96f5ff532..4a9e0f44afc9c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
@@ -16,6 +16,7 @@
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/CodeGen/MachineInstrBuilder.h"
#include "llvm/CodeGen/MachineRegisterInfo.h"
+#include <optional>
using namespace llvm;
@@ -99,10 +100,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 +119,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 +152,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 +164,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 +182,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 +190,203 @@ 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;
+}
+
+bool NVPTXInstrInfo::invertPredicateBranchInstr(MachineBasicBlock &MBB) const {
+ MachineBasicBlock *TBB = nullptr, *FBB = nullptr;
+ SmallVector<MachineOperand, 4> Cond;
+ if (analyzeBranch(MBB, TBB, FBB, Cond, /*AllowModify=*/false))
+ return false;
+ if (Cond.empty())
+ return false;
+ if (reverseBranchCondition(Cond))
+ return false;
+ DebugLoc DL = MBB.findBranchDebugLoc();
+ removeBranch(MBB);
+ insertBranch(MBB, TBB, FBB, Cond, DL);
+ return true;
+}
+
+static bool isIntegerSetp(const MachineInstr &MI) {
+ switch (MI.getOpcode()) {
+ case NVPTX::SETP_i16rr:
+ case NVPTX::SETP_i16ri:
+ case NVPTX::SETP_i16ir:
+ case NVPTX::SETP_i32rr:
+ case NVPTX::SETP_i32ri:
+ case NVPTX::SETP_i32ir:
+ case NVPTX::SETP_i64rr:
+ case NVPTX::SETP_i64ri:
+ case NVPTX::SETP_i64ir:
+ return true;
+ default:
+ return false;
+ }
+}
+
+static bool isFloatSetp(const MachineInstr &MI) {
+ switch (MI.getOpcode()) {
+ // TODO: Because these generate 2 predicates ISel doesn't duplicate the
+ // computation but rather directly lowers it to `not.pred` if we need the
+ // inverse. This means for patterns like these tests we would want to
+ // speculatively invert the `not.pred` instruction rather than the comparison.
+ // case NVPTX::SETP_bf16x2rr:
+ // case NVPTX::SETP_f16x2rr:
+ case NVPTX::SETP_bf16rr:
+ case NVPTX::SETP_f16rr:
+ case NVPTX::SETP_f32rr:
+ case NVPTX::SETP_f32ri:
+ case NVPTX::SETP_f32ir:
+ case NVPTX::SETP_f64rr:
+ case NVPTX::SETP_f64ri:
+ case NVPTX::SETP_f64ir:
+ return true;
+ default:
+ return false;
+ }
+}
+
+static std::optional<int64_t> invertIntegerCmpMode(int64_t Mode) {
+ switch (Mode) {
+ case NVPTX::PTXCmpMode::EQ:
+ return NVPTX::PTXCmpMode::NE;
+ case NVPTX::PTXCmpMode::NE:
+ return NVPTX::PTXCmpMode::EQ;
+ case NVPTX::PTXCmpMode::LT:
+ return NVPTX::PTXCmpMode::GE;
+ case NVPTX::PTXCmpMode::LE:
+ return NVPTX::PTXCmpMode::GT;
+ case NVPTX::PTXCmpMode::GT:
+ return NVPTX::PTXCmpMode::LE;
+ case NVPTX::PTXCmpMode::GE:
+ return NVPTX::PTXCmpMode::LT;
+ case NVPTX::PTXCmpMode::LTU:
+ return NVPTX::PTXCmpMode::GEU;
+ case NVPTX::PTXCmpMode::LEU:
+ return NVPTX::PTXCmpMode::GTU;
+ case NVPTX::PTXCmpMode::GTU:
+ return NVPTX::PTXCmpMode::LEU;
+ case NVPTX::PTXCmpMode::GEU:
+ return NVPTX::PTXCmpMode::LTU;
+ default:
+ return std::nullopt;
+ }
+}
+
+static std::optional<int64_t> invertFloatCmpMode(int64_t Mode) {
+ switch (Mode) {
+ case NVPTX::PTXCmpMode::EQ:
+ return NVPTX::PTXCmpMode::NEU;
+ case NVPTX::PTXCmpMode::NE:
+ return NVPTX::PTXCmpMode::EQU;
+ case NVPTX::PTXCmpMode::EQU:
+ return NVPTX::PTXCmpMode::NE;
+ case NVPTX::PTXCmpMode::NEU:
+ return NVPTX::PTXCmpMode::EQ;
+ case NVPTX::PTXCmpMode::LT:
+ return NVPTX::PTXCmpMode::GEU;
+ case NVPTX::PTXCmpMode::LE:
+ return NVPTX::PTXCmpMode::GTU;
+ case NVPTX::PTXCmpMode::GT:
+ return NVPTX::PTXCmpMode::LEU;
+ case NVPTX::PTXCmpMode::GE:
+ return NVPTX::PTXCmpMode::LTU;
+ case NVPTX::PTXCmpMode::LTU:
+ return NVPTX::PTXCmpMode::GE;
+ case NVPTX::PTXCmpMode::LEU:
+ return NVPTX::PTXCmpMode::GT;
+ case NVPTX::PTXCmpMode::GTU:
+ return NVPTX::PTXCmpMode::LE;
+ case NVPTX::PTXCmpMode::GEU:
+ return NVPTX::PTXCmpMode::LT;
+ case NVPTX::PTXCmpMode::NUM:
+ return NVPTX::PTXCmpMode::NotANumber;
+ case NVPTX::PTXCmpMode::NotANumber:
+ return NVPTX::PTXCmpMode::NUM;
+ default:
+ return std::nullopt;
+ }
+}
+
+static bool invertCompareInstr(MachineInstr &MI) {
+ MachineOperand &ModeOp = MI.getOperand(3);
+ assert(ModeOp.isImm() && "SETP mode operand must be an immediate");
+
+ std::optional<int64_t> Inverted;
+ if (isIntegerSetp(MI))
+ Inverted = invertIntegerCmpMode(ModeOp.getImm());
+ else if (isFloatSetp(MI))
+ Inverted = invertFloatCmpMode(ModeOp.getImm());
+
+ if (!Inverted)
+ return false;
+
+ ModeOp.setImm(*Inverted);
+ return true;
+}
+
+bool NVPTXInstrInfo::findCommutedOpIndices(const MachineInstr &MI,
+ unsigned &SrcOpIdx1,
+ unsigned &SrcOpIdx2) const {
+ if (isIntegerSetp(MI) || isFloatSetp(MI))
+ return fixCommutedOpIndices(SrcOpIdx1, SrcOpIdx2, 1, 2);
+ return TargetInstrInfo::findCommutedOpIndices(MI, SrcOpIdx1, SrcOpIdx2);
+}
+
+MachineInstr *NVPTXInstrInfo::commuteInstructionImpl(MachineInstr &MI,
+ bool NewMI,
+ unsigned OpIdx1,
+ unsigned OpIdx2) const {
+ assert(!NewMI && "this should never be used");
+
+ if (!isIntegerSetp(MI) && !isFloatSetp(MI))
+ return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
+
+ if (!invertCompareInstr(MI))
+ return nullptr;
+
+ // For now all users must be invertible conditional branches.
+ // TODO: support other users such as selects.
+ bool AllInverted = true;
+ MachineRegisterInfo &MRI = MI.getParent()->getParent()->getRegInfo();
+ for (MachineInstr &UseMI :
+ MRI.use_nodbg_instructions(MI.getOperand(0).getReg())) {
+ if (!(UseMI.isConditionalBranch() &&
+ invertPredicateBranchInstr(*UseMI.getParent()))) {
+ AllInverted = false;
+ break;
+ }
+ }
+
+ if (!AllInverted) {
+ for (MachineInstr &UseMI :
+ MRI.use_nodbg_instructions(MI.getOperand(0).getReg())) {
+ if (!(UseMI.isConditionalBranch() &&
+ invertPredicateBranchInstr(*UseMI.getParent())))
+ break;
+ }
+ invertCompareInstr(MI);
+ return nullptr;
+ }
+ return &MI;
+}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
index 3c150b1e26785..957fcf5cae69a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.h
@@ -67,6 +67,16 @@ class NVPTXInstrInfo : public NVPTXGenInstrInfo {
MachineBasicBlock *FBB, ArrayRef<MachineOperand> Cond,
const DebugLoc &DL,
int *BytesAdded = nullptr) const override;
+ bool
+ reverseBranchCondition(SmallVectorImpl<MachineOperand> &Cond) const override;
+ bool findCommutedOpIndices(const MachineInstr &MI, unsigned &SrcOpIdx1,
+ unsigned &SrcOpIdx2) const override;
+ MachineInstr *commuteInstructionImpl(MachineInstr &MI, bool NewMI,
+ unsigned OpIdx1,
+ unsigned OpIdx2) const override;
+
+private:
+ bool invertPredicateBranchInstr(MachineBasicBlock &MBB) const;
};
} // namespace llvm
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 858635bc99815..13672097b072f 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1568,7 +1568,7 @@ multiclass FSETP<RegTyInfo t, bit allow_ftz = true> {
defvar ftz_str = !if(allow_ftz, "$ftz", "");
defvar op_str = "setp.${cmp:FCmp}" # ftz_str # "." # t.PtxType;
defvar flags = !con((ins CmpMode:$cmp), !if(allow_ftz, (ins FTZFlag:$ftz), (ins)));
- let hasSideEffects = false in {
+ let hasSideEffects = false, isCompare = 1, isCommutable = 1 in {
def rr :
BasicFlagsNVPTXInst<(outs B1:$dst), (ins t.RC:$a, t.RC:$b),
flags, op_str>;
@@ -1594,7 +1594,7 @@ multiclass FSETP<RegTyInfo t, bit allow_ftz = true> {
multiclass ISETP<RegTyInfo t> {
defvar op_str = "setp.${cmp:ICmp}.${cmp:IType}" # t.Size;
- let hasSideEffects = false in {
+ let hasSideEffects = false, isCompare = 1, isCommutable = 1 in {
def rr :
BasicFlagsNVPTXInst<(outs B1:$dst), (ins t.RC:$a, t.RC:$b),
(ins CmpMode:$cmp), op_str>;
@@ -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: // ...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/191888
More information about the llvm-commits
mailing list