[llvm] [SLSR] Fix slsr gep stride delta miscompile (PR #204278)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Jun 30 04:32:03 PDT 2026
https://github.com/fengfeng09 updated https://github.com/llvm/llvm-project/pull/204278
>From 680480f77c3c4cf91d069c301df9cb65c77165a8 Mon Sep 17 00:00:00 2001
From: "feng.feng" <feng.feng at iluvatar.com>
Date: Tue, 16 Jun 2026 19:01:53 +0800
Subject: [PATCH 1/2] [SLSR][NFC] Precommit miscompiled cases.
Signed-off-by: feng.feng <feng.feng at iluvatar.com>
---
.../slsr-gep-sext-wrap.ll | 134 ++++++++++++++++++
1 file changed, 134 insertions(+)
create mode 100644 llvm/test/Transforms/StraightLineStrengthReduce/slsr-gep-sext-wrap.ll
diff --git a/llvm/test/Transforms/StraightLineStrengthReduce/slsr-gep-sext-wrap.ll b/llvm/test/Transforms/StraightLineStrengthReduce/slsr-gep-sext-wrap.ll
new file mode 100644
index 0000000000000..83ca1c1602f29
--- /dev/null
+++ b/llvm/test/Transforms/StraightLineStrengthReduce/slsr-gep-sext-wrap.ll
@@ -0,0 +1,134 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 4
+; RUN: opt < %s -passes=slsr,gvn -S | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64-p:64:64:64"
+
+; The GEP index is sign-extended from i32 to the i64 pointer-index width. SLSR
+; must not factor the narrow stride, fold a Stride-delta in i32, and then
+; sign-extend that constant: sext does not distribute over a wrapping i32
+; subtraction. ScalarEvolution models `xor X, signmask` as a non-nsw
+; `add X, signmask`, so the two strides differ by a constant in i32 even though
+; sext(stride2) - sext(stride1) is not that constant. The candidates must stay
+; independent instead of being rewritten to g1 + 4 * sext(0x80000000).
+define void @slsr_gep_xor_signmask_no_rewrite(ptr %p, i32 %n, float %r) {
+; CHECK-LABEL: define void @slsr_gep_xor_signmask_no_rewrite(
+; CHECK-SAME: ptr [[P:%.*]], i32 [[N:%.*]], float [[R:%.*]]) {
+; CHECK-NEXT: [[I1:%.*]] = sext i32 [[N]] to i64
+; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
+; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
+; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 -8589934592
+; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
+; CHECK-NEXT: ret void
+;
+ %i1 = sext i32 %n to i64
+ %g1 = getelementptr float, ptr %p, i64 %i1
+ store float %r, ptr %g1, align 4
+
+ %x = xor i32 %n, -2147483648
+ %i2 = sext i32 %x to i64
+ %g2 = getelementptr float, ptr %p, i64 %i2
+ store float %r, ptr %g2, align 4
+ ret void
+}
+
+; The candidate stride `%x = add %m, %n` has no nsw flag, so %m + %n may
+; signed-overflow and sext(%x) != sext(%m) + sext(%n). The narrow stride must
+; not be factored, otherwise g2 is wrongly rewritten to g1 + 4 * sext(%m).
+define void @slsr_gep_add_no_rewrite(ptr %p, i32 %m, i32 %n, float %r) {
+; CHECK-LABEL: define void @slsr_gep_add_no_rewrite(
+; CHECK-SAME: ptr [[P:%.*]], i32 [[M:%.*]], i32 [[N:%.*]], float [[R:%.*]]) {
+; CHECK-NEXT: [[I1:%.*]] = sext i32 [[N]] to i64
+; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
+; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
+; CHECK-NEXT: [[TMP1:%.*]] = sext i32 [[M]] to i64
+; CHECK-NEXT: [[TMP2:%.*]] = shl i64 [[TMP1]], 2
+; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 [[TMP2]]
+; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
+; CHECK-NEXT: ret void
+;
+ %i1 = sext i32 %n to i64
+ %g1 = getelementptr float, ptr %p, i64 %i1
+ store float %r, ptr %g1, align 4
+
+ %x = add i32 %m, %n
+ %i2 = sext i32 %x to i64
+ %g2 = getelementptr float, ptr %p, i64 %i2
+ store float %r, ptr %g2, align 4
+ ret void
+}
+
+; Sanity: when the narrow stride difference provably does not wrap (nsw add),
+; sext distributes and the Stride-delta rewrite is still performed.
+define void @slsr_gep_nsw_still_rewrites(ptr %p, i32 %n, float %r) {
+; CHECK-LABEL: define void @slsr_gep_nsw_still_rewrites(
+; CHECK-SAME: ptr [[P:%.*]], i32 [[N:%.*]], float [[R:%.*]]) {
+; CHECK-NEXT: [[I1:%.*]] = sext i32 [[N]] to i64
+; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
+; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
+; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 64
+; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
+; CHECK-NEXT: ret void
+;
+ %i1 = sext i32 %n to i64
+ %g1 = getelementptr float, ptr %p, i64 %i1
+ store float %r, ptr %g1, align 4
+
+ %x = add nsw i32 %n, 16
+ %i2 = sext i32 %x to i64
+ %g2 = getelementptr float, ptr %p, i64 %i2
+ store float %r, ptr %g2, align 4
+ ret void
+}
+
+; Same miscompile with the roles swapped: the `xor X, signmask` stride is the
+; basis (g1, dominates) and the safe stride %n is the candidate (g2). The stride
+; difference still sign-extends non-linearly, so neither stride being the basis
+; nor the candidate makes the Stride-delta rewrite valid. g2 must stay
+; independent instead of being rewritten to g1 + 4 * sext(0x80000000).
+define void @slsr_gep_xor_signmask_as_basis(ptr %p, i32 %n, float %r) {
+; CHECK-LABEL: define void @slsr_gep_xor_signmask_as_basis(
+; CHECK-SAME: ptr [[P:%.*]], i32 [[N:%.*]], float [[R:%.*]]) {
+; CHECK-NEXT: [[X:%.*]] = xor i32 [[N]], -2147483648
+; CHECK-NEXT: [[I1:%.*]] = sext i32 [[X]] to i64
+; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
+; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
+; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 -8589934592
+; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
+; CHECK-NEXT: ret void
+;
+ %x = xor i32 %n, -2147483648
+ %i1 = sext i32 %x to i64
+ %g1 = getelementptr float, ptr %p, i64 %i1
+ store float %r, ptr %g1, align 4
+
+ %i2 = sext i32 %n to i64
+ %g2 = getelementptr float, ptr %p, i64 %i2
+ store float %r, ptr %g2, align 4
+ ret void
+}
+
+; A non-nsw `add X, C` basis with a constant stride delta is unsound too: %n+100
+; may signed-overflow, so sext(%n+100) != sext(%n)+100 and rewriting g2 to
+; g1 - 400 is wrong. The constant delta makes the basis stride participate, so
+; it must be screened even though the candidate stride %n cannot wrap.
+define void @slsr_gep_add_const_as_basis(ptr %p, i32 %n, float %r) {
+; CHECK-LABEL: define void @slsr_gep_add_const_as_basis(
+; CHECK-SAME: ptr [[P:%.*]], i32 [[N:%.*]], float [[R:%.*]]) {
+; CHECK-NEXT: [[X:%.*]] = add i32 [[N]], 100
+; CHECK-NEXT: [[I1:%.*]] = sext i32 [[X]] to i64
+; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
+; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
+; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 -400
+; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
+; CHECK-NEXT: ret void
+;
+ %x = add i32 %n, 100
+ %i1 = sext i32 %x to i64
+ %g1 = getelementptr float, ptr %p, i64 %i1
+ store float %r, ptr %g1, align 4
+
+ %i2 = sext i32 %n to i64
+ %g2 = getelementptr float, ptr %p, i64 %i2
+ store float %r, ptr %g2, align 4
+ ret void
+}
>From 77769ae9fa52a918f7a759a068f1d0a03094e1ec Mon Sep 17 00:00:00 2001
From: "feng.feng" <feng.feng at iluvatar.com>
Date: Tue, 16 Jun 2026 19:08:28 +0800
Subject: [PATCH 2/2] [SLSR] Fix GEP Stride-delta miscompile when factoring a
sign-extended wrapping index
A narrow GEP index is sign-extended to the index width before the
address computation, and the Stride-delta rewrite reconstructs the
stride difference in that sign-extended domain, requiring
sext(Sc) == sext(Sb) + sext(Delta). This fails when a stride does not
sign-extend linearly: an add/sub/mul/shl without `nsw`, or an
`xor X, signmask` (which ScalarEvolution models as a non-nsw
`add X, signmask`). For example `gep float, p, sext(xor n, signmask)`
was rewritten relative to `gep float, p, sext(n)` into
`g1 + 4 * sext(signmask)`, wrong for half of the inputs.
Add a wrap check (mayHaveSignedWrap) and screen it where each stride
participates:
- the rewritten candidate stride, at candidate allocation;
- the basis stride in candidatePredicate, only when the delta is a
constant (a variable delta is fine, e.g. stride_var);
- the deeper basis in compressPath, which re-derives a constant delta.
Assisted-by: Opus 4.8
Signed-off-by: feng.feng <feng.feng at iluvatar.com>
---
.../Scalar/StraightLineStrengthReduce.cpp | 81 ++++++++++++++++++-
.../CodeGen/AMDGPU/asyncmark-gfx12plus.ll | 37 +++++----
.../AMDGPU/barrier-signal-wait-latency.ll | 52 ++++++++----
.../NVPTX/slsr-var-delta.ll | 23 +++---
.../slsr-gep-sext-wrap.ll | 16 ++--
5 files changed, 156 insertions(+), 53 deletions(-)
diff --git a/llvm/lib/Transforms/Scalar/StraightLineStrengthReduce.cpp b/llvm/lib/Transforms/Scalar/StraightLineStrengthReduce.cpp
index abca7020e1c4f..44da66092f9ce 100644
--- a/llvm/lib/Transforms/Scalar/StraightLineStrengthReduce.cpp
+++ b/llvm/lib/Transforms/Scalar/StraightLineStrengthReduce.cpp
@@ -639,6 +639,51 @@ static void unifyBitWidth(APInt &A, APInt &B) {
B = B.sext(A.getBitWidth());
}
+// Whether sign-extending V to a wider type may not distribute over arithmetic,
+// i.e. the narrow value does not sign-extend linearly. This is true for:
+// - an add/sub/mul/shl that does not carry the `nsw` flag, or
+// - an `xor X, signmask`: it merely flips the sign bit, but ScalarEvolution
+// models it as a non-nsw `add X, signmask`. So a stride `xor X, signmask`
+// and a stride X look like they differ by the constant signmask in the
+// narrow type, even though sext(xor X, signmask) - sext(X) is not that
+// constant (it is +signmask or -signmask depending on the sign of X).
+static bool mayHaveSignedWrap(const Value *V) {
+ if (match(V, m_Xor(m_Value(), m_SignMask())))
+ return true;
+ // OverflowingBinaryOperator covers exactly add/sub/mul/shl.
+ const auto *OBO = dyn_cast<OverflowingBinaryOperator>(V);
+ if (!OBO)
+ return true;
+ return !OBO->hasNoSignedWrap();
+}
+
+// True when the GEP index is narrower than the index width, i.e. it is
+// implicitly sign-extended to the index width (not the pointer width) of the
+// address space before the address computation. A value already at or wider
+// than the index width is not sign-extended (it is used as-is or truncated), so
+// it cannot trigger the non-distributing-sext problem.
+static bool isSignExtendedGepIndex(const Value *Idx, GetElementPtrInst *GEP,
+ const DataLayout *DL) {
+ return Idx->getType()->getIntegerBitWidth() <
+ DL->getIndexSizeInBits(GEP->getAddressSpace());
+}
+
+// A narrow GEP index is sign-extended to the index width before the address
+// computation. SLSR's Stride-delta rewrite turns two such GEPs into
+// Basis + Index * (Sc - Sb), so the stride difference Sc - Sb is reconstructed
+// in the sign-extended domain. This requires sext(Sc) == sext(Sb) +
+// sext(Delta).
+//
+// This screens the rewritten candidate's stride Sc = Sb + Delta: if Sc is
+// computed by a possibly-wrapping op, sext(Sc) does not equal sext(Sb) +
+// sext(Delta) and the rewrite would produce a wrong pointer.
+static bool isSafeToFactorGepIndex(const Value *Idx, GetElementPtrInst *GEP,
+ const DataLayout *DL) {
+ if (isSignExtendedGepIndex(Idx, GEP, DL))
+ return !mayHaveSignedWrap(Idx);
+ return true;
+}
+
Value *StraightLineStrengthReduce::getDelta(const Candidate &C,
const Candidate &Basis,
Candidate::DKind K) const {
@@ -694,6 +739,26 @@ bool StraightLineStrengthReduce::candidatePredicate(Candidate *Basis,
if (!Delta)
return false;
+ // For a GEP Stride-delta rewrite g2 = g1 + Index * Delta, the addresses are
+ // computed from the sign-extended strides, so this requires
+ // sext(Sc) == sext(Sb) + sext(Delta).
+ //
+ // The rewritten candidate's stride Sc = Sb + Delta is already screened
+ // broadly at allocation time (allocateCandidatesAndFindBasis): a wrapping Sc
+ // breaks the identity for any Delta. The basis's stride Sb = Sc - Delta only
+ // needs screening when Delta folds to a *constant*: then sext(Sb) + C can
+ // differ from sext(Sc) if Sb wraps. For a *variable* Delta the basis may wrap
+ // and still be sound, because the candidate stride carries the no-wrap
+ // guarantee (e.g. Sc is an `add nsw`, as in stride_var); rejecting it would
+ // pessimize those.
+ if (K == Candidate::StrideDelta && C.CandidateKind == Candidate::GEP &&
+ isa<ConstantInt>(Delta)) {
+ auto *BasisGEP = cast<GetElementPtrInst>(Basis->Ins);
+ if (isSignExtendedGepIndex(Basis->Stride, BasisGEP, DL) &&
+ mayHaveSignedWrap(Basis->Stride))
+ return false;
+ }
+
// IndexDelta rewrite is not always profitable, e.g.,
// X = B + 8 * S
// Y = B + S,
@@ -839,6 +904,18 @@ auto StraightLineStrengthReduce::compressPath(Candidate &C,
if (!isSimilar(C, *NextRoot, CurrKind))
break;
+ // Path compression folds a constant Stride-delta directly against the
+ // deeper basis NextRoot, bypassing candidatePredicate's wrap guard. With a
+ // constant delta sext(Sb) + C can differ from sext(Sc) if the deeper
+ // basis's stride wraps, so do not compress past such a basis (mirrors the
+ // check in candidatePredicate).
+ if (CurrKind == Candidate::StrideDelta &&
+ C.CandidateKind == Candidate::GEP &&
+ isSignExtendedGepIndex(NextRoot->Stride,
+ cast<GetElementPtrInst>(NextRoot->Ins), DL) &&
+ mayHaveSignedWrap(NextRoot->Stride))
+ break;
+
if (auto DeltaVal =
dyn_cast<SCEVConstant>(SE->getMinusSCEV(CandPart, BasisPart))) {
Root = NextRoot;
@@ -950,6 +1027,8 @@ bool StraightLineStrengthReduce::isFoldable(const Candidate &C,
void StraightLineStrengthReduce::allocateCandidatesAndFindBasis(
Candidate::Kind CT, const SCEV *B, ConstantInt *Idx, Value *S,
Instruction *I) {
+ bool IsSafe = CT != Candidate::GEP ||
+ isSafeToFactorGepIndex(S, cast<GetElementPtrInst>(I), DL);
// Record the SCEV of S that we may use it as a variable delta.
// Ensure that we rewrite C with a existing IR that reproduces delta value.
@@ -964,7 +1043,7 @@ void StraightLineStrengthReduce::allocateCandidatesAndFindBasis(
// mode if the constant falls within a certain range.
// So, we also check if the instruction is already high efficient enough
// for the strength reduction algorithm.
- if (!isFoldable(C, TTI) && !C.isHighEfficiency()) {
+ if (!isFoldable(C, TTI) && !C.isHighEfficiency() && IsSafe) {
setBasisAndDeltaFor(C);
// Compress unnecessary rewrite to improve ILP
diff --git a/llvm/test/CodeGen/AMDGPU/asyncmark-gfx12plus.ll b/llvm/test/CodeGen/AMDGPU/asyncmark-gfx12plus.ll
index 2a27d7ff34bd5..763596fa0cffe 100644
--- a/llvm/test/CodeGen/AMDGPU/asyncmark-gfx12plus.ll
+++ b/llvm/test/CodeGen/AMDGPU/asyncmark-gfx12plus.ll
@@ -373,8 +373,10 @@ define amdgpu_kernel void @test_pipelined_loop_with_global(ptr addrspace(1) %foo
; SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; SDAG-NEXT: v_dual_add_nc_u32 v2, v8, v7 :: v_dual_mov_b32 v0, s1
; SDAG-NEXT: ds_load_b32 v1, v0
+; SDAG-NEXT: v_mov_b32_e32 v5, s0
+; SDAG-NEXT: s_add_co_i32 s0, s11, -1
; SDAG-NEXT: s_wait_dscnt 0x0
-; SDAG-NEXT: v_dual_mov_b32 v5, s0 :: v_dual_add_nc_u32 v1, v2, v1
+; SDAG-NEXT: v_add_nc_u32_e32 v1, v2, v1
; SDAG-NEXT: global_store_b32 v5, v1, s[2:3] scale_offset
; SDAG-NEXT: ; wait_asyncmark(0)
; SDAG-NEXT: s_wait_asynccnt 0x0
@@ -384,8 +386,8 @@ define amdgpu_kernel void @test_pipelined_loop_with_global(ptr addrspace(1) %foo
; SDAG-NEXT: v_add_nc_u32_e32 v1, v3, v4
; SDAG-NEXT: s_wait_dscnt 0x0
; SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
-; SDAG-NEXT: v_add_nc_u32_e32 v0, v1, v0
-; SDAG-NEXT: global_store_b32 v5, v0, s[2:3] offset:4 scale_offset
+; SDAG-NEXT: v_dual_mov_b32 v1, s0 :: v_dual_add_nc_u32 v0, v1, v0
+; SDAG-NEXT: global_store_b32 v1, v0, s[2:3] scale_offset
; SDAG-NEXT: s_endpgm
;
; GISEL-LABEL: test_pipelined_loop_with_global:
@@ -422,20 +424,20 @@ define amdgpu_kernel void @test_pipelined_loop_with_global(ptr addrspace(1) %foo
; GISEL-NEXT: v_readfirstlane_b32 s15, v1
; GISEL-NEXT: s_wait_loadcnt 0x0
; GISEL-NEXT: v_readfirstlane_b32 s16, v2
-; GISEL-NEXT: s_mov_b32 s9, s15
+; GISEL-NEXT: s_mov_b32 s8, s15
; GISEL-NEXT: s_mov_b32 s19, s16
; GISEL-NEXT: .LBB2_1: ; %loop_body
; GISEL-NEXT: ; =>This Inner Loop Header: Depth=1
-; GISEL-NEXT: s_add_co_u32 s8, s14, 8
+; GISEL-NEXT: s_add_co_u32 s9, s14, 8
; GISEL-NEXT: s_clause 0x1
; GISEL-NEXT: global_load_b32 v1, v0, s[6:7]
; GISEL-NEXT: global_load_b32 v2, v0, s[0:1]
-; GISEL-NEXT: v_mov_b32_e32 v3, s8
-; GISEL-NEXT: s_mov_b32 s18, s9
+; GISEL-NEXT: v_mov_b32_e32 v3, s9
+; GISEL-NEXT: s_mov_b32 s18, s8
; GISEL-NEXT: s_wait_kmcnt 0x0
-; GISEL-NEXT: s_add_co_i32 s9, s12, s17
+; GISEL-NEXT: s_add_co_i32 s8, s12, s17
; GISEL-NEXT: s_add_co_i32 s13, s13, 1
-; GISEL-NEXT: s_mov_b32 s8, s19
+; GISEL-NEXT: s_mov_b32 s9, s19
; GISEL-NEXT: global_load_async_to_lds_b32 v3, v0, s[6:7] offset:4 nv
; GISEL-NEXT: v_mov_b32_e32 v3, s14
; GISEL-NEXT: ; asyncmark
@@ -445,10 +447,10 @@ define amdgpu_kernel void @test_pipelined_loop_with_global(ptr addrspace(1) %foo
; GISEL-NEXT: ds_load_b32 v3, v3
; GISEL-NEXT: s_wait_dscnt 0x0
; GISEL-NEXT: v_readfirstlane_b32 s17, v3
-; GISEL-NEXT: s_add_co_i32 s9, s9, s17
+; GISEL-NEXT: s_add_co_i32 s8, s8, s17
; GISEL-NEXT: s_wait_xcnt 0x0
; GISEL-NEXT: s_add_co_u32 s0, s0, 4
-; GISEL-NEXT: v_mov_b32_e32 v3, s9
+; GISEL-NEXT: v_mov_b32_e32 v3, s8
; GISEL-NEXT: s_add_co_ci_u32 s1, s1, 0
; GISEL-NEXT: s_add_co_u32 s6, s6, 4
; GISEL-NEXT: s_add_co_ci_u32 s7, s7, 0
@@ -460,7 +462,7 @@ define amdgpu_kernel void @test_pipelined_loop_with_global(ptr addrspace(1) %foo
; GISEL-NEXT: s_add_co_u32 s14, s14, 4
; GISEL-NEXT: s_cmp_lt_i32 s13, s11
; GISEL-NEXT: s_wait_loadcnt 0x1
-; GISEL-NEXT: v_readfirstlane_b32 s9, v1
+; GISEL-NEXT: v_readfirstlane_b32 s8, v1
; GISEL-NEXT: s_wait_loadcnt 0x0
; GISEL-NEXT: v_readfirstlane_b32 s19, v2
; GISEL-NEXT: s_cbranch_scc1 .LBB2_1
@@ -469,7 +471,7 @@ define amdgpu_kernel void @test_pipelined_loop_with_global(ptr addrspace(1) %foo
; GISEL-NEXT: ; wait_asyncmark(1)
; GISEL-NEXT: s_wait_asynccnt 0x1
; GISEL-NEXT: s_lshl_b32 s1, s0, 2
-; GISEL-NEXT: s_add_co_i32 s4, s18, s8
+; GISEL-NEXT: s_add_co_i32 s4, s18, s9
; GISEL-NEXT: s_add_co_u32 s1, s10, s1
; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; GISEL-NEXT: v_mov_b32_e32 v0, s1
@@ -480,7 +482,7 @@ define amdgpu_kernel void @test_pipelined_loop_with_global(ptr addrspace(1) %foo
; GISEL-NEXT: s_add_co_i32 s1, s4, s1
; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; GISEL-NEXT: v_mov_b32_e32 v2, s1
-; GISEL-NEXT: s_add_co_i32 s1, s9, s19
+; GISEL-NEXT: s_add_co_i32 s1, s8, s19
; GISEL-NEXT: global_store_b32 v1, v2, s[2:3] scale_offset
; GISEL-NEXT: ; wait_asyncmark(0)
; GISEL-NEXT: s_wait_asynccnt 0x0
@@ -488,9 +490,10 @@ define amdgpu_kernel void @test_pipelined_loop_with_global(ptr addrspace(1) %foo
; GISEL-NEXT: s_wait_dscnt 0x0
; GISEL-NEXT: v_readfirstlane_b32 s0, v0
; GISEL-NEXT: s_add_co_i32 s0, s1, s0
-; GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
-; GISEL-NEXT: v_mov_b32_e32 v0, s0
-; GISEL-NEXT: global_store_b32 v1, v0, s[2:3] offset:4 scale_offset
+; GISEL-NEXT: s_add_co_i32 s1, s11, -1
+; GISEL-NEXT: s_wait_xcnt 0x0
+; GISEL-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GISEL-NEXT: global_store_b32 v1, v0, s[2:3] scale_offset
; GISEL-NEXT: s_endpgm
prolog:
; Load first iteration
diff --git a/llvm/test/CodeGen/AMDGPU/barrier-signal-wait-latency.ll b/llvm/test/CodeGen/AMDGPU/barrier-signal-wait-latency.ll
index 0b512b00c3fbc..3193b0df876e4 100644
--- a/llvm/test/CodeGen/AMDGPU/barrier-signal-wait-latency.ll
+++ b/llvm/test/CodeGen/AMDGPU/barrier-signal-wait-latency.ll
@@ -115,26 +115,35 @@ define amdgpu_kernel void @test_barrier_multiple(ptr addrspace(1) %out, i32 %siz
; OPT-NEXT: s_load_b96 s[0:2], s[4:5], 0x24
; OPT-NEXT: v_and_b32_e32 v1, 0x3ff, v0
; OPT-NEXT: s_delay_alu instid0(VALU_DEP_1)
-; OPT-NEXT: v_lshlrev_b32_e32 v2, 2, v1
+; OPT-NEXT: v_lshlrev_b32_e32 v3, 2, v1
; OPT-NEXT: s_wait_kmcnt 0x0
-; OPT-NEXT: v_xad_u32 v0, v1, -1, s2
-; OPT-NEXT: global_store_b32 v2, v1, s[0:1]
+; OPT-NEXT: v_sub_nc_u32_e32 v2, s2, v1
+; OPT-NEXT: global_store_b32 v3, v1, s[0:1]
; OPT-NEXT: s_barrier_signal -1
+; OPT-NEXT: v_add_nc_u32_e32 v0, -1, v2
+; OPT-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
; OPT-NEXT: v_ashrrev_i32_e32 v1, 31, v0
+; OPT-NEXT: v_lshlrev_b64_e32 v[0:1], 2, v[0:1]
+; OPT-NEXT: s_barrier_wait -1
; OPT-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; OPT-NEXT: v_add_co_u32 v0, vcc_lo, s0, v0
+; OPT-NEXT: v_add_co_ci_u32_e64 v1, null, s1, v1, vcc_lo
+; OPT-NEXT: global_load_b32 v4, v[0:1], off
+; OPT-NEXT: v_add_nc_u32_e32 v0, -2, v2
+; OPT-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; OPT-NEXT: v_ashrrev_i32_e32 v1, 31, v0
; OPT-NEXT: v_lshlrev_b64_e32 v[0:1], 2, v[0:1]
+; OPT-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
; OPT-NEXT: v_add_co_u32 v0, vcc_lo, s0, v0
-; OPT-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; OPT-NEXT: s_wait_alu depctr_va_vcc(0)
; OPT-NEXT: v_add_co_ci_u32_e64 v1, null, s1, v1, vcc_lo
-; OPT-NEXT: s_barrier_wait -1
-; OPT-NEXT: global_load_b32 v3, v[0:1], off
; OPT-NEXT: s_wait_loadcnt 0x0
-; OPT-NEXT: global_store_b32 v2, v3, s[0:1]
+; OPT-NEXT: global_store_b32 v3, v4, s[0:1]
; OPT-NEXT: s_barrier_signal -1
; OPT-NEXT: s_barrier_wait -1
-; OPT-NEXT: global_load_b32 v0, v[0:1], off offset:-4
+; OPT-NEXT: global_load_b32 v0, v[0:1], off
; OPT-NEXT: s_wait_loadcnt 0x0
-; OPT-NEXT: global_store_b32 v2, v0, s[0:1]
+; OPT-NEXT: global_store_b32 v3, v0, s[0:1]
; OPT-NEXT: s_endpgm
;
; NOOPT-LABEL: test_barrier_multiple:
@@ -142,26 +151,35 @@ define amdgpu_kernel void @test_barrier_multiple(ptr addrspace(1) %out, i32 %siz
; NOOPT-NEXT: s_load_b96 s[0:2], s[4:5], 0x24
; NOOPT-NEXT: v_and_b32_e32 v2, 0x3ff, v0
; NOOPT-NEXT: s_delay_alu instid0(VALU_DEP_1)
-; NOOPT-NEXT: v_lshlrev_b32_e32 v3, 2, v2
+; NOOPT-NEXT: v_lshlrev_b32_e32 v4, 2, v2
; NOOPT-NEXT: s_wait_kmcnt 0x0
-; NOOPT-NEXT: v_xad_u32 v0, v2, -1, s2
-; NOOPT-NEXT: global_store_b32 v3, v2, s[0:1]
+; NOOPT-NEXT: v_sub_nc_u32_e32 v3, s2, v2
+; NOOPT-NEXT: global_store_b32 v4, v2, s[0:1]
; NOOPT-NEXT: s_barrier_signal -1
; NOOPT-NEXT: s_barrier_wait -1
-; NOOPT-NEXT: v_ashrrev_i32_e32 v1, 31, v0
+; NOOPT-NEXT: v_add_nc_u32_e32 v0, -1, v3
; NOOPT-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; NOOPT-NEXT: v_ashrrev_i32_e32 v1, 31, v0
; NOOPT-NEXT: v_lshlrev_b64_e32 v[0:1], 2, v[0:1]
+; NOOPT-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
; NOOPT-NEXT: v_add_co_u32 v0, vcc_lo, s0, v0
-; NOOPT-NEXT: s_delay_alu instid0(VALU_DEP_1)
; NOOPT-NEXT: v_add_co_ci_u32_e64 v1, null, s1, v1, vcc_lo
; NOOPT-NEXT: global_load_b32 v2, v[0:1], off
+; NOOPT-NEXT: v_add_nc_u32_e32 v0, -2, v3
+; NOOPT-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; NOOPT-NEXT: v_ashrrev_i32_e32 v1, 31, v0
+; NOOPT-NEXT: v_lshlrev_b64_e32 v[0:1], 2, v[0:1]
+; NOOPT-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
+; NOOPT-NEXT: v_add_co_u32 v0, vcc_lo, s0, v0
+; NOOPT-NEXT: s_wait_alu depctr_va_vcc(0)
+; NOOPT-NEXT: v_add_co_ci_u32_e64 v1, null, s1, v1, vcc_lo
; NOOPT-NEXT: s_wait_loadcnt 0x0
-; NOOPT-NEXT: global_store_b32 v3, v2, s[0:1]
+; NOOPT-NEXT: global_store_b32 v4, v2, s[0:1]
; NOOPT-NEXT: s_barrier_signal -1
; NOOPT-NEXT: s_barrier_wait -1
-; NOOPT-NEXT: global_load_b32 v0, v[0:1], off offset:-4
+; NOOPT-NEXT: global_load_b32 v0, v[0:1], off
; NOOPT-NEXT: s_wait_loadcnt 0x0
-; NOOPT-NEXT: global_store_b32 v3, v0, s[0:1]
+; NOOPT-NEXT: global_store_b32 v4, v0, s[0:1]
; NOOPT-NEXT: s_endpgm
entry:
%tid = call i32 @llvm.amdgcn.workitem.id.x()
diff --git a/llvm/test/Transforms/StraightLineStrengthReduce/NVPTX/slsr-var-delta.ll b/llvm/test/Transforms/StraightLineStrengthReduce/NVPTX/slsr-var-delta.ll
index ce6f720cc4352..6d5b4dcae03d9 100644
--- a/llvm/test/Transforms/StraightLineStrengthReduce/NVPTX/slsr-var-delta.ll
+++ b/llvm/test/Transforms/StraightLineStrengthReduce/NVPTX/slsr-var-delta.ll
@@ -11,7 +11,7 @@ define void @foo(ptr %a, ptr %b, i32 %j) {
; PTX-LABEL: foo(
; PTX: {
; PTX-NEXT: .reg .b32 %r<4>;
-; PTX-NEXT: .reg .b64 %rd<9>;
+; PTX-NEXT: .reg .b64 %rd<8>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0:
; PTX-NEXT: ld.param.b64 %rd1, [foo_param_0];
@@ -24,42 +24,41 @@ define void @foo(ptr %a, ptr %b, i32 %j) {
; PTX-NEXT: st.b32 [%rd4], 0;
; PTX-NEXT: add.s64 %rd5, %rd4, %rd3;
; PTX-NEXT: st.b32 [%rd5], 1;
-; PTX-NEXT: add.s64 %rd6, %rd5, 4;
; PTX-NEXT: st.b32 [%rd5+4], 2;
+; PTX-NEXT: add.s64 %rd6, %rd5, %rd3;
+; PTX-NEXT: st.b32 [%rd6+4], 3;
; PTX-NEXT: add.s64 %rd7, %rd6, %rd3;
-; PTX-NEXT: st.b32 [%rd7], 3;
-; PTX-NEXT: add.s64 %rd8, %rd7, %rd3;
-; PTX-NEXT: st.b32 [%rd8], 4;
+; PTX-NEXT: st.b32 [%rd7+4], 4;
; PTX-NEXT: ret;
%i.0 = load i32, ptr %a, align 8
- %i = add i32 %i.0, %j
+ %i = add nsw i32 %i.0, %j
; CHECK: [[L:%.*]] = load i32, ptr %a, align 8
- ; CHECK: [[I:%.*]] = add i32 [[L]], %j
+ ; CHECK: [[I:%.*]] = add nsw i32 [[L]], %j
%gep.24 = getelementptr float, ptr %b, i32 %i
; CHECK: [[GEP0:%.*]] = getelementptr float, ptr %b, i32 [[I]]
; CHECK: store i32 0, ptr [[GEP0]]
store i32 0, ptr %gep.24
- %gep.24.sum1 = add i32 %i, %i
+ %gep.24.sum1 = add nsw i32 %i, %i
%gep.25 = getelementptr float, ptr %b, i32 %gep.24.sum1
; CHECK: [[EXT1:%.*]] = sext i32 [[I]] to i64
; CHECK: [[MUL1:%.*]] = shl i64 [[EXT1]], 2
; CHECK: [[GEP1:%.*]] = getelementptr i8, ptr [[GEP0]], i64 [[MUL1]]
; CHECK: store i32 1, ptr [[GEP1]]
store i32 1, ptr %gep.25
- %gep.26.sum3 = add i32 1, %i
- %gep.27.sum = add i32 %gep.26.sum3, %i
+ %gep.26.sum3 = add nsw i32 1, %i
+ %gep.27.sum = add nsw i32 %gep.26.sum3, %i
%gep.28 = getelementptr float, ptr %b, i32 %gep.27.sum
; CHECK: [[GEP2:%.*]] = getelementptr i8, ptr [[GEP1]], i64 4
; CHECK: store i32 2, ptr [[GEP2]]
store i32 2, ptr %gep.28
- %gep.28.sum = add i32 %gep.27.sum, %i
+ %gep.28.sum = add nsw i32 %gep.27.sum, %i
%gep.29 = getelementptr float, ptr %b, i32 %gep.28.sum
; CHECK: [[EXT2:%.*]] = sext i32 [[I]] to i64
; CHECK: [[MUL2:%.*]] = shl i64 [[EXT2]], 2
; CHECK: [[GEP3:%.*]] = getelementptr i8, ptr [[GEP2]], i64 [[MUL2]]
; CHECK: store i32 3, ptr [[GEP3]]
store i32 3, ptr %gep.29
- %gep.29.sum = add i32 %gep.28.sum, %i
+ %gep.29.sum = add nsw i32 %gep.28.sum, %i
%gep.30 = getelementptr float, ptr %b, i32 %gep.29.sum
; CHECK: [[EXT3:%.*]] = sext i32 [[I]] to i64
; CHECK: [[MUL3:%.*]] = shl i64 [[EXT3]], 2
diff --git a/llvm/test/Transforms/StraightLineStrengthReduce/slsr-gep-sext-wrap.ll b/llvm/test/Transforms/StraightLineStrengthReduce/slsr-gep-sext-wrap.ll
index 83ca1c1602f29..5ea89bdaf979e 100644
--- a/llvm/test/Transforms/StraightLineStrengthReduce/slsr-gep-sext-wrap.ll
+++ b/llvm/test/Transforms/StraightLineStrengthReduce/slsr-gep-sext-wrap.ll
@@ -16,7 +16,9 @@ define void @slsr_gep_xor_signmask_no_rewrite(ptr %p, i32 %n, float %r) {
; CHECK-NEXT: [[I1:%.*]] = sext i32 [[N]] to i64
; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
-; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 -8589934592
+; CHECK-NEXT: [[X:%.*]] = xor i32 [[N]], -2147483648
+; CHECK-NEXT: [[I2:%.*]] = sext i32 [[X]] to i64
+; CHECK-NEXT: [[G2:%.*]] = getelementptr float, ptr [[P]], i64 [[I2]]
; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
; CHECK-NEXT: ret void
;
@@ -40,9 +42,9 @@ define void @slsr_gep_add_no_rewrite(ptr %p, i32 %m, i32 %n, float %r) {
; CHECK-NEXT: [[I1:%.*]] = sext i32 [[N]] to i64
; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
-; CHECK-NEXT: [[TMP1:%.*]] = sext i32 [[M]] to i64
-; CHECK-NEXT: [[TMP2:%.*]] = shl i64 [[TMP1]], 2
-; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 [[TMP2]]
+; CHECK-NEXT: [[X:%.*]] = add i32 [[M]], [[N]]
+; CHECK-NEXT: [[I2:%.*]] = sext i32 [[X]] to i64
+; CHECK-NEXT: [[G2:%.*]] = getelementptr float, ptr [[P]], i64 [[I2]]
; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
; CHECK-NEXT: ret void
;
@@ -92,7 +94,8 @@ define void @slsr_gep_xor_signmask_as_basis(ptr %p, i32 %n, float %r) {
; CHECK-NEXT: [[I1:%.*]] = sext i32 [[X]] to i64
; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
-; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 -8589934592
+; CHECK-NEXT: [[I2:%.*]] = sext i32 [[N]] to i64
+; CHECK-NEXT: [[G2:%.*]] = getelementptr float, ptr [[P]], i64 [[I2]]
; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
; CHECK-NEXT: ret void
;
@@ -118,7 +121,8 @@ define void @slsr_gep_add_const_as_basis(ptr %p, i32 %n, float %r) {
; CHECK-NEXT: [[I1:%.*]] = sext i32 [[X]] to i64
; CHECK-NEXT: [[G1:%.*]] = getelementptr float, ptr [[P]], i64 [[I1]]
; CHECK-NEXT: store float [[R]], ptr [[G1]], align 4
-; CHECK-NEXT: [[G2:%.*]] = getelementptr i8, ptr [[G1]], i64 -400
+; CHECK-NEXT: [[I2:%.*]] = sext i32 [[N]] to i64
+; CHECK-NEXT: [[G2:%.*]] = getelementptr float, ptr [[P]], i64 [[I2]]
; CHECK-NEXT: store float [[R]], ptr [[G2]], align 4
; CHECK-NEXT: ret void
;
More information about the llvm-commits
mailing list