[llvm] [NVPTX] Expose CSE for shifts of logic operations (PR #213091)
Hao Ren via llvm-commits
llvm-commits at lists.llvm.org
Thu Aug 20 11:01:53 PDT 2026
https://github.com/nvidia-moomoo updated https://github.com/llvm/llvm-project/pull/213091
>From 45b94214c8e5ffccc695a6938d90cd332f8dfb55 Mon Sep 17 00:00:00 2001
From: Hao Ren <haor at nvidia.com>
Date: Thu, 6 Aug 2026 18:40:12 +0000
Subject: [PATCH 1/6] [DAGCombiner][NFC] Add pre-commit tests for
shift-of-logic CSE
Capture current NVPTX codegen for shifts of AND/OR/XOR operations that share
an input and shift amount, ahead of the combine that exposes the shared shift
for CSE. Covers no-extend, zext, and sext variants, the address-arithmetic
(getelementptr) form, and an isolated negative case.
---
llvm/test/CodeGen/NVPTX/shift-logic-cse.ll | 155 +++++++++++++++++++++
1 file changed, 155 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
diff --git a/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll b/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
new file mode 100644
index 0000000000000..17d8124afb683
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
@@ -0,0 +1,155 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s
+
+; Three logic ops (and/or/xor) share %x and the same shift amount, with no
+; extend. Commuting the shifts through the logic ops exposes a single shifted
+; %x that is CSEd, replacing three shifts with one.
+define void @mixed_logic_matches(i32 %x, ptr %p0, ptr %p1, ptr %p2) {
+; CHECK-LABEL: mixed_logic_matches(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<8>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [mixed_logic_matches_param_0];
+; CHECK-NEXT: and.b32 %r2, %r1, 5;
+; CHECK-NEXT: ld.param.b64 %rd1, [mixed_logic_matches_param_1];
+; CHECK-NEXT: shl.b32 %r3, %r2, 2;
+; CHECK-NEXT: ld.param.b64 %rd2, [mixed_logic_matches_param_2];
+; CHECK-NEXT: st.b32 [%rd1], %r3;
+; CHECK-NEXT: ld.param.b64 %rd3, [mixed_logic_matches_param_3];
+; CHECK-NEXT: shl.b32 %r4, %r1, 2;
+; CHECK-NEXT: or.b32 %r5, %r4, 28;
+; CHECK-NEXT: st.b32 [%rd2], %r5;
+; CHECK-NEXT: xor.b32 %r6, %r1, 17;
+; CHECK-NEXT: shl.b32 %r7, %r6, 2;
+; CHECK-NEXT: st.b32 [%rd3], %r7;
+; CHECK-NEXT: ret;
+ %and = and i32 %x, 5
+ %and.shl = shl i32 %and, 2
+ store i32 %and.shl, ptr %p0, align 4
+ %or = or i32 %x, 7
+ %or.shl = shl i32 %or, 2
+ store i32 %or.shl, ptr %p1, align 4
+ %xor = xor i32 %x, 17
+ %xor.shl = shl i32 %xor, 2
+ store i32 %xor.shl, ptr %p2, align 4
+ ret void
+}
+
+; Matching zext shifts: shl (zext (and X, C)), S. The shared shl (zext X), S is
+; exposed and CSEd.
+define void @zext_and_matches(i32 %x, ptr %p0, ptr %p1) {
+; CHECK-LABEL: zext_and_matches(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [zext_and_matches_param_0];
+; CHECK-NEXT: and.b32 %r2, %r1, 5;
+; CHECK-NEXT: ld.param.b64 %rd1, [zext_and_matches_param_1];
+; CHECK-NEXT: ld.param.b64 %rd2, [zext_and_matches_param_2];
+; CHECK-NEXT: mul.wide.u32 %rd3, %r2, 8;
+; CHECK-NEXT: st.b64 [%rd1], %rd3;
+; CHECK-NEXT: and.b32 %r3, %r1, 6;
+; CHECK-NEXT: mul.wide.u32 %rd4, %r3, 8;
+; CHECK-NEXT: st.b64 [%rd2], %rd4;
+; CHECK-NEXT: ret;
+ %and0 = and i32 %x, 5
+ %z0 = zext i32 %and0 to i64
+ %s0 = shl i64 %z0, 3
+ store i64 %s0, ptr %p0, align 8
+ %and1 = and i32 %x, 6
+ %z1 = zext i32 %and1 to i64
+ %s1 = shl i64 %z1, 3
+ store i64 %s1, ptr %p1, align 8
+ ret void
+}
+
+; Matching sext shifts: shl (sext (or X, C)), S.
+define void @sext_or_matches(i32 %x, ptr %p0, ptr %p1) {
+; CHECK-LABEL: sext_or_matches(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [sext_or_matches_param_0];
+; CHECK-NEXT: or.b32 %r2, %r1, 5;
+; CHECK-NEXT: ld.param.b64 %rd1, [sext_or_matches_param_1];
+; CHECK-NEXT: ld.param.b64 %rd2, [sext_or_matches_param_2];
+; CHECK-NEXT: mul.wide.s32 %rd3, %r2, 4;
+; CHECK-NEXT: st.b64 [%rd1], %rd3;
+; CHECK-NEXT: or.b32 %r3, %r1, 6;
+; CHECK-NEXT: mul.wide.s32 %rd4, %r3, 4;
+; CHECK-NEXT: st.b64 [%rd2], %rd4;
+; CHECK-NEXT: ret;
+ %or0 = or i32 %x, 5
+ %e0 = sext i32 %or0 to i64
+ %s0 = shl i64 %e0, 2
+ store i64 %s0, ptr %p0, align 8
+ %or1 = or i32 %x, 6
+ %e1 = sext i32 %or1 to i64
+ %s1 = shl i64 %e1, 2
+ store i64 %s1, ptr %p1, align 8
+ ret void
+}
+
+; The shift is implicit here: the getelementptr's element-size scaling lowers
+; to shl during SelectionDAG, so ordinary address arithmetic exercises the same
+; pattern. NVPTX does not fold the scaled index into the load, so the combine
+; still applies and the shared index computation is CSEd.
+define void @gep_index_matches(i32 %base, ptr %data, ptr %p0, ptr %p1) {
+; CHECK-LABEL: gep_index_matches(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<6>;
+; CHECK-NEXT: .reg .b64 %rd<8>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [gep_index_matches_param_0];
+; CHECK-NEXT: xor.b32 %r2, %r1, 5;
+; CHECK-NEXT: ld.param.b64 %rd1, [gep_index_matches_param_1];
+; CHECK-NEXT: ld.param.b64 %rd2, [gep_index_matches_param_2];
+; CHECK-NEXT: mul.wide.u32 %rd3, %r2, 4;
+; CHECK-NEXT: add.s64 %rd4, %rd1, %rd3;
+; CHECK-NEXT: ld.param.b64 %rd5, [gep_index_matches_param_3];
+; CHECK-NEXT: ld.b32 %r3, [%rd4];
+; CHECK-NEXT: st.b32 [%rd2], %r3;
+; CHECK-NEXT: xor.b32 %r4, %r1, 6;
+; CHECK-NEXT: mul.wide.u32 %rd6, %r4, 4;
+; CHECK-NEXT: add.s64 %rd7, %rd1, %rd6;
+; CHECK-NEXT: ld.b32 %r5, [%rd7];
+; CHECK-NEXT: st.b32 [%rd5], %r5;
+; CHECK-NEXT: ret;
+ %x0 = xor i32 %base, 5
+ %i0 = zext i32 %x0 to i64
+ %a0 = getelementptr i32, ptr %data, i64 %i0
+ %v0 = load i32, ptr %a0, align 4
+ store i32 %v0, ptr %p0, align 4
+ %x1 = xor i32 %base, 6
+ %i1 = zext i32 %x1 to i64
+ %a1 = getelementptr i32, ptr %data, i64 %i1
+ %v1 = load i32, ptr %a1, align 4
+ store i32 %v1, ptr %p1, align 4
+ ret void
+}
+
+; Negative test: a lone shl of a logic op has no matching shift to share with,
+; so it is left alone. Preserving it lets targets that support it fold the shift
+; into an addressing mode; NVPTX simply keeps the single shift.
+define i32 @isolated_no_match(i32 %x) {
+; CHECK-LABEL: isolated_no_match(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [isolated_no_match_param_0];
+; CHECK-NEXT: xor.b32 %r2, %r1, 5;
+; CHECK-NEXT: shl.b32 %r3, %r2, 2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %xor = xor i32 %x, 5
+ %shl = shl i32 %xor, 2
+ ret i32 %shl
+}
>From 7ebed00dba3514050d93f9ef0df8f4bc78887301 Mon Sep 17 00:00:00 2001
From: Hao Ren <haor at nvidia.com>
Date: Fri, 24 Jul 2026 20:18:44 +0000
Subject: [PATCH 2/6] [DAGCombiner] Expose CSE for shifts of logic operations
When multiple AND, OR, or XOR operations with constant operands share the same input, commute outer SHLs through the operations so the common shifted input can be CSEd. Handle optional integer extends atomically across all eligible siblings.
Gate the combine behind a target hook because commuting the shifts can disrupt target-specific patterns such as addressing modes. Enable it for NVPTX, where reducing N shifts to one is profitable.
---
llvm/include/llvm/CodeGen/TargetLowering.h | 9 ++
llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp | 126 ++++++++++++++++++
llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 5 +
llvm/test/CodeGen/NVPTX/shift-logic-cse.ll | 71 +++++-----
4 files changed, 173 insertions(+), 38 deletions(-)
diff --git a/llvm/include/llvm/CodeGen/TargetLowering.h b/llvm/include/llvm/CodeGen/TargetLowering.h
index 46b52c30cded6..5ee77e93c79a1 100644
--- a/llvm/include/llvm/CodeGen/TargetLowering.h
+++ b/llvm/include/llvm/CodeGen/TargetLowering.h
@@ -4662,6 +4662,15 @@ class LLVM_ABI TargetLowering : public TargetLoweringBase {
return true;
}
+ /// Return true if it is profitable to commute a constant SHL through
+ /// multiple bitwise logic operations that share an input. This can expose a
+ /// common shifted operand, but may also prevent shifts from folding into
+ /// addressing modes, so targets must opt in.
+ virtual bool shouldCommuteShiftOfLogicToExposeCSE(const SDNode *N,
+ CombineLevel Level) const {
+ return false;
+ }
+
/// GlobalISel - return true if it is profitable to move this shift by a
/// constant amount through its operand, adjusting any immediate operands as
/// necessary to preserve semantics. This transformation may not be desirable
diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
index 88fe89274b0b2..5dd7275ab8b85 100644
--- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
@@ -590,6 +590,7 @@ namespace {
EVT VT, SDValue N0, SDValue N1,
SDNodeFlags Flags = SDNodeFlags());
+ SDValue foldShiftOfLogicWithSharedInput(SDNode *N);
SDValue visitShiftByConstant(SDNode *N);
SDValue foldSelectOfConstants(SDNode *N);
@@ -10779,6 +10780,127 @@ static SDValue combineShiftOfShiftedLogic(SDNode *Shift, SelectionDAG &DAG) {
LogicOp->getFlags());
}
+/// Commute SHL with a bitwise logic operation when doing so exposes a common
+/// shifted operand. For example:
+///
+/// Before:
+/// N = shl (zext (Logic X, C)), ShiftAmount
+/// SiblingShift = shl (zext (SiblingLogic X, SiblingC)), ShiftAmount
+///
+/// After:
+/// ShiftedX = shl (zext X), ShiftAmount
+/// N = Logic ShiftedX, ShiftedC
+/// SiblingShift = SiblingLogic ShiftedX, SiblingShiftedC
+///
+/// ShiftedC = (zext C) << ShiftAmount and SiblingShiftedC =
+/// (zext SiblingC) << ShiftAmount are folded constants. This replaces two
+/// variable shifts with the single shared ShiftedX. Requiring a matching
+/// sibling avoids disrupting isolated address calculations where a shift may be
+/// folded into the addressing mode.
+SDValue DAGCombiner::foldShiftOfLogicWithSharedInput(SDNode *N) {
+ struct ShiftOfLogic {
+ SDNode *Shift;
+ SDValue Logic;
+ SDValue X;
+ SDValue Constant;
+ unsigned ExtendOpcode;
+ };
+
+ // Match a logic operation, with an optional extension, inside a SHL.
+ auto matchShiftOfLogic = [&](SDNode *Shift, ShiftOfLogic &Match) {
+ if (Shift->getOpcode() != ISD::SHL || !Shift->getOperand(0).hasOneUse())
+ return false;
+
+ SDValue ShiftInput = Shift->getOperand(0);
+ unsigned ExtendOpcode = 0;
+ SDValue Logic = ShiftInput;
+ if (ISD::isExtOpcode(ShiftInput.getOpcode())) {
+ ExtendOpcode = ShiftInput.getOpcode();
+ Logic = ShiftInput.getOperand(0);
+ }
+
+ SDValue X, C;
+ if (!sd_match(Logic, m_OneUse(m_BitwiseLogic(m_Value(X),
+ m_Value(C, m_ConstInt())))))
+ return false;
+
+ Match = {Shift, Logic, X, C, ExtendOpcode};
+ return true;
+ };
+
+ // Match N as the root shift-of-logic; bail if it does not fit the pattern or
+ // should not be commuted.
+ ShiftOfLogic Root;
+ if (!matchShiftOfLogic(N, Root) ||
+ !TLI.shouldCommuteShiftOfLogicToExposeCSE(N, Level))
+ return SDValue();
+
+ EVT VT = N->getValueType(0);
+ SDValue ShiftAmount = N->getOperand(1);
+
+ // Collect candidate shifts that share X. Reached through another user of X,
+ // the logic result feeds the shift directly or through an optional extend.
+ SmallVector<SDNode *, 4> CandidateShifts;
+ for (SDNode *CandidateLogic : Root.X->users()) {
+ if (CandidateLogic == Root.Logic.getNode())
+ continue;
+ for (SDNode *LogicUser : CandidateLogic->users()) {
+ if (ISD::isExtOpcode(LogicUser->getOpcode())) {
+ // shl (ext (logic X, C)): step through the extend to find the shift.
+ for (SDNode *ExtendUser : LogicUser->users())
+ if (ExtendUser->getOpcode() == ISD::SHL)
+ CandidateShifts.push_back(ExtendUser);
+ } else if (LogicUser->getOpcode() == ISD::SHL) {
+ // shl (logic X, C): the user is already the shift.
+ CandidateShifts.push_back(LogicUser);
+ }
+ }
+ }
+
+ // Verify each candidate against the root's pattern; those that match (same X,
+ // extend, type, and shift amount) become siblings.
+ SmallVector<ShiftOfLogic, 4> Siblings;
+ for (SDNode *CandidateShift : CandidateShifts) {
+ ShiftOfLogic Candidate;
+ if (!matchShiftOfLogic(CandidateShift, Candidate) ||
+ Candidate.X != Root.X || Candidate.ExtendOpcode != Root.ExtendOpcode ||
+ CandidateShift->getValueType(0) != VT ||
+ CandidateShift->getOperand(1) != ShiftAmount ||
+ !TLI.isDesirableToCommuteWithShift(CandidateShift, Level))
+ continue;
+ Siblings.push_back(Candidate);
+ }
+ if (Siblings.empty())
+ return SDValue();
+
+ // Build the shared shifted X once, then rewrite the root and every sibling
+ // into a logic op over it so the shift is CSE'd.
+ SDValue ShiftedX = DAG.getNode(
+ ISD::SHL, SDLoc(N), VT,
+ Root.ExtendOpcode ? DAG.getNode(Root.ExtendOpcode, SDLoc(N), VT, Root.X)
+ : Root.X,
+ ShiftAmount);
+
+ // Rebuild the logic op from shared ShiftedX and a folded constant shift.
+ auto rebuild = [&](SDValue OldLogic, SDValue OldC, const SDLoc &NodeDL) {
+ SDValue NewC = OldC;
+ if (Root.ExtendOpcode)
+ NewC = DAG.getNode(Root.ExtendOpcode, NodeDL, VT, OldC);
+ SDValue ShiftedC = DAG.getNode(ISD::SHL, NodeDL, VT, NewC, ShiftAmount);
+ return DAG.getNode(OldLogic.getOpcode(), NodeDL, VT, ShiftedX, ShiftedC,
+ OldLogic->getFlags());
+ };
+
+ for (const ShiftOfLogic &Sibling : Siblings) {
+ SDValue NewLogic =
+ rebuild(Sibling.Logic, Sibling.Constant, SDLoc(Sibling.Shift));
+ DAG.ReplaceAllUsesOfValueWith(SDValue(Sibling.Shift, 0), NewLogic);
+ deleteAndRecombine(Sibling.Shift);
+ AddToWorklist(NewLogic.getNode());
+ }
+ return rebuild(Root.Logic, Root.Constant, SDLoc(N));
+}
+
/// Handle transforms common to the three shifts, when the shift amount is a
/// constant.
/// We are looking for: (shift being one of shl/sra/srl)
@@ -10801,6 +10923,10 @@ SDValue DAGCombiner::visitShiftByConstant(SDNode *N) {
if (SDValue R = combineShiftOfShiftedLogic(N, DAG))
return R;
+ // Fold sibling shifts of logic operations to expose a shared shifted input.
+ if (SDValue R = foldShiftOfLogicWithSharedInput(N))
+ return R;
+
// We want to pull some binops through shifts, so that we have (and (shift))
// instead of (shift (and)), likewise for add, or, xor, etc. This sort of
// thing happens with address calculations, so it's important to canonicalize
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 42741ba98b438..92ff7415dd6f5 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -181,6 +181,11 @@ class NVPTXTargetLowering : public TargetLowering {
TargetLoweringOpt &TLO,
unsigned Depth = 0) const override;
+ bool shouldCommuteShiftOfLogicToExposeCSE(const SDNode *,
+ CombineLevel) const override {
+ return true;
+ }
+
private:
const NVPTXSubtarget &STI; // cache the subtarget here
mutable unsigned GlobalUniqueCallSite;
diff --git a/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll b/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
index 17d8124afb683..0aa689c2ea2ed 100644
--- a/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
+++ b/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
@@ -7,23 +7,21 @@
define void @mixed_logic_matches(i32 %x, ptr %p0, ptr %p1, ptr %p2) {
; CHECK-LABEL: mixed_logic_matches(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<8>;
+; CHECK-NEXT: .reg .b32 %r<6>;
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [mixed_logic_matches_param_0];
-; CHECK-NEXT: and.b32 %r2, %r1, 5;
; CHECK-NEXT: ld.param.b64 %rd1, [mixed_logic_matches_param_1];
-; CHECK-NEXT: shl.b32 %r3, %r2, 2;
+; CHECK-NEXT: shl.b32 %r2, %r1, 2;
+; CHECK-NEXT: and.b32 %r3, %r2, 20;
; CHECK-NEXT: ld.param.b64 %rd2, [mixed_logic_matches_param_2];
; CHECK-NEXT: st.b32 [%rd1], %r3;
; CHECK-NEXT: ld.param.b64 %rd3, [mixed_logic_matches_param_3];
-; CHECK-NEXT: shl.b32 %r4, %r1, 2;
-; CHECK-NEXT: or.b32 %r5, %r4, 28;
-; CHECK-NEXT: st.b32 [%rd2], %r5;
-; CHECK-NEXT: xor.b32 %r6, %r1, 17;
-; CHECK-NEXT: shl.b32 %r7, %r6, 2;
-; CHECK-NEXT: st.b32 [%rd3], %r7;
+; CHECK-NEXT: or.b32 %r4, %r2, 28;
+; CHECK-NEXT: st.b32 [%rd2], %r4;
+; CHECK-NEXT: xor.b32 %r5, %r2, 68;
+; CHECK-NEXT: st.b32 [%rd3], %r5;
; CHECK-NEXT: ret;
%and = and i32 %x, 5
%and.shl = shl i32 %and, 2
@@ -42,19 +40,18 @@ define void @mixed_logic_matches(i32 %x, ptr %p0, ptr %p1, ptr %p2) {
define void @zext_and_matches(i32 %x, ptr %p0, ptr %p1) {
; CHECK-LABEL: zext_and_matches(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [zext_and_matches_param_0];
-; CHECK-NEXT: and.b32 %r2, %r1, 5;
; CHECK-NEXT: ld.param.b64 %rd1, [zext_and_matches_param_1];
; CHECK-NEXT: ld.param.b64 %rd2, [zext_and_matches_param_2];
-; CHECK-NEXT: mul.wide.u32 %rd3, %r2, 8;
-; CHECK-NEXT: st.b64 [%rd1], %rd3;
-; CHECK-NEXT: and.b32 %r3, %r1, 6;
-; CHECK-NEXT: mul.wide.u32 %rd4, %r3, 8;
-; CHECK-NEXT: st.b64 [%rd2], %rd4;
+; CHECK-NEXT: mul.wide.u32 %rd3, %r1, 8;
+; CHECK-NEXT: and.b64 %rd4, %rd3, 40;
+; CHECK-NEXT: st.b64 [%rd1], %rd4;
+; CHECK-NEXT: and.b64 %rd5, %rd3, 48;
+; CHECK-NEXT: st.b64 [%rd2], %rd5;
; CHECK-NEXT: ret;
%and0 = and i32 %x, 5
%z0 = zext i32 %and0 to i64
@@ -71,19 +68,18 @@ define void @zext_and_matches(i32 %x, ptr %p0, ptr %p1) {
define void @sext_or_matches(i32 %x, ptr %p0, ptr %p1) {
; CHECK-LABEL: sext_or_matches(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [sext_or_matches_param_0];
-; CHECK-NEXT: or.b32 %r2, %r1, 5;
; CHECK-NEXT: ld.param.b64 %rd1, [sext_or_matches_param_1];
; CHECK-NEXT: ld.param.b64 %rd2, [sext_or_matches_param_2];
-; CHECK-NEXT: mul.wide.s32 %rd3, %r2, 4;
-; CHECK-NEXT: st.b64 [%rd1], %rd3;
-; CHECK-NEXT: or.b32 %r3, %r1, 6;
-; CHECK-NEXT: mul.wide.s32 %rd4, %r3, 4;
-; CHECK-NEXT: st.b64 [%rd2], %rd4;
+; CHECK-NEXT: mul.wide.s32 %rd3, %r1, 4;
+; CHECK-NEXT: or.b64 %rd4, %rd3, 20;
+; CHECK-NEXT: st.b64 [%rd1], %rd4;
+; CHECK-NEXT: or.b64 %rd5, %rd3, 24;
+; CHECK-NEXT: st.b64 [%rd2], %rd5;
; CHECK-NEXT: ret;
%or0 = or i32 %x, 5
%e0 = sext i32 %or0 to i64
@@ -103,24 +99,23 @@ define void @sext_or_matches(i32 %x, ptr %p0, ptr %p1) {
define void @gep_index_matches(i32 %base, ptr %data, ptr %p0, ptr %p1) {
; CHECK-LABEL: gep_index_matches(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<6>;
-; CHECK-NEXT: .reg .b64 %rd<8>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<9>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [gep_index_matches_param_0];
-; CHECK-NEXT: xor.b32 %r2, %r1, 5;
; CHECK-NEXT: ld.param.b64 %rd1, [gep_index_matches_param_1];
; CHECK-NEXT: ld.param.b64 %rd2, [gep_index_matches_param_2];
-; CHECK-NEXT: mul.wide.u32 %rd3, %r2, 4;
-; CHECK-NEXT: add.s64 %rd4, %rd1, %rd3;
-; CHECK-NEXT: ld.param.b64 %rd5, [gep_index_matches_param_3];
-; CHECK-NEXT: ld.b32 %r3, [%rd4];
-; CHECK-NEXT: st.b32 [%rd2], %r3;
-; CHECK-NEXT: xor.b32 %r4, %r1, 6;
-; CHECK-NEXT: mul.wide.u32 %rd6, %r4, 4;
-; CHECK-NEXT: add.s64 %rd7, %rd1, %rd6;
-; CHECK-NEXT: ld.b32 %r5, [%rd7];
-; CHECK-NEXT: st.b32 [%rd5], %r5;
+; CHECK-NEXT: mul.wide.u32 %rd3, %r1, 4;
+; CHECK-NEXT: xor.b64 %rd4, %rd3, 20;
+; CHECK-NEXT: add.s64 %rd5, %rd1, %rd4;
+; CHECK-NEXT: ld.param.b64 %rd6, [gep_index_matches_param_3];
+; CHECK-NEXT: ld.b32 %r2, [%rd5];
+; CHECK-NEXT: st.b32 [%rd2], %r2;
+; CHECK-NEXT: xor.b64 %rd7, %rd3, 24;
+; CHECK-NEXT: add.s64 %rd8, %rd1, %rd7;
+; CHECK-NEXT: ld.b32 %r3, [%rd8];
+; CHECK-NEXT: st.b32 [%rd6], %r3;
; CHECK-NEXT: ret;
%x0 = xor i32 %base, 5
%i0 = zext i32 %x0 to i64
>From ccc97da9cb41de88a22a53e100828fe3512bd778 Mon Sep 17 00:00:00 2001
From: Hao Ren <haor at nvidia.com>
Date: Fri, 7 Aug 2026 17:11:15 +0000
Subject: [PATCH 3/6] [NVPTX] Move shift-of-logic CSE combine into the target
DAG combiner
Move the shift-of-logic CSE rewrite out of the generic DAGCombiner and into
NVPTX's PerformSHLCombine, and drop the shouldCommuteShiftOfLogicToExposeCSE
target hook it was gated behind.
Because the combine now runs only on NVPTX -- which does not fold these shifts
into addressing modes and has no bitwise-logic tree combines to disrupt -- the
addressing-mode and logic-tree profitability guards are unnecessary. The only
remaining condition is that a matching sibling exists to share the shift. It
runs before mul.wide folding, which would otherwise consume the shift.
---
llvm/include/llvm/CodeGen/TargetLowering.h | 9 --
llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp | 126 ------------------
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 126 ++++++++++++++++++
llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 5 -
4 files changed, 126 insertions(+), 140 deletions(-)
diff --git a/llvm/include/llvm/CodeGen/TargetLowering.h b/llvm/include/llvm/CodeGen/TargetLowering.h
index 5ee77e93c79a1..46b52c30cded6 100644
--- a/llvm/include/llvm/CodeGen/TargetLowering.h
+++ b/llvm/include/llvm/CodeGen/TargetLowering.h
@@ -4662,15 +4662,6 @@ class LLVM_ABI TargetLowering : public TargetLoweringBase {
return true;
}
- /// Return true if it is profitable to commute a constant SHL through
- /// multiple bitwise logic operations that share an input. This can expose a
- /// common shifted operand, but may also prevent shifts from folding into
- /// addressing modes, so targets must opt in.
- virtual bool shouldCommuteShiftOfLogicToExposeCSE(const SDNode *N,
- CombineLevel Level) const {
- return false;
- }
-
/// GlobalISel - return true if it is profitable to move this shift by a
/// constant amount through its operand, adjusting any immediate operands as
/// necessary to preserve semantics. This transformation may not be desirable
diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
index 5dd7275ab8b85..88fe89274b0b2 100644
--- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
@@ -590,7 +590,6 @@ namespace {
EVT VT, SDValue N0, SDValue N1,
SDNodeFlags Flags = SDNodeFlags());
- SDValue foldShiftOfLogicWithSharedInput(SDNode *N);
SDValue visitShiftByConstant(SDNode *N);
SDValue foldSelectOfConstants(SDNode *N);
@@ -10780,127 +10779,6 @@ static SDValue combineShiftOfShiftedLogic(SDNode *Shift, SelectionDAG &DAG) {
LogicOp->getFlags());
}
-/// Commute SHL with a bitwise logic operation when doing so exposes a common
-/// shifted operand. For example:
-///
-/// Before:
-/// N = shl (zext (Logic X, C)), ShiftAmount
-/// SiblingShift = shl (zext (SiblingLogic X, SiblingC)), ShiftAmount
-///
-/// After:
-/// ShiftedX = shl (zext X), ShiftAmount
-/// N = Logic ShiftedX, ShiftedC
-/// SiblingShift = SiblingLogic ShiftedX, SiblingShiftedC
-///
-/// ShiftedC = (zext C) << ShiftAmount and SiblingShiftedC =
-/// (zext SiblingC) << ShiftAmount are folded constants. This replaces two
-/// variable shifts with the single shared ShiftedX. Requiring a matching
-/// sibling avoids disrupting isolated address calculations where a shift may be
-/// folded into the addressing mode.
-SDValue DAGCombiner::foldShiftOfLogicWithSharedInput(SDNode *N) {
- struct ShiftOfLogic {
- SDNode *Shift;
- SDValue Logic;
- SDValue X;
- SDValue Constant;
- unsigned ExtendOpcode;
- };
-
- // Match a logic operation, with an optional extension, inside a SHL.
- auto matchShiftOfLogic = [&](SDNode *Shift, ShiftOfLogic &Match) {
- if (Shift->getOpcode() != ISD::SHL || !Shift->getOperand(0).hasOneUse())
- return false;
-
- SDValue ShiftInput = Shift->getOperand(0);
- unsigned ExtendOpcode = 0;
- SDValue Logic = ShiftInput;
- if (ISD::isExtOpcode(ShiftInput.getOpcode())) {
- ExtendOpcode = ShiftInput.getOpcode();
- Logic = ShiftInput.getOperand(0);
- }
-
- SDValue X, C;
- if (!sd_match(Logic, m_OneUse(m_BitwiseLogic(m_Value(X),
- m_Value(C, m_ConstInt())))))
- return false;
-
- Match = {Shift, Logic, X, C, ExtendOpcode};
- return true;
- };
-
- // Match N as the root shift-of-logic; bail if it does not fit the pattern or
- // should not be commuted.
- ShiftOfLogic Root;
- if (!matchShiftOfLogic(N, Root) ||
- !TLI.shouldCommuteShiftOfLogicToExposeCSE(N, Level))
- return SDValue();
-
- EVT VT = N->getValueType(0);
- SDValue ShiftAmount = N->getOperand(1);
-
- // Collect candidate shifts that share X. Reached through another user of X,
- // the logic result feeds the shift directly or through an optional extend.
- SmallVector<SDNode *, 4> CandidateShifts;
- for (SDNode *CandidateLogic : Root.X->users()) {
- if (CandidateLogic == Root.Logic.getNode())
- continue;
- for (SDNode *LogicUser : CandidateLogic->users()) {
- if (ISD::isExtOpcode(LogicUser->getOpcode())) {
- // shl (ext (logic X, C)): step through the extend to find the shift.
- for (SDNode *ExtendUser : LogicUser->users())
- if (ExtendUser->getOpcode() == ISD::SHL)
- CandidateShifts.push_back(ExtendUser);
- } else if (LogicUser->getOpcode() == ISD::SHL) {
- // shl (logic X, C): the user is already the shift.
- CandidateShifts.push_back(LogicUser);
- }
- }
- }
-
- // Verify each candidate against the root's pattern; those that match (same X,
- // extend, type, and shift amount) become siblings.
- SmallVector<ShiftOfLogic, 4> Siblings;
- for (SDNode *CandidateShift : CandidateShifts) {
- ShiftOfLogic Candidate;
- if (!matchShiftOfLogic(CandidateShift, Candidate) ||
- Candidate.X != Root.X || Candidate.ExtendOpcode != Root.ExtendOpcode ||
- CandidateShift->getValueType(0) != VT ||
- CandidateShift->getOperand(1) != ShiftAmount ||
- !TLI.isDesirableToCommuteWithShift(CandidateShift, Level))
- continue;
- Siblings.push_back(Candidate);
- }
- if (Siblings.empty())
- return SDValue();
-
- // Build the shared shifted X once, then rewrite the root and every sibling
- // into a logic op over it so the shift is CSE'd.
- SDValue ShiftedX = DAG.getNode(
- ISD::SHL, SDLoc(N), VT,
- Root.ExtendOpcode ? DAG.getNode(Root.ExtendOpcode, SDLoc(N), VT, Root.X)
- : Root.X,
- ShiftAmount);
-
- // Rebuild the logic op from shared ShiftedX and a folded constant shift.
- auto rebuild = [&](SDValue OldLogic, SDValue OldC, const SDLoc &NodeDL) {
- SDValue NewC = OldC;
- if (Root.ExtendOpcode)
- NewC = DAG.getNode(Root.ExtendOpcode, NodeDL, VT, OldC);
- SDValue ShiftedC = DAG.getNode(ISD::SHL, NodeDL, VT, NewC, ShiftAmount);
- return DAG.getNode(OldLogic.getOpcode(), NodeDL, VT, ShiftedX, ShiftedC,
- OldLogic->getFlags());
- };
-
- for (const ShiftOfLogic &Sibling : Siblings) {
- SDValue NewLogic =
- rebuild(Sibling.Logic, Sibling.Constant, SDLoc(Sibling.Shift));
- DAG.ReplaceAllUsesOfValueWith(SDValue(Sibling.Shift, 0), NewLogic);
- deleteAndRecombine(Sibling.Shift);
- AddToWorklist(NewLogic.getNode());
- }
- return rebuild(Root.Logic, Root.Constant, SDLoc(N));
-}
-
/// Handle transforms common to the three shifts, when the shift amount is a
/// constant.
/// We are looking for: (shift being one of shl/sra/srl)
@@ -10923,10 +10801,6 @@ SDValue DAGCombiner::visitShiftByConstant(SDNode *N) {
if (SDValue R = combineShiftOfShiftedLogic(N, DAG))
return R;
- // Fold sibling shifts of logic operations to expose a shared shifted input.
- if (SDValue R = foldShiftOfLogicWithSharedInput(N))
- return R;
-
// We want to pull some binops through shifts, so that we have (and (shift))
// instead of (shift (and)), likewise for add, or, xor, etc. This sort of
// thing happens with address calculations, so it's important to canonicalize
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index e788b0e44041f..21ec90a1f6420 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -6616,11 +6616,137 @@ static SDValue PerformMULCombine(SDNode *N,
return PerformMULCombineWithOperands(N, N0, N1, DCI);
}
+/// Commute SHL with a bitwise logic operation when doing so exposes a common
+/// shifted operand. For example:
+///
+/// Before:
+/// N = shl (zext (Logic X, C)), ShiftAmount
+/// SiblingShift = shl (zext (SiblingLogic X, SiblingC)), ShiftAmount
+///
+/// After:
+/// ShiftedX = shl (zext X), ShiftAmount
+/// N = Logic ShiftedX, ShiftedC
+/// SiblingShift = SiblingLogic ShiftedX, SiblingShiftedC
+///
+/// ShiftedC = (zext C) << ShiftAmount and SiblingShiftedC =
+/// (zext SiblingC) << ShiftAmount are folded constants. This replaces two
+/// variable shifts with the single shared ShiftedX. Requiring a matching
+/// sibling avoids disrupting isolated address calculations where a shift may be
+/// folded into the addressing mode.
+static SDValue
+PerformShiftOfLogicCSECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
+ using namespace SDPatternMatch;
+ SelectionDAG &DAG = DCI.DAG;
+
+ struct ShiftOfLogic {
+ SDNode *Shift;
+ SDValue Logic;
+ SDValue X;
+ SDValue Constant;
+ unsigned ExtendOpcode;
+ };
+
+ // Match a logic operation, with an optional extension, inside a SHL.
+ auto matchShiftOfLogic = [&](SDNode *Shift, ShiftOfLogic &Match) {
+ if (Shift->getOpcode() != ISD::SHL || !Shift->getOperand(0).hasOneUse())
+ return false;
+
+ SDValue ShiftInput = Shift->getOperand(0);
+ unsigned ExtendOpcode = 0;
+ SDValue Logic = ShiftInput;
+ if (ISD::isExtOpcode(ShiftInput.getOpcode())) {
+ ExtendOpcode = ShiftInput.getOpcode();
+ Logic = ShiftInput.getOperand(0);
+ }
+
+ SDValue X, C;
+ if (!sd_match(Logic, m_OneUse(m_BitwiseLogic(m_Value(X),
+ m_Value(C, m_ConstInt())))))
+ return false;
+
+ Match = {Shift, Logic, X, C, ExtendOpcode};
+ return true;
+ };
+
+ // Match N as the root shift-of-logic; bail if it does not fit the pattern.
+ ShiftOfLogic Root;
+ if (!matchShiftOfLogic(N, Root))
+ return SDValue();
+
+ EVT VT = N->getValueType(0);
+ SDValue ShiftAmount = N->getOperand(1);
+
+ // Collect candidate shifts that share X. Reached through another user of X,
+ // the logic result feeds the shift directly or through an optional extend.
+ SmallVector<SDNode *, 4> CandidateShifts;
+ for (SDNode *CandidateLogic : Root.X->users()) {
+ if (CandidateLogic == Root.Logic.getNode())
+ continue;
+ for (SDNode *LogicUser : CandidateLogic->users()) {
+ if (ISD::isExtOpcode(LogicUser->getOpcode())) {
+ // shl (ext (logic X, C)): step through the extend to find the shift.
+ for (SDNode *ExtendUser : LogicUser->users())
+ if (ExtendUser->getOpcode() == ISD::SHL)
+ CandidateShifts.push_back(ExtendUser);
+ } else if (LogicUser->getOpcode() == ISD::SHL) {
+ // shl (logic X, C): the user is already the shift.
+ CandidateShifts.push_back(LogicUser);
+ }
+ }
+ }
+
+ // Verify each candidate against the root's pattern; those that match (same X,
+ // extend, type, and shift amount) become siblings.
+ SmallVector<ShiftOfLogic, 4> Siblings;
+ for (SDNode *CandidateShift : CandidateShifts) {
+ ShiftOfLogic Candidate;
+ if (!matchShiftOfLogic(CandidateShift, Candidate) ||
+ Candidate.X != Root.X || Candidate.ExtendOpcode != Root.ExtendOpcode ||
+ CandidateShift->getValueType(0) != VT ||
+ CandidateShift->getOperand(1) != ShiftAmount)
+ continue;
+ Siblings.push_back(Candidate);
+ }
+ if (Siblings.empty())
+ return SDValue();
+
+ // Build the shared shifted X once, then rewrite the root and every sibling
+ // into a logic op over it so the shift is CSE'd.
+ SDValue ShiftedX = DAG.getNode(
+ ISD::SHL, SDLoc(N), VT,
+ Root.ExtendOpcode ? DAG.getNode(Root.ExtendOpcode, SDLoc(N), VT, Root.X)
+ : Root.X,
+ ShiftAmount);
+
+ // Rebuild the logic op from shared ShiftedX and a folded constant shift.
+ auto rebuild = [&](SDValue OldLogic, SDValue OldC, const SDLoc &NodeDL) {
+ SDValue NewC = OldC;
+ if (Root.ExtendOpcode)
+ NewC = DAG.getNode(Root.ExtendOpcode, NodeDL, VT, OldC);
+ SDValue ShiftedC = DAG.getNode(ISD::SHL, NodeDL, VT, NewC, ShiftAmount);
+ return DAG.getNode(OldLogic.getOpcode(), NodeDL, VT, ShiftedX, ShiftedC,
+ OldLogic->getFlags());
+ };
+
+ for (const ShiftOfLogic &Sibling : Siblings) {
+ SDValue NewLogic =
+ rebuild(Sibling.Logic, Sibling.Constant, SDLoc(Sibling.Shift));
+ DAG.ReplaceAllUsesOfValueWith(SDValue(Sibling.Shift, 0), NewLogic);
+ DCI.AddToWorklist(NewLogic.getNode());
+ }
+ return rebuild(Root.Logic, Root.Constant, SDLoc(N));
+}
+
/// PerformSHLCombine - Runs PTX-specific DAG combine patterns on SHL nodes.
static SDValue PerformSHLCombine(SDNode *N,
TargetLowering::DAGCombinerInfo &DCI,
CodeGenOptLevel OptLevel) {
if (OptLevel > CodeGenOptLevel::None) {
+ // Expose a shared shifted operand for CSE before mul.wide folding, which
+ // would otherwise consume the shift.
+ if (SDValue Ret = PerformShiftOfLogicCSECombine(N, DCI))
+ return Ret;
+
// Try mul.wide combining at OptLevel > 0
if (SDValue Ret = TryMULWIDECombine(N, DCI))
return Ret;
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 92ff7415dd6f5..42741ba98b438 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -181,11 +181,6 @@ class NVPTXTargetLowering : public TargetLowering {
TargetLoweringOpt &TLO,
unsigned Depth = 0) const override;
- bool shouldCommuteShiftOfLogicToExposeCSE(const SDNode *,
- CombineLevel) const override {
- return true;
- }
-
private:
const NVPTXSubtarget &STI; // cache the subtarget here
mutable unsigned GlobalUniqueCallSite;
>From 736e435e310060cbe3aa85c6d22062dc0373d2fe Mon Sep 17 00:00:00 2001
From: Hao Ren <haor at nvidia.com>
Date: Tue, 11 Aug 2026 18:40:13 +0000
Subject: [PATCH 4/6] [NVPTX] Address review feedback on shift-of-logic CSE
combine
- Require a constant shift amount; with a variable amount the per-op constant
shift cannot fold and would add a shift instead of removing one. Add a
negative test for it.
- matchShiftOfLogicOp returns std::optional<ShiftOfLogicOp>, populated in place.
- Rename ShiftOfLogic/Logic to ShiftOfLogicOp/LogicOp.
- Express the sibling check as a single positive condition.
- Use DCI.CombineTo for the sibling rewrites.
- Declare locals near first use and extend the doc-comment example.
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 121 ++++++++++----------
llvm/test/CodeGen/NVPTX/shift-logic-cse.ll | 30 +++++
2 files changed, 93 insertions(+), 58 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 21ec90a1f6420..92856bc99391a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -6620,69 +6620,73 @@ static SDValue PerformMULCombine(SDNode *N,
/// shifted operand. For example:
///
/// Before:
-/// N = shl (zext (Logic X, C)), ShiftAmount
-/// SiblingShift = shl (zext (SiblingLogic X, SiblingC)), ShiftAmount
+/// N = shl (zext (LogicOp X, C)), ShiftAmount
+/// SiblingShift = shl (zext (SiblingLogicOp X, SiblingC)), ShiftAmount
+/// OtherShift = shl (zext (OtherLogicOp X, OtherC)), ShiftAmount
///
/// After:
/// ShiftedX = shl (zext X), ShiftAmount
-/// N = Logic ShiftedX, ShiftedC
-/// SiblingShift = SiblingLogic ShiftedX, SiblingShiftedC
+/// N = LogicOp ShiftedX, ShiftedC
+/// SiblingShift = SiblingLogicOp ShiftedX, ShiftedSiblingC
+/// OtherShift = OtherLogicOp ShiftedX, ShiftedOtherC
///
-/// ShiftedC = (zext C) << ShiftAmount and SiblingShiftedC =
-/// (zext SiblingC) << ShiftAmount are folded constants. This replaces two
-/// variable shifts with the single shared ShiftedX. Requiring a matching
-/// sibling avoids disrupting isolated address calculations where a shift may be
-/// folded into the addressing mode.
+/// ShiftedC = (zext C) << ShiftAmount, and likewise ShiftedSiblingC and
+/// ShiftedOtherC, are folded constants. This replaces three variable shifts
+/// with the single shared ShiftedX. Requiring a matching sibling avoids
+/// disrupting isolated address calculations where a shift may be folded into
+/// the addressing mode.
static SDValue
PerformShiftOfLogicCSECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
using namespace SDPatternMatch;
- SelectionDAG &DAG = DCI.DAG;
- struct ShiftOfLogic {
+ struct ShiftOfLogicOp {
SDNode *Shift;
- SDValue Logic;
+ SDValue LogicOp;
SDValue X;
SDValue Constant;
unsigned ExtendOpcode;
};
// Match a logic operation, with an optional extension, inside a SHL.
- auto matchShiftOfLogic = [&](SDNode *Shift, ShiftOfLogic &Match) {
+ auto matchShiftOfLogicOp =
+ [&](SDNode *Shift) -> std::optional<ShiftOfLogicOp> {
if (Shift->getOpcode() != ISD::SHL || !Shift->getOperand(0).hasOneUse())
- return false;
+ return std::nullopt;
- SDValue ShiftInput = Shift->getOperand(0);
- unsigned ExtendOpcode = 0;
- SDValue Logic = ShiftInput;
- if (ISD::isExtOpcode(ShiftInput.getOpcode())) {
- ExtendOpcode = ShiftInput.getOpcode();
- Logic = ShiftInput.getOperand(0);
+ ShiftOfLogicOp Match;
+ Match.Shift = Shift;
+ Match.LogicOp = Shift->getOperand(0);
+ Match.ExtendOpcode = 0;
+ if (ISD::isExtOpcode(Match.LogicOp.getOpcode())) {
+ Match.ExtendOpcode = Match.LogicOp.getOpcode();
+ Match.LogicOp = Match.LogicOp.getOperand(0);
}
- SDValue X, C;
- if (!sd_match(Logic, m_OneUse(m_BitwiseLogic(m_Value(X),
- m_Value(C, m_ConstInt())))))
- return false;
+ if (!sd_match(Match.LogicOp, m_OneUse(m_BitwiseLogic(
+ m_Value(Match.X),
+ m_Value(Match.Constant, m_ConstInt())))))
+ return std::nullopt;
- Match = {Shift, Logic, X, C, ExtendOpcode};
- return true;
+ return Match;
};
// Match N as the root shift-of-logic; bail if it does not fit the pattern.
- ShiftOfLogic Root;
- if (!matchShiftOfLogic(N, Root))
+ std::optional<ShiftOfLogicOp> Root = matchShiftOfLogicOp(N);
+ if (!Root)
return SDValue();
- EVT VT = N->getValueType(0);
- SDValue ShiftAmount = N->getOperand(1);
+ // Only profitable for a constant shift amount: the per-op constant shift then
+ // folds away instead of becoming an extra variable shift.
+ if (!isConstOrConstSplat(N->getOperand(1)))
+ return SDValue();
// Collect candidate shifts that share X. Reached through another user of X,
// the logic result feeds the shift directly or through an optional extend.
SmallVector<SDNode *, 4> CandidateShifts;
- for (SDNode *CandidateLogic : Root.X->users()) {
- if (CandidateLogic == Root.Logic.getNode())
+ for (SDNode *CandidateLogicOp : Root->X->users()) {
+ if (CandidateLogicOp == Root->LogicOp.getNode())
continue;
- for (SDNode *LogicUser : CandidateLogic->users()) {
+ for (SDNode *LogicUser : CandidateLogicOp->users()) {
if (ISD::isExtOpcode(LogicUser->getOpcode())) {
// shl (ext (logic X, C)): step through the extend to find the shift.
for (SDNode *ExtendUser : LogicUser->users())
@@ -6697,44 +6701,45 @@ PerformShiftOfLogicCSECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
// Verify each candidate against the root's pattern; those that match (same X,
// extend, type, and shift amount) become siblings.
- SmallVector<ShiftOfLogic, 4> Siblings;
+ EVT VT = N->getValueType(0);
+ SDValue ShiftAmount = N->getOperand(1);
+ SmallVector<ShiftOfLogicOp, 4> Siblings;
for (SDNode *CandidateShift : CandidateShifts) {
- ShiftOfLogic Candidate;
- if (!matchShiftOfLogic(CandidateShift, Candidate) ||
- Candidate.X != Root.X || Candidate.ExtendOpcode != Root.ExtendOpcode ||
- CandidateShift->getValueType(0) != VT ||
- CandidateShift->getOperand(1) != ShiftAmount)
- continue;
- Siblings.push_back(Candidate);
+ std::optional<ShiftOfLogicOp> Candidate =
+ matchShiftOfLogicOp(CandidateShift);
+ if (Candidate && Candidate->X == Root->X &&
+ Candidate->ExtendOpcode == Root->ExtendOpcode &&
+ CandidateShift->getValueType(0) == VT &&
+ CandidateShift->getOperand(1) == ShiftAmount)
+ Siblings.push_back(*Candidate);
}
if (Siblings.empty())
return SDValue();
// Build the shared shifted X once, then rewrite the root and every sibling
// into a logic op over it so the shift is CSE'd.
- SDValue ShiftedX = DAG.getNode(
- ISD::SHL, SDLoc(N), VT,
- Root.ExtendOpcode ? DAG.getNode(Root.ExtendOpcode, SDLoc(N), VT, Root.X)
- : Root.X,
- ShiftAmount);
+ SelectionDAG &DAG = DCI.DAG;
+ SDValue ShiftedX =
+ DAG.getNode(ISD::SHL, SDLoc(N), VT,
+ Root->ExtendOpcode
+ ? DAG.getNode(Root->ExtendOpcode, SDLoc(N), VT, Root->X)
+ : Root->X,
+ ShiftAmount);
// Rebuild the logic op from shared ShiftedX and a folded constant shift.
- auto rebuild = [&](SDValue OldLogic, SDValue OldC, const SDLoc &NodeDL) {
+ auto rebuild = [&](SDValue OldLogicOp, SDValue OldC, const SDLoc &NodeDL) {
SDValue NewC = OldC;
- if (Root.ExtendOpcode)
- NewC = DAG.getNode(Root.ExtendOpcode, NodeDL, VT, OldC);
+ if (Root->ExtendOpcode)
+ NewC = DAG.getNode(Root->ExtendOpcode, NodeDL, VT, OldC);
SDValue ShiftedC = DAG.getNode(ISD::SHL, NodeDL, VT, NewC, ShiftAmount);
- return DAG.getNode(OldLogic.getOpcode(), NodeDL, VT, ShiftedX, ShiftedC,
- OldLogic->getFlags());
+ return DAG.getNode(OldLogicOp.getOpcode(), NodeDL, VT, ShiftedX, ShiftedC,
+ OldLogicOp->getFlags());
};
- for (const ShiftOfLogic &Sibling : Siblings) {
- SDValue NewLogic =
- rebuild(Sibling.Logic, Sibling.Constant, SDLoc(Sibling.Shift));
- DAG.ReplaceAllUsesOfValueWith(SDValue(Sibling.Shift, 0), NewLogic);
- DCI.AddToWorklist(NewLogic.getNode());
- }
- return rebuild(Root.Logic, Root.Constant, SDLoc(N));
+ for (const ShiftOfLogicOp &Sibling : Siblings)
+ DCI.CombineTo(Sibling.Shift, rebuild(Sibling.LogicOp, Sibling.Constant,
+ SDLoc(Sibling.Shift)));
+ return rebuild(Root->LogicOp, Root->Constant, SDLoc(N));
}
/// PerformSHLCombine - Runs PTX-specific DAG combine patterns on SHL nodes.
diff --git a/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll b/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
index 0aa689c2ea2ed..9892da6e94f35 100644
--- a/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
+++ b/llvm/test/CodeGen/NVPTX/shift-logic-cse.ll
@@ -148,3 +148,33 @@ define i32 @isolated_no_match(i32 %x) {
%shl = shl i32 %xor, 2
ret i32 %shl
}
+
+; Negative test: with a variable shift amount the folded constant shift cannot
+; be constant-folded, so it would add a shift rather than remove one. The
+; combine requires a constant shift amount and leaves these alone.
+define void @variable_shift_amount(i32 %x, i32 %s, ptr %p0, ptr %p1) {
+; CHECK-LABEL: variable_shift_amount(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<7>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [variable_shift_amount_param_0];
+; CHECK-NEXT: and.b32 %r2, %r1, 5;
+; CHECK-NEXT: ld.param.b32 %r3, [variable_shift_amount_param_1];
+; CHECK-NEXT: shl.b32 %r4, %r2, %r3;
+; CHECK-NEXT: ld.param.b64 %rd1, [variable_shift_amount_param_2];
+; CHECK-NEXT: st.b32 [%rd1], %r4;
+; CHECK-NEXT: ld.param.b64 %rd2, [variable_shift_amount_param_3];
+; CHECK-NEXT: or.b32 %r5, %r1, 6;
+; CHECK-NEXT: shl.b32 %r6, %r5, %r3;
+; CHECK-NEXT: st.b32 [%rd2], %r6;
+; CHECK-NEXT: ret;
+ %and = and i32 %x, 5
+ %and.shl = shl i32 %and, %s
+ store i32 %and.shl, ptr %p0, align 4
+ %or = or i32 %x, 6
+ %or.shl = shl i32 %or, %s
+ store i32 %or.shl, ptr %p1, align 4
+ ret void
+}
>From c320c703969b83feeb931b56490c78028b517182 Mon Sep 17 00:00:00 2001
From: Hao Ren <haor at nvidia.com>
Date: Wed, 12 Aug 2026 20:30:07 +0000
Subject: [PATCH 5/6] [NVPTX] Address remaining shift-of-logic review feedback
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 72 ++++++++++-----------
1 file changed, 35 insertions(+), 37 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 92856bc99391a..514c81d71214a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -6620,23 +6620,21 @@ static SDValue PerformMULCombine(SDNode *N,
/// shifted operand. For example:
///
/// Before:
-/// N = shl (zext (LogicOp X, C)), ShiftAmount
-/// SiblingShift = shl (zext (SiblingLogicOp X, SiblingC)), ShiftAmount
-/// OtherShift = shl (zext (OtherLogicOp X, OtherC)), ShiftAmount
+/// N = shl (zext (LogicOp X, C)), ShiftAmount
+/// OtherShift = shl (zext (OtherLogicOp X, OtherC)), ShiftAmount
///
/// After:
-/// ShiftedX = shl (zext X), ShiftAmount
-/// N = LogicOp ShiftedX, ShiftedC
-/// SiblingShift = SiblingLogicOp ShiftedX, ShiftedSiblingC
-/// OtherShift = OtherLogicOp ShiftedX, ShiftedOtherC
+/// ShiftedX = shl (zext X), ShiftAmount
+/// N = LogicOp ShiftedX, ShiftedC
+/// OtherShift = OtherLogicOp ShiftedX, ShiftedOtherC
///
-/// ShiftedC = (zext C) << ShiftAmount, and likewise ShiftedSiblingC and
-/// ShiftedOtherC, are folded constants. This replaces three variable shifts
-/// with the single shared ShiftedX. Requiring a matching sibling avoids
-/// disrupting isolated address calculations where a shift may be folded into
-/// the addressing mode.
-static SDValue
-PerformShiftOfLogicCSECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
+/// ShiftedC = (zext C) << ShiftAmount and ShiftedOtherC =
+/// (zext OtherC) << ShiftAmount are folded constants. This replaces two
+/// variable shifts with the single shared ShiftedX. Requiring another matching
+/// shift avoids disrupting isolated address calculations where a shift may be
+/// folded into the addressing mode.
+static SDValue combineShiftOfLogicOp(SDNode *N,
+ TargetLowering::DAGCombinerInfo &DCI) {
using namespace SDPatternMatch;
struct ShiftOfLogicOp {
@@ -6652,7 +6650,6 @@ PerformShiftOfLogicCSECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
[&](SDNode *Shift) -> std::optional<ShiftOfLogicOp> {
if (Shift->getOpcode() != ISD::SHL || !Shift->getOperand(0).hasOneUse())
return std::nullopt;
-
ShiftOfLogicOp Match;
Match.Shift = Shift;
Match.LogicOp = Shift->getOperand(0);
@@ -6671,7 +6668,7 @@ PerformShiftOfLogicCSECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
};
// Match N as the root shift-of-logic; bail if it does not fit the pattern.
- std::optional<ShiftOfLogicOp> Root = matchShiftOfLogicOp(N);
+ const std::optional<ShiftOfLogicOp> Root = matchShiftOfLogicOp(N);
if (!Root)
return SDValue();
@@ -6699,27 +6696,27 @@ PerformShiftOfLogicCSECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
}
}
- // Verify each candidate against the root's pattern; those that match (same X,
- // extend, type, and shift amount) become siblings.
- EVT VT = N->getValueType(0);
- SDValue ShiftAmount = N->getOperand(1);
- SmallVector<ShiftOfLogicOp, 4> Siblings;
+ // Verify each candidate against the root's pattern: the same X, extension,
+ // type, and shift amount.
+ const EVT VT = N->getValueType(0);
+ const SDValue ShiftAmount = N->getOperand(1);
+ SmallVector<ShiftOfLogicOp, 4> Matches;
for (SDNode *CandidateShift : CandidateShifts) {
- std::optional<ShiftOfLogicOp> Candidate =
+ const std::optional<ShiftOfLogicOp> Candidate =
matchShiftOfLogicOp(CandidateShift);
if (Candidate && Candidate->X == Root->X &&
Candidate->ExtendOpcode == Root->ExtendOpcode &&
CandidateShift->getValueType(0) == VT &&
CandidateShift->getOperand(1) == ShiftAmount)
- Siblings.push_back(*Candidate);
+ Matches.push_back(*Candidate);
}
- if (Siblings.empty())
+ if (Matches.empty())
return SDValue();
- // Build the shared shifted X once, then rewrite the root and every sibling
+ // Build the shared shifted X once, then rewrite the root and every match
// into a logic op over it so the shift is CSE'd.
SelectionDAG &DAG = DCI.DAG;
- SDValue ShiftedX =
+ const SDValue ShiftedX =
DAG.getNode(ISD::SHL, SDLoc(N), VT,
Root->ExtendOpcode
? DAG.getNode(Root->ExtendOpcode, SDLoc(N), VT, Root->X)
@@ -6727,19 +6724,20 @@ PerformShiftOfLogicCSECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
ShiftAmount);
// Rebuild the logic op from shared ShiftedX and a folded constant shift.
- auto rebuild = [&](SDValue OldLogicOp, SDValue OldC, const SDLoc &NodeDL) {
- SDValue NewC = OldC;
+ auto buildCommutedLogicOp = [&](const SDValue LogicOp, SDValue C,
+ const SDLoc &DL) {
if (Root->ExtendOpcode)
- NewC = DAG.getNode(Root->ExtendOpcode, NodeDL, VT, OldC);
- SDValue ShiftedC = DAG.getNode(ISD::SHL, NodeDL, VT, NewC, ShiftAmount);
- return DAG.getNode(OldLogicOp.getOpcode(), NodeDL, VT, ShiftedX, ShiftedC,
- OldLogicOp->getFlags());
+ C = DAG.getNode(Root->ExtendOpcode, DL, VT, C);
+ const SDValue ShiftedC = DAG.getNode(ISD::SHL, DL, VT, C, ShiftAmount);
+ return DAG.getNode(LogicOp.getOpcode(), DL, VT, ShiftedX, ShiftedC,
+ LogicOp->getFlags());
};
- for (const ShiftOfLogicOp &Sibling : Siblings)
- DCI.CombineTo(Sibling.Shift, rebuild(Sibling.LogicOp, Sibling.Constant,
- SDLoc(Sibling.Shift)));
- return rebuild(Root->LogicOp, Root->Constant, SDLoc(N));
+ for (const ShiftOfLogicOp &Match : Matches)
+ DCI.CombineTo(Match.Shift,
+ buildCommutedLogicOp(Match.LogicOp, Match.Constant,
+ SDLoc(Match.Shift)));
+ return buildCommutedLogicOp(Root->LogicOp, Root->Constant, SDLoc(N));
}
/// PerformSHLCombine - Runs PTX-specific DAG combine patterns on SHL nodes.
@@ -6749,7 +6747,7 @@ static SDValue PerformSHLCombine(SDNode *N,
if (OptLevel > CodeGenOptLevel::None) {
// Expose a shared shifted operand for CSE before mul.wide folding, which
// would otherwise consume the shift.
- if (SDValue Ret = PerformShiftOfLogicCSECombine(N, DCI))
+ if (SDValue Ret = combineShiftOfLogicOp(N, DCI))
return Ret;
// Try mul.wide combining at OptLevel > 0
>From d909effb93998b7a6b303c5c1759e82466e0aac4 Mon Sep 17 00:00:00 2001
From: Hao Ren <haor at nvidia.com>
Date: Thu, 20 Aug 2026 18:00:27 +0000
Subject: [PATCH 6/6] [NVPTX][NFC] Improve const correctness in shift-of-logic
combine
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 10 +++++-----
1 file changed, 5 insertions(+), 5 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 514c81d71214a..0bb0e28b2de92 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -6633,7 +6633,7 @@ static SDValue PerformMULCombine(SDNode *N,
/// variable shifts with the single shared ShiftedX. Requiring another matching
/// shift avoids disrupting isolated address calculations where a shift may be
/// folded into the addressing mode.
-static SDValue combineShiftOfLogicOp(SDNode *N,
+static SDValue combineShiftOfLogicOp(SDNode *const N,
TargetLowering::DAGCombinerInfo &DCI) {
using namespace SDPatternMatch;
@@ -6647,7 +6647,7 @@ static SDValue combineShiftOfLogicOp(SDNode *N,
// Match a logic operation, with an optional extension, inside a SHL.
auto matchShiftOfLogicOp =
- [&](SDNode *Shift) -> std::optional<ShiftOfLogicOp> {
+ [&](SDNode *const Shift) -> std::optional<ShiftOfLogicOp> {
if (Shift->getOpcode() != ISD::SHL || !Shift->getOperand(0).hasOneUse())
return std::nullopt;
ShiftOfLogicOp Match;
@@ -6680,13 +6680,13 @@ static SDValue combineShiftOfLogicOp(SDNode *N,
// Collect candidate shifts that share X. Reached through another user of X,
// the logic result feeds the shift directly or through an optional extend.
SmallVector<SDNode *, 4> CandidateShifts;
- for (SDNode *CandidateLogicOp : Root->X->users()) {
+ for (const SDNode *CandidateLogicOp : Root->X->users()) {
if (CandidateLogicOp == Root->LogicOp.getNode())
continue;
for (SDNode *LogicUser : CandidateLogicOp->users()) {
if (ISD::isExtOpcode(LogicUser->getOpcode())) {
// shl (ext (logic X, C)): step through the extend to find the shift.
- for (SDNode *ExtendUser : LogicUser->users())
+ for (SDNode *const ExtendUser : LogicUser->users())
if (ExtendUser->getOpcode() == ISD::SHL)
CandidateShifts.push_back(ExtendUser);
} else if (LogicUser->getOpcode() == ISD::SHL) {
@@ -6701,7 +6701,7 @@ static SDValue combineShiftOfLogicOp(SDNode *N,
const EVT VT = N->getValueType(0);
const SDValue ShiftAmount = N->getOperand(1);
SmallVector<ShiftOfLogicOp, 4> Matches;
- for (SDNode *CandidateShift : CandidateShifts) {
+ for (SDNode *const CandidateShift : CandidateShifts) {
const std::optional<ShiftOfLogicOp> Candidate =
matchShiftOfLogicOp(CandidateShift);
if (Candidate && Candidate->X == Root->X &&
More information about the llvm-commits
mailing list