[llvm] [AMDGPU][SCEV] Tighten max backedge-taken count for shift recurrence … (PR #197292)
via llvm-commits
llvm-commits at lists.llvm.org
Wed May 27 13:56:40 PDT 2026
https://github.com/carlobertolli updated https://github.com/llvm/llvm-project/pull/197292
>From 677fe52e57008f7b64a4c6e756e7b1316600614e Mon Sep 17 00:00:00 2001
From: carlobertolli <carlo.bertolli at amd.com>
Date: Tue, 12 May 2026 12:28:27 -0500
Subject: [PATCH] [AMDGPU][SCEV] Tighten max backedge-taken count for shift
recurrence loops. Application code on AMDGPUs often uses the following
pattern:
for (int i = blockDim.x / 2; i >= 1; i >>= 1)
if (threadIdx.x < i) {
<STMT based on threadIdx.x>
}
__syncthreads();
}
but this is currently not unrolled by LLVM for two reasons:
1. We overestimate the loop count with the number of bits in the iteration variable (32).
2. Runtime unrolling is not applicable because of convergent instructions.
For the specific pattern above, we can compute a tighter bound by using
range information, which for blockDim.x is (0,1025) and already emitted by clang,
resulting in 10 as upper bound as opposed to 32.
The next step in unrolling the loop is to adjust the threshold used to control
full loop unrolling with convergent instructions.
Assisted-by: Cursor (Claude)
---
llvm/lib/Analysis/ScalarEvolution.cpp | 51 ++++++---
.../test/Analysis/ScalarEvolution/shift-op.ll | 102 +++++++++++++++++-
2 files changed, 139 insertions(+), 14 deletions(-)
diff --git a/llvm/lib/Analysis/ScalarEvolution.cpp b/llvm/lib/Analysis/ScalarEvolution.cpp
index 855ab3bb5d621..f87245e575c04 100644
--- a/llvm/lib/Analysis/ScalarEvolution.cpp
+++ b/llvm/lib/Analysis/ScalarEvolution.cpp
@@ -9683,10 +9683,11 @@ ScalarEvolution::ExitLimit ScalarEvolution::computeShiftCompareExitLimit(
return getCouldNotCompute();
// Return true if V is of the form "LHS `shift_op` <positive constant>".
- // Return LHS in OutLHS and shift_opt in OutOpCode.
- auto MatchPositiveShift =
- [](Value *V, Value *&OutLHS, Instruction::BinaryOps &OutOpCode) {
-
+ // Return LHS in OutLHS, shift_op in OutOpCode, and the shift amount in
+ // OutShiftAmt.
+ auto MatchPositiveShift = [](Value *V, Value *&OutLHS,
+ Instruction::BinaryOps &OutOpCode,
+ unsigned &OutShiftAmt) {
using namespace PatternMatch;
ConstantInt *ShiftAmt;
@@ -9699,7 +9700,11 @@ ScalarEvolution::ExitLimit ScalarEvolution::computeShiftCompareExitLimit(
else
return false;
- return ShiftAmt->getValue().isStrictlyPositive();
+ uint64_t Amt = ShiftAmt->getValue().getLimitedValue();
+ if (Amt == 0 || Amt >= OutLHS->getType()->getScalarSizeInBits())
+ return false;
+ OutShiftAmt = Amt;
+ return true;
};
// Recognize a "shift recurrence" either of the form %iv or of %iv.shifted in
@@ -9709,14 +9714,17 @@ ScalarEvolution::ExitLimit ScalarEvolution::computeShiftCompareExitLimit(
// %iv.shifted = lshr i32 %iv, <positive constant>
//
// Return true on a successful match. Return the corresponding PHI node (%iv
- // above) in PNOut and the opcode of the shift operation in OpCodeOut.
- auto MatchShiftRecurrence =
- [&](Value *V, PHINode *&PNOut, Instruction::BinaryOps &OpCodeOut) {
+ // above) in PNOut, the opcode of the shift operation in OpCodeOut, and the
+ // shift amount in ShiftAmtOut.
+ auto MatchShiftRecurrence = [&](Value *V, PHINode *&PNOut,
+ Instruction::BinaryOps &OpCodeOut,
+ unsigned &ShiftAmtOut) {
std::optional<Instruction::BinaryOps> PostShiftOpCode;
{
Instruction::BinaryOps OpC;
Value *V;
+ unsigned Amt;
// If we encounter a shift instruction, "peel off" the shift operation,
// and remember that we did so. Later when we inspect %iv's backedge
@@ -9727,7 +9735,7 @@ ScalarEvolution::ExitLimit ScalarEvolution::computeShiftCompareExitLimit(
// instruction as the one feeding into the PHI's backedge value. We only
// really care about it being the same *kind* of shift instruction --
// that's all that is required for our later inferences to hold.
- if (MatchPositiveShift(LHS, V, OpC)) {
+ if (MatchPositiveShift(LHS, V, OpC, Amt)) {
PostShiftOpCode = OpC;
LHS = V;
}
@@ -9743,7 +9751,7 @@ ScalarEvolution::ExitLimit ScalarEvolution::computeShiftCompareExitLimit(
return
// The backedge value for the PHI node must be a shift by a positive
// amount
- MatchPositiveShift(BEValue, OpLHS, OpCodeOut) &&
+ MatchPositiveShift(BEValue, OpLHS, OpCodeOut, ShiftAmtOut) &&
// of the PHI node itself
OpLHS == PNOut &&
@@ -9755,7 +9763,8 @@ ScalarEvolution::ExitLimit ScalarEvolution::computeShiftCompareExitLimit(
PHINode *PN;
Instruction::BinaryOps OpCode;
- if (!MatchShiftRecurrence(LHS, PN, OpCode))
+ unsigned ShiftAmt;
+ if (!MatchShiftRecurrence(LHS, PN, OpCode, ShiftAmt))
return getCouldNotCompute();
const DataLayout &DL = getDataLayout();
@@ -9803,8 +9812,26 @@ ScalarEvolution::ExitLimit ScalarEvolution::computeShiftCompareExitLimit(
if (Result->isNullValue()) {
unsigned BitWidth = getTypeSizeInBits(RHS->getType());
+ unsigned MaxBTC = BitWidth;
+
+ // For right-shift recurrences (lshr/ashr with non-negative start), we can
+ // compute a tighter max backedge-taken count from the range of the start
+ // value. After k shifts of ShiftAmt, value = start >> (k * ShiftAmt).
+ // The value reaches 0 (the stable value) when k * ShiftAmt >=
+ // activeBits(start), so max BTC = ceil(activeBits(maxStart) / ShiftAmt).
+ if (OpCode == Instruction::LShr || OpCode == Instruction::AShr) {
+ Value *StartValue = PN->getIncomingValueForBlock(Predecessor);
+ const SCEV *StartSCEV = getSCEV(StartValue);
+ APInt MaxStart = getUnsignedRangeMax(StartSCEV);
+ if (MaxStart.isStrictlyPositive()) {
+ unsigned ActiveBits = MaxStart.getActiveBits();
+ unsigned RangeBTC = divideCeil(ActiveBits, ShiftAmt);
+ MaxBTC = std::min(MaxBTC, RangeBTC);
+ }
+ }
+
const SCEV *UpperBound =
- getConstant(getEffectiveSCEVType(RHS->getType()), BitWidth);
+ getConstant(getEffectiveSCEVType(RHS->getType()), MaxBTC);
return ExitLimit(getCouldNotCompute(), UpperBound, UpperBound, false);
}
diff --git a/llvm/test/Analysis/ScalarEvolution/shift-op.ll b/llvm/test/Analysis/ScalarEvolution/shift-op.ll
index 150e4c996716b..36ec0341d7560 100644
--- a/llvm/test/Analysis/ScalarEvolution/shift-op.ll
+++ b/llvm/test/Analysis/ScalarEvolution/shift-op.ll
@@ -71,8 +71,8 @@ define void @test3(ptr %init.ptr) {
; CHECK-LABEL: 'test3'
; CHECK-NEXT: Determining loop execution counts for: @test3
; CHECK-NEXT: Loop %loop: Unpredictable backedge-taken count.
-; CHECK-NEXT: Loop %loop: constant max backedge-taken count is i32 32
-; CHECK-NEXT: Loop %loop: symbolic max backedge-taken count is i32 32
+; CHECK-NEXT: Loop %loop: constant max backedge-taken count is i32 16
+; CHECK-NEXT: Loop %loop: symbolic max backedge-taken count is i32 16
;
entry:
%init = load i32, ptr %init.ptr, !range !0
@@ -222,5 +222,103 @@ loop:
br i1 %exit.cond, label %leave, label %loop
}
+; Shift amount equals bit width (32 for i32) - produces poison per LangRef,
+; so we do not recognize this as a valid shift recurrence.
+define void @test_shift_amount_eq_bitwidth(i32 %init) {
+;
+; CHECK-LABEL: 'test_shift_amount_eq_bitwidth'
+; CHECK-NEXT: Determining loop execution counts for: @test_shift_amount_eq_bitwidth
+; CHECK-NEXT: Loop %loop: Unpredictable backedge-taken count.
+; CHECK-NEXT: Loop %loop: Unpredictable constant max backedge-taken count.
+; CHECK-NEXT: Loop %loop: Unpredictable symbolic max backedge-taken count.
+;
+ entry:
+ br label %loop
+
+ loop:
+ %iv = phi i32 [ %init, %entry ], [ %iv.shift, %loop ]
+ %iv.shift = lshr i32 %iv, 32
+ %exit.cond = icmp eq i32 %iv, 0
+ br i1 %exit.cond, label %leave, label %loop
+
+ leave:
+ ret void
+}
+
+; lshr by 1 with range metadata: activeBits(49999) = 16, MaxBTC = 16.
+; Mirrors test3 but with lshr instead of ashr.
+define void @test_lshr_range(ptr %init.ptr) {
+;
+; CHECK-LABEL: 'test_lshr_range'
+; CHECK-NEXT: Determining loop execution counts for: @test_lshr_range
+; CHECK-NEXT: Loop %loop: Unpredictable backedge-taken count.
+; CHECK-NEXT: Loop %loop: constant max backedge-taken count is i32 16
+; CHECK-NEXT: Loop %loop: symbolic max backedge-taken count is i32 16
+;
+ entry:
+ %init = load i32, ptr %init.ptr, !range !0
+ br label %loop
+
+ loop:
+ %iv = phi i32 [ %init, %entry ], [ %iv.shift, %loop ]
+ %iv.shift = lshr i32 %iv, 1
+ %exit.cond = icmp eq i32 %iv, 0
+ br i1 %exit.cond, label %leave, label %loop
+
+ leave:
+ ret void
+}
+
+; lshr by 3 with range metadata: activeBits(49999) = 16,
+; MaxBTC = ceil(16/3) = 6.
+define void @test_lshr_shift3_range(ptr %init.ptr) {
+;
+; CHECK-LABEL: 'test_lshr_shift3_range'
+; CHECK-NEXT: Determining loop execution counts for: @test_lshr_shift3_range
+; CHECK-NEXT: Loop %loop: Unpredictable backedge-taken count.
+; CHECK-NEXT: Loop %loop: constant max backedge-taken count is i32 6
+; CHECK-NEXT: Loop %loop: symbolic max backedge-taken count is i32 6
+;
+ entry:
+ %init = load i32, ptr %init.ptr, !range !0
+ br label %loop
+
+ loop:
+ %iv = phi i32 [ %init, %entry ], [ %iv.shift, %loop ]
+ %iv.shift = lshr i32 %iv, 3
+ %exit.cond = icmp eq i32 %iv, 0
+ br i1 %exit.cond, label %leave, label %loop
+
+ leave:
+ ret void
+}
+
+; GPU-style reduction: for (i = blockDim.x/2; i >= 1; i >>= 1)
+; %blockdim has !range [1, 1025), so %start = lshr %blockdim, 1 has max 512.
+; activeBits(512) = 10, MaxBTC = 10.
+define void @test_gpu_reduction(ptr %blockdim.ptr) {
+;
+; CHECK-LABEL: 'test_gpu_reduction'
+; CHECK-NEXT: Determining loop execution counts for: @test_gpu_reduction
+; CHECK-NEXT: Loop %loop: Unpredictable backedge-taken count.
+; CHECK-NEXT: Loop %loop: constant max backedge-taken count is i32 10
+; CHECK-NEXT: Loop %loop: symbolic max backedge-taken count is i32 10
+;
+ entry:
+ %blockdim = load i32, ptr %blockdim.ptr, !range !2
+ %start = lshr i32 %blockdim, 1
+ br label %loop
+
+ loop:
+ %iv = phi i32 [ %start, %entry ], [ %iv.shift, %loop ]
+ %iv.shift = lshr i32 %iv, 1
+ %exit.cond = icmp eq i32 %iv, 0
+ br i1 %exit.cond, label %leave, label %loop
+
+ leave:
+ ret void
+}
+
!0 = !{i32 0, i32 50000}
!1 = !{i32 -5000, i32 -1}
+!2 = !{i32 1, i32 1025}
More information about the llvm-commits
mailing list