30
30
#include " compiler/utils/work_item_loops_pass.h"
31
31
#include " vecz/pass.h"
32
32
#include " vecz/vecz_target_info.h"
33
+ #include " llvm/IR/PassManager.h"
34
+ #include " llvm/Transforms/AggressiveInstCombine/AggressiveInstCombine.h"
33
35
#include " llvm/Transforms/IPO/AlwaysInliner.h"
36
+ #include " llvm/Transforms/Scalar/DCE.h"
37
+ #include " llvm/Transforms/Scalar/GVN.h"
38
+ #include " llvm/Transforms/Scalar/SROA.h"
39
+ #include " llvm/Transforms/Scalar/SimplifyCFG.h"
34
40
#endif
35
41
36
42
using namespace llvm ;
@@ -65,6 +71,7 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
65
71
MPM.addPass (SpecConstantsPass (SpecConstantsPass::HandlingMode::emulation));
66
72
MPM.addPass (ConvertToMuxBuiltinsSYCLNativeCPUPass ());
67
73
#ifdef NATIVECPU_USE_OCK
74
+ MPM.addPass (compiler::utils::PrepareBarriersPass ());
68
75
MPM.addPass (compiler::utils::TransferKernelMetadataPass ());
69
76
MPM.addPass (FixABIMuxBuiltinsPass ());
70
77
// Always enable vectorizer, unless explictly disabled or -O0 is set.
@@ -86,13 +93,20 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
86
93
MAM.registerPass (
87
94
[QueryFunc] { return vecz::VeczPassOptionsAnalysis (QueryFunc); });
88
95
MPM.addPass (vecz::RunVeczPass ());
96
+ FunctionPassManager FPM;
97
+ FPM.addPass (SimplifyCFGPass ());
98
+ FPM.addPass (SROAPass (SROAOptions::ModifyCFG));
99
+ FPM.addPass (AggressiveInstCombinePass ());
100
+ FPM.addPass (GVNPass (GVNOptions ().setMemDep (true )));
101
+ FPM.addPass (DCEPass ());
102
+ FPM.addPass (SimplifyCFGPass ());
103
+ MPM.addPass (createModuleToFunctionPassAdaptor (std::move (FPM)));
89
104
}
90
105
compiler::utils::WorkItemLoopsPassOptions Opts;
91
106
Opts.IsDebug = IsDebug;
92
107
Opts.ForceNoTail = ForceNoTail;
93
108
MAM.registerPass ([] { return compiler::utils::BuiltinInfoAnalysis (); });
94
109
MAM.registerPass ([] { return compiler::utils::SubgroupAnalysis (); });
95
- MPM.addPass (compiler::utils::PrepareBarriersPass ());
96
110
MPM.addPass (compiler::utils::WorkItemLoopsPass (Opts));
97
111
MPM.addPass (compiler::utils::ReplaceLocalModuleScopeVariablesPass ());
98
112
MPM.addPass (AlwaysInlinerPass ());
0 commit comments