3030#include " compiler/utils/work_item_loops_pass.h"
3131#include " vecz/pass.h"
3232#include " vecz/vecz_target_info.h"
33+ #include " llvm/IR/PassManager.h"
34+ #include " llvm/Transforms/AggressiveInstCombine/AggressiveInstCombine.h"
3335#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"
3440#endif
3541
3642using namespace llvm ;
@@ -65,6 +71,7 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
6571 MPM.addPass (SpecConstantsPass (SpecConstantsPass::HandlingMode::emulation));
6672 MPM.addPass (ConvertToMuxBuiltinsSYCLNativeCPUPass ());
6773#ifdef NATIVECPU_USE_OCK
74+ MPM.addPass (compiler::utils::PrepareBarriersPass ());
6875 MPM.addPass (compiler::utils::TransferKernelMetadataPass ());
6976 MPM.addPass (FixABIMuxBuiltinsPass ());
7077 // Always enable vectorizer, unless explictly disabled or -O0 is set.
@@ -86,13 +93,20 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
8693 MAM.registerPass (
8794 [QueryFunc] { return vecz::VeczPassOptionsAnalysis (QueryFunc); });
8895 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)));
89104 }
90105 compiler::utils::WorkItemLoopsPassOptions Opts;
91106 Opts.IsDebug = IsDebug;
92107 Opts.ForceNoTail = ForceNoTail;
93108 MAM.registerPass ([] { return compiler::utils::BuiltinInfoAnalysis (); });
94109 MAM.registerPass ([] { return compiler::utils::SubgroupAnalysis (); });
95- MPM.addPass (compiler::utils::PrepareBarriersPass ());
96110 MPM.addPass (compiler::utils::WorkItemLoopsPass (Opts));
97111 MPM.addPass (compiler::utils::ReplaceLocalModuleScopeVariablesPass ());
98112 MPM.addPass (AlwaysInlinerPass ());
0 commit comments