[llvm] [GVN] Add code to disable scalar PRE (PR #190386)
Daniel Donenfeld via llvm-commits
llvm-commits at lists.llvm.org
Thu May 7 07:22:01 PDT 2026
https://github.com/daniel-donenfeld updated https://github.com/llvm/llvm-project/pull/190386
>From 2befcf4f99c02c6f3d6a55a1fe64bafabe12a461 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Mon, 16 Mar 2026 16:50:53 +0000
Subject: [PATCH 1/6] Add code to disable scalar PRE in GVN
---
llvm/include/llvm/Transforms/Scalar/GVN.h | 9 +++++
llvm/lib/Passes/PassBuilder.cpp | 2 ++
llvm/lib/Passes/PassRegistry.def | 2 +-
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 3 +-
llvm/lib/Transforms/Scalar/GVN.cpp | 34 ++++++++++++++-----
llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll | 26 ++++++++++++++
6 files changed, 65 insertions(+), 11 deletions(-)
create mode 100644 llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
diff --git a/llvm/include/llvm/Transforms/Scalar/GVN.h b/llvm/include/llvm/Transforms/Scalar/GVN.h
index f79896e993d9b..ccc608eb0c42b 100644
--- a/llvm/include/llvm/Transforms/Scalar/GVN.h
+++ b/llvm/include/llvm/Transforms/Scalar/GVN.h
@@ -77,6 +77,7 @@ class GVNLegacyPass;
/// additional setters and then pass it to GVN.
struct GVNOptions {
std::optional<bool> AllowPRE;
+ std::optional<bool> AllowScalarPRE;
std::optional<bool> AllowLoadPRE;
std::optional<bool> AllowLoadInLoopPRE;
std::optional<bool> AllowLoadPRESplitBackedge;
@@ -91,6 +92,12 @@ struct GVNOptions {
return *this;
}
+ /// Enables or disables PRE of scalars in GVN.
+ GVNOptions &setScalarPRE(bool ScalarPRE) {
+ AllowScalarPRE = ScalarPRE;
+ return *this;
+ }
+
/// Enables or disables PRE of loads in GVN.
GVNOptions &setLoadPRE(bool LoadPRE) {
AllowLoadPRE = LoadPRE;
@@ -149,6 +156,7 @@ class GVNPass : public PassInfoMixin<GVNPass> {
MemoryDependenceResults &getMemDep() const { return *MD; }
LLVM_ABI bool isPREEnabled() const;
+ LLVM_ABI bool isScalarPREEnabled() const;
LLVM_ABI bool isLoadPREEnabled() const;
LLVM_ABI bool isLoadInLoopPREEnabled() const;
LLVM_ABI bool isLoadPRESplitBackedgeEnabled() const;
@@ -408,6 +416,7 @@ class GVNPass : public PassInfoMixin<GVNPass> {
};
/// Create a legacy GVN pass.
+LLVM_ABI FunctionPass *createGVNPass(bool ScalarPRE);
LLVM_ABI FunctionPass *createGVNPass();
/// A simple and fast domtree-based GVN pass to hoist common expressions
diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp
index 55a4e99e7402e..25d5ac2e1e446 100644
--- a/llvm/lib/Passes/PassBuilder.cpp
+++ b/llvm/lib/Passes/PassBuilder.cpp
@@ -1331,6 +1331,8 @@ Expected<GVNOptions> parseGVNOptions(StringRef Params) {
bool Enable = !ParamName.consume_front("no-");
if (ParamName == "pre") {
Result.setPRE(Enable);
+ } else if (ParamName == "scalar-pre") {
+ Result.setScalarPRE(Enable);
} else if (ParamName == "load-pre") {
Result.setLoadPRE(Enable);
} else if (ParamName == "split-backedge-load-pre") {
diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def
index c92d93d7ae396..c64bf9bf729b0 100644
--- a/llvm/lib/Passes/PassRegistry.def
+++ b/llvm/lib/Passes/PassRegistry.def
@@ -598,7 +598,7 @@ FUNCTION_PASS_WITH_PARAMS(
FUNCTION_PASS_WITH_PARAMS(
"gvn", "GVNPass", [](GVNOptions Opts) { return GVNPass(Opts); },
parseGVNOptions,
- "no-pre;pre;no-load-pre;load-pre;no-split-backedge-load-pre;"
+ "no-pre;pre;no-scalar-pre;scalar-pre;no-load-pre;load-pre;no-split-backedge-load-pre;"
"split-backedge-load-pre;no-memdep;memdep;no-memoryssa;memoryssa")
FUNCTION_PASS_WITH_PARAMS(
"hardware-loops", "HardwareLoopsPass",
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index 10e746c502c09..9351c8dde60d4 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -291,7 +291,8 @@ NVPTXTargetMachine::getPredicatedAddrSpace(const Value *V) const {
void NVPTXPassConfig::addEarlyCSEOrGVNPass() {
if (getOptLevel() == CodeGenOptLevel::Aggressive)
- addPass(createGVNPass());
+ // Disable scalar PRE due to Register Pressure increase
+ addPass(createGVNPass(/*ScalarPRE=*/false));
else
addPass(createEarlyCSEPass());
}
diff --git a/llvm/lib/Transforms/Scalar/GVN.cpp b/llvm/lib/Transforms/Scalar/GVN.cpp
index 7cab4be169123..1a592ff160950 100644
--- a/llvm/lib/Transforms/Scalar/GVN.cpp
+++ b/llvm/lib/Transforms/Scalar/GVN.cpp
@@ -106,6 +106,8 @@ STATISTIC(MaxBBSpeculationCutoffReachedTimes,
"preventing further exploration");
static cl::opt<bool> GVNEnablePRE("enable-pre", cl::init(true), cl::Hidden);
+static cl::opt<bool> GVNEnableScalarPRE("enable-scalar-pre", cl::init(true),
+ cl::Hidden);
static cl::opt<bool> GVNEnableLoadPRE("enable-load-pre", cl::init(true));
static cl::opt<bool> GVNEnableLoadInLoopPRE("enable-load-in-loop-pre",
cl::init(true));
@@ -854,6 +856,10 @@ bool GVNPass::isPREEnabled() const {
return Options.AllowPRE.value_or(GVNEnablePRE);
}
+bool GVNPass::isScalarPREEnabled() const {
+ return Options.AllowScalarPRE.value_or(GVNEnableScalarPRE);
+}
+
bool GVNPass::isLoadPREEnabled() const {
return Options.AllowLoadPRE.value_or(GVNEnableLoadPRE);
}
@@ -915,6 +921,8 @@ void GVNPass::printPipeline(
OS << '<';
if (Options.AllowPRE != std::nullopt)
OS << (*Options.AllowPRE ? "" : "no-") << "pre;";
+ if (Options.AllowScalarPRE != std::nullopt)
+ OS << (*Options.AllowScalarPRE ? "" : "no-") << "scalar-pre;";
if (Options.AllowLoadPRE != std::nullopt)
OS << (*Options.AllowLoadPRE ? "" : "no-") << "load-pre;";
if (Options.AllowLoadPRESplitBackedge != std::nullopt)
@@ -2029,12 +2037,15 @@ bool GVNPass::processNonLocalLoad(LoadInst *Load) {
}
bool Changed = false;
- // If this load follows a GEP, see if we can PRE the indices before analyzing.
- if (GetElementPtrInst *GEP =
- dyn_cast<GetElementPtrInst>(Load->getOperand(0))) {
- for (Use &U : GEP->indices())
- if (Instruction *I = dyn_cast<Instruction>(U.get()))
- Changed |= performScalarPRE(I);
+ // This is a limited form of scalar PRE for load indices. If this load follows
+ // a GEP, see if we can PRE the indices before analyzing.
+ if (isPREEnabled() && isScalarPREEnabled()) {
+ if (GetElementPtrInst *GEP =
+ dyn_cast<GetElementPtrInst>(Load->getOperand(0))) {
+ for (Use &U : GEP->indices())
+ if (Instruction *I = dyn_cast<Instruction>(U.get()))
+ Changed |= performScalarPRE(I);
+ }
}
// Step 2: Analyze the availability of the load.
@@ -2842,7 +2853,7 @@ bool GVNPass::runImpl(Function &F, AssumptionCache &RunAC, DominatorTree &RunDT,
++Iteration;
}
- if (isPREEnabled()) {
+ if (isPREEnabled() && isScalarPREEnabled()) {
// Fabricate val-num for dead-code in order to suppress assertion in
// performPRE().
assignValNumForDeadCode();
@@ -3345,10 +3356,12 @@ class llvm::gvn::GVNLegacyPass : public FunctionPass {
static char ID; // Pass identification, replacement for typeid.
explicit GVNLegacyPass(bool MemDepAnalysis = GVNEnableMemDep,
- bool MemSSAAnalysis = GVNEnableMemorySSA)
+ bool MemSSAAnalysis = GVNEnableMemorySSA,
+ bool ScalarPRE = true)
: FunctionPass(ID), Impl(GVNOptions()
.setMemDep(MemDepAnalysis)
- .setMemorySSA(MemSSAAnalysis)) {
+ .setMemorySSA(MemSSAAnalysis)
+ .setScalarPRE(ScalarPRE)) {
initializeGVNLegacyPassPass(*PassRegistry::getPassRegistry());
}
@@ -3410,3 +3423,6 @@ INITIALIZE_PASS_END(GVNLegacyPass, "gvn", "Global Value Numbering", false, false
// The public interface to this file...
FunctionPass *llvm::createGVNPass() { return new GVNLegacyPass(); }
+FunctionPass *llvm::createGVNPass(bool ScalarPRE) {
+ return new GVNLegacyPass(GVNEnableMemDep, GVNEnableMemorySSA, ScalarPRE);
+}
diff --git a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
new file mode 100644
index 0000000000000..f11e9dba2bd6b
--- /dev/null
+++ b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
@@ -0,0 +1,26 @@
+; RUN: opt -enable-scalar-pre=false -enable-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK
+
+define void @kernel(ptr %arr, i8 %cond) {
+entry:
+ %tobool.not = icmp eq i8 %cond, 0
+ %tmp7.pre = load i32, ptr %arr, align 4
+ br i1 %tobool.not, label %if.end, label %if.then
+
+; CHECK: if.then:
+; CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[LOAD:%.*]], 2
+
+if.then: ; preds = %entry
+ %add = add nsw i32 %tmp7.pre, 2
+ %getElem = getelementptr inbounds nuw i8, ptr %arr, i64 8
+ store i32 %add, ptr %getElem, align 4
+ br label %if.end
+
+; CHECK: if.end:
+; CHECK-NEXT: [[ADD2:%.*]] = add nsw i32 [[LOAD:%.*]], 2
+
+if.end: ; preds = %if.then, %entry
+ %add8 = add nsw i32 %tmp7.pre, 2
+ %getElem1 = getelementptr inbounds nuw i8, ptr %arr, i64 12
+ store i32 %add8, ptr %getElem1, align 4
+ ret void
+}
\ No newline at end of file
>From 3d8429fce987a51b9c4bad7365f3700af5215b02 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Mon, 6 Apr 2026 19:34:44 +0000
Subject: [PATCH 2/6] Address PR feedback, update lit test
---
llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll | 26 ++++++++++++++-----
1 file changed, 19 insertions(+), 7 deletions(-)
diff --git a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
index f11e9dba2bd6b..cd108a5e6c3b5 100644
--- a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
+++ b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
@@ -1,26 +1,38 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 6
; RUN: opt -enable-scalar-pre=false -enable-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK
define void @kernel(ptr %arr, i8 %cond) {
+; CHECK-LABEL: define void @kernel(
+; CHECK-SAME: ptr [[ARR:%.*]], i8 [[COND:%.*]]) {
+; CHECK-NEXT: [[ENTRY:.*:]]
+; CHECK-NEXT: [[TOBOOL_NOT:%.*]] = icmp eq i8 [[COND]], 0
+; CHECK-NEXT: [[TMP7_PRE:%.*]] = load i32, ptr [[ARR]], align 4
+; CHECK-NEXT: br i1 [[TOBOOL_NOT]], label %[[IF_END:.*]], label %[[IF_THEN:.*]]
+; CHECK: [[IF_THEN]]:
+; CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP7_PRE]], 2
+; CHECK-NEXT: [[GETELEM:%.*]] = getelementptr inbounds nuw i8, ptr [[ARR]], i64 8
+; CHECK-NEXT: store i32 [[ADD]], ptr [[GETELEM]], align 4
+; CHECK-NEXT: br label %[[IF_END]]
+; CHECK: [[IF_END]]:
+; CHECK-NEXT: [[ADD8:%.*]] = add nsw i32 [[TMP7_PRE]], 2
+; CHECK-NEXT: [[GETELEM1:%.*]] = getelementptr inbounds nuw i8, ptr [[ARR]], i64 12
+; CHECK-NEXT: store i32 [[ADD8]], ptr [[GETELEM1]], align 4
+; CHECK-NEXT: ret void
+;
entry:
%tobool.not = icmp eq i8 %cond, 0
%tmp7.pre = load i32, ptr %arr, align 4
br i1 %tobool.not, label %if.end, label %if.then
-; CHECK: if.then:
-; CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[LOAD:%.*]], 2
-
if.then: ; preds = %entry
%add = add nsw i32 %tmp7.pre, 2
%getElem = getelementptr inbounds nuw i8, ptr %arr, i64 8
store i32 %add, ptr %getElem, align 4
br label %if.end
-; CHECK: if.end:
-; CHECK-NEXT: [[ADD2:%.*]] = add nsw i32 [[LOAD:%.*]], 2
-
if.end: ; preds = %if.then, %entry
%add8 = add nsw i32 %tmp7.pre, 2
%getElem1 = getelementptr inbounds nuw i8, ptr %arr, i64 12
store i32 %add8, ptr %getElem1, align 4
ret void
-}
\ No newline at end of file
+}
>From 1394cd1d8a54034eee5d39e867c9f431dc382b6c Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Mon, 6 Apr 2026 19:52:16 +0000
Subject: [PATCH 3/6] Update test to check that enabling scalar pre performs
the optimization and that pass parsing is working as expected
---
llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll | 23 +++++++++++++++++++
1 file changed, 23 insertions(+)
diff --git a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
index cd108a5e6c3b5..1c215d5874d37 100644
--- a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
+++ b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
@@ -1,5 +1,8 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 6
; RUN: opt -enable-scalar-pre=false -enable-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK
+; RUN: opt -enable-scalar-pre=true -enable-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK-ENABLED
+; RUN: opt -passes='gvn<no-scalar-pre>' -S < %s | FileCheck %s --check-prefixes=CHECK
+; RUN: opt -passes='gvn<scalar-pre>' -S < %s | FileCheck %s --check-prefixes=CHECK-ENABLED
define void @kernel(ptr %arr, i8 %cond) {
; CHECK-LABEL: define void @kernel(
@@ -19,6 +22,26 @@ define void @kernel(ptr %arr, i8 %cond) {
; CHECK-NEXT: store i32 [[ADD8]], ptr [[GETELEM1]], align 4
; CHECK-NEXT: ret void
;
+; CHECK-ENABLED-LABEL: define void @kernel(
+; CHECK-ENABLED-SAME: ptr [[ARR:%.*]], i8 [[COND:%.*]]) {
+; CHECK-ENABLED-NEXT: [[ENTRY:.*:]]
+; CHECK-ENABLED-NEXT: [[TOBOOL_NOT:%.*]] = icmp eq i8 [[COND]], 0
+; CHECK-ENABLED-NEXT: [[TMP7_PRE:%.*]] = load i32, ptr [[ARR]], align 4
+; CHECK-ENABLED-NEXT: br i1 [[TOBOOL_NOT]], label %[[ENTRY_IF_END_CRIT_EDGE:.*]], label %[[IF_THEN:.*]]
+; CHECK-ENABLED: [[ENTRY_IF_END_CRIT_EDGE]]:
+; CHECK-ENABLED-NEXT: [[DOTPRE:%.*]] = add nsw i32 [[TMP7_PRE]], 2
+; CHECK-ENABLED-NEXT: br label %[[IF_END:.*]]
+; CHECK-ENABLED: [[IF_THEN]]:
+; CHECK-ENABLED-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP7_PRE]], 2
+; CHECK-ENABLED-NEXT: [[GETELEM:%.*]] = getelementptr inbounds nuw i8, ptr [[ARR]], i64 8
+; CHECK-ENABLED-NEXT: store i32 [[ADD]], ptr [[GETELEM]], align 4
+; CHECK-ENABLED-NEXT: br label %[[IF_END]]
+; CHECK-ENABLED: [[IF_END]]:
+; CHECK-ENABLED-NEXT: [[ADD8_PRE_PHI:%.*]] = phi i32 [ [[DOTPRE]], %[[ENTRY_IF_END_CRIT_EDGE]] ], [ [[ADD]], %[[IF_THEN]] ]
+; CHECK-ENABLED-NEXT: [[GETELEM1:%.*]] = getelementptr inbounds nuw i8, ptr [[ARR]], i64 12
+; CHECK-ENABLED-NEXT: store i32 [[ADD8_PRE_PHI]], ptr [[GETELEM1]], align 4
+; CHECK-ENABLED-NEXT: ret void
+;
entry:
%tobool.not = icmp eq i8 %cond, 0
%tmp7.pre = load i32, ptr %arr, align 4
>From 096654b83727894d9c49a17f97e4d08947a1b470 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Wed, 29 Apr 2026 14:20:26 +0000
Subject: [PATCH 4/6] Rename PRE flag to scalar-pre and fix missing usage
---
llvm/include/llvm/Transforms/Scalar/GVN.h | 8 --------
llvm/lib/Passes/PassBuilder.cpp | 4 +---
llvm/lib/Passes/PassRegistry.def | 2 +-
llvm/lib/Transforms/Scalar/GVN.cpp | 13 +++----------
llvm/test/Other/new-pm-print-pipeline.ll | 4 ++--
llvm/test/Transforms/GVN/PRE/local-pre.ll | 4 ++--
llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll | 4 ++--
llvm/test/Transforms/GVN/PRE/pre-aliasning-path.ll | 4 ++--
llvm/test/Transforms/GVN/PRE/pre-basic-add.ll | 6 +++---
llvm/test/Transforms/GVN/PRE/pre-jt-add.ll | 4 ++--
.../test/Transforms/GVN/PRE/pre-loop-load-new-pm.ll | 4 ++--
llvm/test/Transforms/GVN/PRE/pre-loop-load.ll | 2 +-
llvm/test/Transforms/GVN/PRE/pre-poison-add.ll | 4 ++--
13 files changed, 23 insertions(+), 40 deletions(-)
diff --git a/llvm/include/llvm/Transforms/Scalar/GVN.h b/llvm/include/llvm/Transforms/Scalar/GVN.h
index ccc608eb0c42b..67bf49e01348c 100644
--- a/llvm/include/llvm/Transforms/Scalar/GVN.h
+++ b/llvm/include/llvm/Transforms/Scalar/GVN.h
@@ -76,7 +76,6 @@ class GVNLegacyPass;
/// Intended use is to create a default object, modify parameters with
/// additional setters and then pass it to GVN.
struct GVNOptions {
- std::optional<bool> AllowPRE;
std::optional<bool> AllowScalarPRE;
std::optional<bool> AllowLoadPRE;
std::optional<bool> AllowLoadInLoopPRE;
@@ -86,12 +85,6 @@ struct GVNOptions {
GVNOptions() = default;
- /// Enables or disables PRE in GVN.
- GVNOptions &setPRE(bool PRE) {
- AllowPRE = PRE;
- return *this;
- }
-
/// Enables or disables PRE of scalars in GVN.
GVNOptions &setScalarPRE(bool ScalarPRE) {
AllowScalarPRE = ScalarPRE;
@@ -155,7 +148,6 @@ class GVNPass : public PassInfoMixin<GVNPass> {
AAResults *getAliasAnalysis() const { return VN.getAliasAnalysis(); }
MemoryDependenceResults &getMemDep() const { return *MD; }
- LLVM_ABI bool isPREEnabled() const;
LLVM_ABI bool isScalarPREEnabled() const;
LLVM_ABI bool isLoadPREEnabled() const;
LLVM_ABI bool isLoadInLoopPREEnabled() const;
diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp
index 25d5ac2e1e446..a0499625931a0 100644
--- a/llvm/lib/Passes/PassBuilder.cpp
+++ b/llvm/lib/Passes/PassBuilder.cpp
@@ -1329,9 +1329,7 @@ Expected<GVNOptions> parseGVNOptions(StringRef Params) {
std::tie(ParamName, Params) = Params.split(';');
bool Enable = !ParamName.consume_front("no-");
- if (ParamName == "pre") {
- Result.setPRE(Enable);
- } else if (ParamName == "scalar-pre") {
+ if (ParamName == "scalar-pre") {
Result.setScalarPRE(Enable);
} else if (ParamName == "load-pre") {
Result.setLoadPRE(Enable);
diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def
index c64bf9bf729b0..e462d95b249a3 100644
--- a/llvm/lib/Passes/PassRegistry.def
+++ b/llvm/lib/Passes/PassRegistry.def
@@ -598,7 +598,7 @@ FUNCTION_PASS_WITH_PARAMS(
FUNCTION_PASS_WITH_PARAMS(
"gvn", "GVNPass", [](GVNOptions Opts) { return GVNPass(Opts); },
parseGVNOptions,
- "no-pre;pre;no-scalar-pre;scalar-pre;no-load-pre;load-pre;no-split-backedge-load-pre;"
+ "no-scalar-pre;scalar-pre;no-load-pre;load-pre;no-split-backedge-load-pre;"
"split-backedge-load-pre;no-memdep;memdep;no-memoryssa;memoryssa")
FUNCTION_PASS_WITH_PARAMS(
"hardware-loops", "HardwareLoopsPass",
diff --git a/llvm/lib/Transforms/Scalar/GVN.cpp b/llvm/lib/Transforms/Scalar/GVN.cpp
index 1a592ff160950..90eba1ee10f0e 100644
--- a/llvm/lib/Transforms/Scalar/GVN.cpp
+++ b/llvm/lib/Transforms/Scalar/GVN.cpp
@@ -105,7 +105,6 @@ STATISTIC(MaxBBSpeculationCutoffReachedTimes,
"Number of times we we reached gvn-max-block-speculations cut-off "
"preventing further exploration");
-static cl::opt<bool> GVNEnablePRE("enable-pre", cl::init(true), cl::Hidden);
static cl::opt<bool> GVNEnableScalarPRE("enable-scalar-pre", cl::init(true),
cl::Hidden);
static cl::opt<bool> GVNEnableLoadPRE("enable-load-pre", cl::init(true));
@@ -852,10 +851,6 @@ void GVNPass::LeaderMap::verifyRemoved(const Value *V) const {
// GVN Pass
//===----------------------------------------------------------------------===//
-bool GVNPass::isPREEnabled() const {
- return Options.AllowPRE.value_or(GVNEnablePRE);
-}
-
bool GVNPass::isScalarPREEnabled() const {
return Options.AllowScalarPRE.value_or(GVNEnableScalarPRE);
}
@@ -919,8 +914,6 @@ void GVNPass::printPipeline(
OS, MapClassName2PassName);
OS << '<';
- if (Options.AllowPRE != std::nullopt)
- OS << (*Options.AllowPRE ? "" : "no-") << "pre;";
if (Options.AllowScalarPRE != std::nullopt)
OS << (*Options.AllowScalarPRE ? "" : "no-") << "scalar-pre;";
if (Options.AllowLoadPRE != std::nullopt)
@@ -2039,7 +2032,7 @@ bool GVNPass::processNonLocalLoad(LoadInst *Load) {
bool Changed = false;
// This is a limited form of scalar PRE for load indices. If this load follows
// a GEP, see if we can PRE the indices before analyzing.
- if (isPREEnabled() && isScalarPREEnabled()) {
+ if (isScalarPREEnabled()) {
if (GetElementPtrInst *GEP =
dyn_cast<GetElementPtrInst>(Load->getOperand(0))) {
for (Use &U : GEP->indices())
@@ -2089,7 +2082,7 @@ bool GVNPass::processNonLocalLoad(LoadInst *Load) {
}
// Step 4: Eliminate partial redundancy.
- if (!isPREEnabled() || !isLoadPREEnabled())
+ if (!isLoadPREEnabled())
return Changed;
if (!isLoadInLoopPREEnabled() && LI->getLoopFor(Load->getParent()))
return Changed;
@@ -2853,7 +2846,7 @@ bool GVNPass::runImpl(Function &F, AssumptionCache &RunAC, DominatorTree &RunDT,
++Iteration;
}
- if (isPREEnabled() && isScalarPREEnabled()) {
+ if (isScalarPREEnabled()) {
// Fabricate val-num for dead-code in order to suppress assertion in
// performPRE().
assignValNumForDeadCode();
diff --git a/llvm/test/Other/new-pm-print-pipeline.ll b/llvm/test/Other/new-pm-print-pipeline.ll
index 53a17975c512b..9d0aee92e01f3 100644
--- a/llvm/test/Other/new-pm-print-pipeline.ll
+++ b/llvm/test/Other/new-pm-print-pipeline.ll
@@ -31,8 +31,8 @@
; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(loop-unroll<>,loop-unroll<partial;peeling;runtime;upperbound;profile-peeling;full-unroll-max=5;O1>,loop-unroll<no-partial;no-peeling;no-runtime;no-upperbound;no-profile-peeling;full-unroll-max=7;O1>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-10
; CHECK-10: function(loop-unroll<O2>,loop-unroll<partial;peeling;runtime;upperbound;profile-peeling;full-unroll-max=5;O1>,loop-unroll<no-partial;no-peeling;no-runtime;no-upperbound;no-profile-peeling;full-unroll-max=7;O1>)
-; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(gvn<>,gvn<pre;load-pre;split-backedge-load-pre;memdep;memoryssa>,gvn<no-pre;no-load-pre;no-split-backedge-load-pre;no-memdep;no-memoryssa>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-11
-; CHECK-11: function(gvn<>,gvn<pre;load-pre;split-backedge-load-pre;no-memdep;memoryssa>,gvn<no-pre;no-load-pre;no-split-backedge-load-pre;memdep;no-memoryssa>)
+; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(gvn<>,gvn<scalar-pre;load-pre;split-backedge-load-pre;memdep;memoryssa>,gvn<no-pre;no-load-pre;no-split-backedge-load-pre;no-memdep;no-memoryssa>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-11
+; CHECK-11: function(gvn<>,gvn<scalar-pre;load-pre;split-backedge-load-pre;no-memdep;memoryssa>,gvn<no-scalar-pre;no-load-pre;no-split-backedge-load-pre;memdep;no-memoryssa>)
; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(early-cse<>,early-cse<memssa>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-12
; CHECK-12: function(early-cse<>,early-cse<memssa>)
diff --git a/llvm/test/Transforms/GVN/PRE/local-pre.ll b/llvm/test/Transforms/GVN/PRE/local-pre.ll
index 7f465927f48b6..c67a5f1549f80 100644
--- a/llvm/test/Transforms/GVN/PRE/local-pre.ll
+++ b/llvm/test/Transforms/GVN/PRE/local-pre.ll
@@ -1,5 +1,5 @@
-; RUN: opt < %s -passes=gvn -enable-pre -S | FileCheck %s
-; RUN: opt < %s -passes="gvn<pre>" -enable-pre=false -S | FileCheck %s
+; RUN: opt < %s -passes=gvn -enable-scalar-pre -S | FileCheck %s
+; RUN: opt < %s -passes="gvn<scalar-pre>" -enable-scalar-pre=false -S | FileCheck %s
declare void @may_exit() nounwind
diff --git a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
index 1c215d5874d37..5bf79cdff649a 100644
--- a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
+++ b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 6
-; RUN: opt -enable-scalar-pre=false -enable-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK
-; RUN: opt -enable-scalar-pre=true -enable-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK-ENABLED
+; RUN: opt -enable-scalar-pre=false -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK
+; RUN: opt -enable-scalar-pre=true -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK-ENABLED
; RUN: opt -passes='gvn<no-scalar-pre>' -S < %s | FileCheck %s --check-prefixes=CHECK
; RUN: opt -passes='gvn<scalar-pre>' -S < %s | FileCheck %s --check-prefixes=CHECK-ENABLED
diff --git a/llvm/test/Transforms/GVN/PRE/pre-aliasning-path.ll b/llvm/test/Transforms/GVN/PRE/pre-aliasning-path.ll
index 60611a032ded5..45e6bd1bf8b6b 100644
--- a/llvm/test/Transforms/GVN/PRE/pre-aliasning-path.ll
+++ b/llvm/test/Transforms/GVN/PRE/pre-aliasning-path.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
-; RUN: opt -enable-load-pre -enable-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK,MDEP
-; RUN: opt -enable-load-pre -enable-pre -passes='gvn<memoryssa>' -S < %s | FileCheck %s --check-prefixes=CHECK,MSSA
+; RUN: opt -enable-load-pre -enable-scalar-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK,MDEP
+; RUN: opt -enable-load-pre -enable-scalar-pre -passes='gvn<memoryssa>' -S < %s | FileCheck %s --check-prefixes=CHECK,MSSA
declare void @side_effect_0() nofree
diff --git a/llvm/test/Transforms/GVN/PRE/pre-basic-add.ll b/llvm/test/Transforms/GVN/PRE/pre-basic-add.ll
index 9bf64962ecb1f..92306015378cb 100644
--- a/llvm/test/Transforms/GVN/PRE/pre-basic-add.ll
+++ b/llvm/test/Transforms/GVN/PRE/pre-basic-add.ll
@@ -1,7 +1,7 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 5
-; RUN: opt < %s -passes=gvn -enable-pre -S | FileCheck %s --check-prefixes=CHECK,MDEP
-; RUN: opt < %s -passes='gvn<memoryssa>' -enable-pre -S | FileCheck %s --check-prefixes=CHECK,MSSA
-; RUN: opt < %s -passes="gvn<pre>" -enable-pre=false -S | FileCheck %s
+; RUN: opt < %s -passes=gvn -enable-scalar-pre -S | FileCheck %s --check-prefixes=CHECK,MDEP
+; RUN: opt < %s -passes='gvn<memoryssa>' -enable-scalar-pre -S | FileCheck %s --check-prefixes=CHECK,MSSA
+; RUN: opt < %s -passes="gvn<scalar-pre>" -enable-scalar-pre=false -S | FileCheck %s
@H = common global i32 0 ; <ptr> [#uses=2]
@G = common global i32 0 ; <ptr> [#uses=1]
diff --git a/llvm/test/Transforms/GVN/PRE/pre-jt-add.ll b/llvm/test/Transforms/GVN/PRE/pre-jt-add.ll
index f62d06dbf0f84..c6470c1158e39 100644
--- a/llvm/test/Transforms/GVN/PRE/pre-jt-add.ll
+++ b/llvm/test/Transforms/GVN/PRE/pre-jt-add.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 5
-; RUN: opt < %s -passes=gvn,jump-threading -enable-pre -S | FileCheck %s --check-prefixes=CHECK,MDEP
-; RUN: opt < %s -passes='gvn<memoryssa>',jump-threading -enable-pre -S | FileCheck %s --check-prefixes=CHECK,MSSA
+; RUN: opt < %s -passes=gvn,jump-threading -enable-scalar-pre -S | FileCheck %s --check-prefixes=CHECK,MDEP
+; RUN: opt < %s -passes='gvn<memoryssa>',jump-threading -enable-scalar-pre -S | FileCheck %s --check-prefixes=CHECK,MSSA
@H = common global i32 0
@G = common global i32 0
diff --git a/llvm/test/Transforms/GVN/PRE/pre-loop-load-new-pm.ll b/llvm/test/Transforms/GVN/PRE/pre-loop-load-new-pm.ll
index 4cd2e47b3c31f..0472500282f42 100644
--- a/llvm/test/Transforms/GVN/PRE/pre-loop-load-new-pm.ll
+++ b/llvm/test/Transforms/GVN/PRE/pre-loop-load-new-pm.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
-; RUN: opt -aa-pipeline=basic-aa -enable-load-pre -enable-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK,MDEP
-; RUN: opt -aa-pipeline=basic-aa -enable-load-pre -enable-pre -passes='gvn<memoryssa>' -S < %s | FileCheck %s --check-prefixes=CHECK,MSSA
+; RUN: opt -aa-pipeline=basic-aa -enable-load-pre -enable-scalar-pre -passes=gvn -S < %s | FileCheck %s --check-prefixes=CHECK,MDEP
+; RUN: opt -aa-pipeline=basic-aa -enable-load-pre -enable-scalar-pre -passes='gvn<memoryssa>' -S < %s | FileCheck %s --check-prefixes=CHECK,MSSA
declare void @side_effect()
declare i1 @side_effect_cond()
diff --git a/llvm/test/Transforms/GVN/PRE/pre-loop-load.ll b/llvm/test/Transforms/GVN/PRE/pre-loop-load.ll
index a9b69a49005e3..9cd6274100c2a 100644
--- a/llvm/test/Transforms/GVN/PRE/pre-loop-load.ll
+++ b/llvm/test/Transforms/GVN/PRE/pre-loop-load.ll
@@ -1,5 +1,5 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
-; RUN: opt -enable-load-pre -enable-pre -passes=lcssa,gvn -S < %s | FileCheck %s
+; RUN: opt -enable-load-pre -enable-scalar-pre -passes=lcssa,gvn -S < %s | FileCheck %s
declare void @side_effect() nofree
declare i1 @side_effect_cond() nofree
diff --git a/llvm/test/Transforms/GVN/PRE/pre-poison-add.ll b/llvm/test/Transforms/GVN/PRE/pre-poison-add.ll
index 32f149b881d72..a4ee356628f11 100644
--- a/llvm/test/Transforms/GVN/PRE/pre-poison-add.ll
+++ b/llvm/test/Transforms/GVN/PRE/pre-poison-add.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 5
-; RUN: opt < %s -passes=gvn -enable-pre -S | FileCheck %s --check-prefixes=CHECK,MDEP
-; RUN: opt < %s -passes='gvn<memoryssa>' -enable-pre -S | FileCheck %s --check-prefixes=CHECK,MSSA
+; RUN: opt < %s -passes=gvn -enable-scalar-pre -S | FileCheck %s --check-prefixes=CHECK,MDEP
+; RUN: opt < %s -passes='gvn<memoryssa>' -enable-scalar-pre -S | FileCheck %s --check-prefixes=CHECK,MSSA
@H = common global i32 0
@G = common global i32 0
>From 46886edcd461bbe8f6082e23cf9a340740c22d5b Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Thu, 30 Apr 2026 17:57:28 +0000
Subject: [PATCH 5/6] Add test for register pressure
---
.../NVPTX/gvn-scalar-pre-reg-pressure.ll | 33 +++++++++++++++++++
1 file changed, 33 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/gvn-scalar-pre-reg-pressure.ll
diff --git a/llvm/test/CodeGen/NVPTX/gvn-scalar-pre-reg-pressure.ll b/llvm/test/CodeGen/NVPTX/gvn-scalar-pre-reg-pressure.ll
new file mode 100644
index 0000000000000..7813e5422d8ca
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/gvn-scalar-pre-reg-pressure.ll
@@ -0,0 +1,33 @@
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -O3 | FileCheck %s --check-prefix=PIPELINE
+; RUN: opt < %s -passes=gvn -enable-scalar-pre=false -S | llc -mtriple=nvptx64 -mcpu=sm_100 -O0 | FileCheck %s --check-prefix=NO-SCALAR-PRE
+; RUN: opt < %s -passes=gvn -enable-scalar-pre=true -S | llc -mtriple=nvptx64 -mcpu=sm_100 -O0 | FileCheck %s --check-prefix=SCALAR-PRE
+
+; Scalar PRE inserts a critical-edge computation and a PHI for the common add.
+; That shape needs more NVPTX virtual registers than keeping the duplicated adds.
+
+define void @kernel(ptr %arr, i8 %cond) {
+; PIPELINE-LABEL: kernel(
+; PIPELINE: .reg .b32 {{%r<3>;}}
+;
+; NO-SCALAR-PRE-LABEL: kernel(
+; NO-SCALAR-PRE: .reg .b32 {{%r<4>;}}
+;
+; SCALAR-PRE-LABEL: kernel(
+; SCALAR-PRE: .reg .b32 {{%r<6>;}}
+entry:
+ %tobool.not = icmp eq i8 %cond, 0
+ %tmp7.pre = load i32, ptr %arr, align 4
+ br i1 %tobool.not, label %if.end, label %if.then
+
+if.then:
+ %add = add nsw i32 %tmp7.pre, 2
+ %getElem = getelementptr inbounds nuw i8, ptr %arr, i64 8
+ store i32 %add, ptr %getElem, align 4
+ br label %if.end
+
+if.end:
+ %add8 = add nsw i32 %tmp7.pre, 2
+ %getElem1 = getelementptr inbounds nuw i8, ptr %arr, i64 12
+ store i32 %add8, ptr %getElem1, align 4
+ ret void
+}
>From a985c2022549bd606cb21cdeee8f14cf80d6ef6d Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Thu, 7 May 2026 14:21:29 +0000
Subject: [PATCH 6/6] Address feedback
---
.../NVPTX/gvn-scalar-pre-reg-pressure.ll | 81 ++++++++++++++++---
llvm/test/Other/new-pm-print-pipeline.ll | 2 +-
llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll | 6 +-
3 files changed, 76 insertions(+), 13 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/gvn-scalar-pre-reg-pressure.ll b/llvm/test/CodeGen/NVPTX/gvn-scalar-pre-reg-pressure.ll
index 7813e5422d8ca..5b7f893889744 100644
--- a/llvm/test/CodeGen/NVPTX/gvn-scalar-pre-reg-pressure.ll
+++ b/llvm/test/CodeGen/NVPTX/gvn-scalar-pre-reg-pressure.ll
@@ -1,19 +1,82 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -O3 | FileCheck %s --check-prefix=PIPELINE
-; RUN: opt < %s -passes=gvn -enable-scalar-pre=false -S | llc -mtriple=nvptx64 -mcpu=sm_100 -O0 | FileCheck %s --check-prefix=NO-SCALAR-PRE
-; RUN: opt < %s -passes=gvn -enable-scalar-pre=true -S | llc -mtriple=nvptx64 -mcpu=sm_100 -O0 | FileCheck %s --check-prefix=SCALAR-PRE
+; RUN: opt < %s -passes='gvn<no-scalar-pre>' -S | llc -mtriple=nvptx64 -mcpu=sm_100 -O0 | FileCheck %s --check-prefix=NO-SCALAR-PRE
+; RUN: opt < %s -passes='gvn<scalar-pre>' -S | llc -mtriple=nvptx64 -mcpu=sm_100 -O0 | FileCheck %s --check-prefix=SCALAR-PRE
; Scalar PRE inserts a critical-edge computation and a PHI for the common add.
; That shape needs more NVPTX virtual registers than keeping the duplicated adds.
-define void @kernel(ptr %arr, i8 %cond) {
-; PIPELINE-LABEL: kernel(
-; PIPELINE: .reg .b32 {{%r<3>;}}
+define void @test_scalar_pre_option(ptr %arr, i8 %cond) {
+; PIPELINE-LABEL: test_scalar_pre_option(
+; PIPELINE: {
+; PIPELINE-NEXT: .reg .pred %p<2>;
+; PIPELINE-NEXT: .reg .b16 %rs<2>;
+; PIPELINE-NEXT: .reg .b32 %r<3>;
+; PIPELINE-NEXT: .reg .b64 %rd<2>;
+; PIPELINE-EMPTY:
+; PIPELINE-NEXT: // %bb.0: // %entry
+; PIPELINE-NEXT: ld.param.b64 %rd1, [test_scalar_pre_option_param_0];
+; PIPELINE-NEXT: ld.param.b8 %rs1, [test_scalar_pre_option_param_1];
+; PIPELINE-NEXT: setp.eq.b16 %p1, %rs1, 0;
+; PIPELINE-NEXT: ld.b32 %r2, [%rd1];
+; PIPELINE-NEXT: add.s32 %r1, %r2, 2;
+; PIPELINE-NEXT: @%p1 bra $L__BB0_2;
+; PIPELINE-NEXT: // %bb.1: // %if.then
+; PIPELINE-NEXT: st.b32 [%rd1+8], %r1;
+; PIPELINE-NEXT: $L__BB0_2: // %if.end
+; PIPELINE-NEXT: st.b32 [%rd1+12], %r1;
+; PIPELINE-NEXT: ret;
;
-; NO-SCALAR-PRE-LABEL: kernel(
-; NO-SCALAR-PRE: .reg .b32 {{%r<4>;}}
+; NO-SCALAR-PRE-LABEL: test_scalar_pre_option(
+; NO-SCALAR-PRE: {
+; NO-SCALAR-PRE-NEXT: .reg .pred %p<2>;
+; NO-SCALAR-PRE-NEXT: .reg .b16 %rs<2>;
+; NO-SCALAR-PRE-NEXT: .reg .b32 %r<4>;
+; NO-SCALAR-PRE-NEXT: .reg .b64 %rd<2>;
+; NO-SCALAR-PRE-EMPTY:
+; NO-SCALAR-PRE-NEXT: // %bb.0: // %entry
+; NO-SCALAR-PRE-NEXT: ld.param.b8 %rs1, [test_scalar_pre_option_param_1];
+; NO-SCALAR-PRE-NEXT: ld.param.b64 %rd1, [test_scalar_pre_option_param_0];
+; NO-SCALAR-PRE-NEXT: setp.eq.b16 %p1, %rs1, 0;
+; NO-SCALAR-PRE-NEXT: ld.b32 %r1, [%rd1];
+; NO-SCALAR-PRE-NEXT: @%p1 bra $L__BB0_2;
+; NO-SCALAR-PRE-NEXT: bra.uni $L__BB0_1;
+; NO-SCALAR-PRE-NEXT: $L__BB0_1: // %if.then
+; NO-SCALAR-PRE-NEXT: add.s32 %r2, %r1, 2;
+; NO-SCALAR-PRE-NEXT: st.b32 [%rd1+8], %r2;
+; NO-SCALAR-PRE-NEXT: bra.uni $L__BB0_2;
+; NO-SCALAR-PRE-NEXT: $L__BB0_2: // %if.end
+; NO-SCALAR-PRE-NEXT: add.s32 %r3, %r1, 2;
+; NO-SCALAR-PRE-NEXT: st.b32 [%rd1+12], %r3;
+; NO-SCALAR-PRE-NEXT: ret;
;
-; SCALAR-PRE-LABEL: kernel(
-; SCALAR-PRE: .reg .b32 {{%r<6>;}}
+; SCALAR-PRE-LABEL: test_scalar_pre_option(
+; SCALAR-PRE: {
+; SCALAR-PRE-NEXT: .reg .pred %p<2>;
+; SCALAR-PRE-NEXT: .reg .b16 %rs<2>;
+; SCALAR-PRE-NEXT: .reg .b32 %r<6>;
+; SCALAR-PRE-NEXT: .reg .b64 %rd<2>;
+; SCALAR-PRE-EMPTY:
+; SCALAR-PRE-NEXT: // %bb.0: // %entry
+; SCALAR-PRE-NEXT: ld.param.b8 %rs1, [test_scalar_pre_option_param_1];
+; SCALAR-PRE-NEXT: ld.param.b64 %rd1, [test_scalar_pre_option_param_0];
+; SCALAR-PRE-NEXT: setp.ne.b16 %p1, %rs1, 0;
+; SCALAR-PRE-NEXT: ld.b32 %r1, [%rd1];
+; SCALAR-PRE-NEXT: @%p1 bra $L__BB0_2;
+; SCALAR-PRE-NEXT: bra.uni $L__BB0_1;
+; SCALAR-PRE-NEXT: $L__BB0_1: // %entry.if.end_crit_edge
+; SCALAR-PRE-NEXT: add.s32 %r2, %r1, 2;
+; SCALAR-PRE-NEXT: mov.b32 %r5, %r2;
+; SCALAR-PRE-NEXT: bra.uni $L__BB0_3;
+; SCALAR-PRE-NEXT: $L__BB0_2: // %if.then
+; SCALAR-PRE-NEXT: add.s32 %r3, %r1, 2;
+; SCALAR-PRE-NEXT: st.b32 [%rd1+8], %r3;
+; SCALAR-PRE-NEXT: mov.b32 %r5, %r3;
+; SCALAR-PRE-NEXT: bra.uni $L__BB0_3;
+; SCALAR-PRE-NEXT: $L__BB0_3: // %if.end
+; SCALAR-PRE-NEXT: mov.b32 %r4, %r5;
+; SCALAR-PRE-NEXT: st.b32 [%rd1+12], %r4;
+; SCALAR-PRE-NEXT: ret;
entry:
%tobool.not = icmp eq i8 %cond, 0
%tmp7.pre = load i32, ptr %arr, align 4
diff --git a/llvm/test/Other/new-pm-print-pipeline.ll b/llvm/test/Other/new-pm-print-pipeline.ll
index 9d0aee92e01f3..21da3eb33dd6d 100644
--- a/llvm/test/Other/new-pm-print-pipeline.ll
+++ b/llvm/test/Other/new-pm-print-pipeline.ll
@@ -31,7 +31,7 @@
; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(loop-unroll<>,loop-unroll<partial;peeling;runtime;upperbound;profile-peeling;full-unroll-max=5;O1>,loop-unroll<no-partial;no-peeling;no-runtime;no-upperbound;no-profile-peeling;full-unroll-max=7;O1>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-10
; CHECK-10: function(loop-unroll<O2>,loop-unroll<partial;peeling;runtime;upperbound;profile-peeling;full-unroll-max=5;O1>,loop-unroll<no-partial;no-peeling;no-runtime;no-upperbound;no-profile-peeling;full-unroll-max=7;O1>)
-; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(gvn<>,gvn<scalar-pre;load-pre;split-backedge-load-pre;memdep;memoryssa>,gvn<no-pre;no-load-pre;no-split-backedge-load-pre;no-memdep;no-memoryssa>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-11
+; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(gvn<>,gvn<scalar-pre;load-pre;split-backedge-load-pre;memdep;memoryssa>,gvn<no-scalar-pre;no-load-pre;no-split-backedge-load-pre;no-memdep;no-memoryssa>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-11
; CHECK-11: function(gvn<>,gvn<scalar-pre;load-pre;split-backedge-load-pre;no-memdep;memoryssa>,gvn<no-scalar-pre;no-load-pre;no-split-backedge-load-pre;memdep;no-memoryssa>)
; RUN: opt -disable-output -disable-verify -print-pipeline-passes -passes='function(early-cse<>,early-cse<memssa>)' < %s | FileCheck %s --match-full-lines --check-prefixes=CHECK-12
diff --git a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
index 5bf79cdff649a..32eaa9a0e8a27 100644
--- a/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
+++ b/llvm/test/Transforms/GVN/PRE/no-scalar-pre.ll
@@ -4,8 +4,8 @@
; RUN: opt -passes='gvn<no-scalar-pre>' -S < %s | FileCheck %s --check-prefixes=CHECK
; RUN: opt -passes='gvn<scalar-pre>' -S < %s | FileCheck %s --check-prefixes=CHECK-ENABLED
-define void @kernel(ptr %arr, i8 %cond) {
-; CHECK-LABEL: define void @kernel(
+define void @test_scalar_pre_option(ptr %arr, i8 %cond) {
+; CHECK-LABEL: define void @test_scalar_pre_option(
; CHECK-SAME: ptr [[ARR:%.*]], i8 [[COND:%.*]]) {
; CHECK-NEXT: [[ENTRY:.*:]]
; CHECK-NEXT: [[TOBOOL_NOT:%.*]] = icmp eq i8 [[COND]], 0
@@ -22,7 +22,7 @@ define void @kernel(ptr %arr, i8 %cond) {
; CHECK-NEXT: store i32 [[ADD8]], ptr [[GETELEM1]], align 4
; CHECK-NEXT: ret void
;
-; CHECK-ENABLED-LABEL: define void @kernel(
+; CHECK-ENABLED-LABEL: define void @test_scalar_pre_option(
; CHECK-ENABLED-SAME: ptr [[ARR:%.*]], i8 [[COND:%.*]]) {
; CHECK-ENABLED-NEXT: [[ENTRY:.*:]]
; CHECK-ENABLED-NEXT: [[TOBOOL_NOT:%.*]] = icmp eq i8 [[COND]], 0
More information about the llvm-commits
mailing list