Skip to content

Commit 6f324d3

Browse files
maksimsabmaksimsab
authored
[ESIMD][NFC] Extract lowerEsimdConstructs into the SYCLPostLink library. (#18587)
This patch allows later functionality reusing in clang-linker-wrapper and sycl-jit.
1 parent 0377e42 commit 6f324d3

5 files changed

Lines changed: 118 additions & 74 deletions

File tree

‎llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h‎

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
//===--------- ESIMDPostSplitProcessing.h - Post split ESIMD Processing ---===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// Post split ESIMD processing contains of lowering ESIMD constructs and
9+
// required optimimizations.
10+
//===----------------------------------------------------------------------===//
11+
12+
#include "llvm/SYCLPostLink/ModuleSplitter.h"
13+
14+
namespace llvm {
15+
namespace sycl {
16+
17+
/// Lowers ESIMD constructs after separation from regular SYCL code.
18+
/// \SplitESIMD identifies that ESIMD splitting is requested in the compilation.
19+
/// Returns true if the given \MD has been modified.
20+
bool lowerESIMDConstructs(llvm::module_split::ModuleDesc &MD, bool OptLevelO0,
21+
bool SplitESIMD);
22+
23+
} // namespace sycl
24+
} // namespace llvm

‎llvm/lib/SYCLPostLink/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,6 @@
11
add_llvm_component_library(LLVMSYCLPostLink
22
ComputeModuleRuntimeInfo.cpp
3+
ESIMDPostSplitProcessing.cpp
34
ModuleSplitter.cpp
45

56
ADDITIONAL_HEADER_DIRS

‎llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp‎

Lines changed: 91 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,91 @@
1+
//===-------- ESIMDPostSplitProcessing.cpp -------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// See comments in the header.
9+
//===----------------------------------------------------------------------===//
10+
11+
#include "llvm/SYCLPostLink/ESIMDPostSplitProcessing.h"
12+
13+
#include "llvm/GenXIntrinsics/GenXSPIRVWriterAdaptor.h"
14+
#include "llvm/Passes/PassBuilder.h"
15+
#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
16+
#include "llvm/SYCLPostLink/ModuleSplitter.h"
17+
#include "llvm/Transforms/IPO/AlwaysInliner.h"
18+
#include "llvm/Transforms/IPO/StripDeadPrototypes.h"
19+
#include "llvm/Transforms/InstCombine/InstCombine.h"
20+
#include "llvm/Transforms/Scalar.h"
21+
#include "llvm/Transforms/Scalar/DCE.h"
22+
#include "llvm/Transforms/Scalar/EarlyCSE.h"
23+
#include "llvm/Transforms/Scalar/SROA.h"
24+
25+
#include <string>
26+
#include <vector>
27+
28+
using namespace llvm;
29+
using namespace llvm::module_split;
30+
31+
namespace {
32+
33+
ModulePassManager buildESIMDLoweringPipeline(bool OptLevelO0, bool SplitESIMD) {
34+
ModulePassManager MPM;
35+
MPM.addPass(SYCLLowerESIMDPass(!SplitESIMD));
36+
37+
if (!OptLevelO0) {
38+
FunctionPassManager FPM;
39+
FPM.addPass(SROAPass(SROAOptions::ModifyCFG));
40+
MPM.addPass(createModuleToFunctionPassAdaptor(std::move(FPM)));
41+
}
42+
MPM.addPass(ESIMDOptimizeVecArgCallConvPass{});
43+
FunctionPassManager MainFPM;
44+
MainFPM.addPass(ESIMDLowerLoadStorePass{});
45+
46+
if (!OptLevelO0) {
47+
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
48+
MainFPM.addPass(EarlyCSEPass(true));
49+
MainFPM.addPass(InstCombinePass{});
50+
MainFPM.addPass(DCEPass{});
51+
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
52+
MainFPM.addPass(EarlyCSEPass(true));
53+
MainFPM.addPass(InstCombinePass{});
54+
MainFPM.addPass(DCEPass{});
55+
}
56+
MPM.addPass(ESIMDLowerSLMReservationCalls{});
57+
MPM.addPass(createModuleToFunctionPassAdaptor(std::move(MainFPM)));
58+
MPM.addPass(GenXSPIRVWriterAdaptor(/*RewriteTypes=*/true,
59+
/*RewriteSingleElementVectorsIn*/ false));
60+
return MPM;
61+
}
62+
63+
} // anonymous namespace
64+
65+
// When ESIMD code was separated from the regular SYCL code,
66+
// we can safely process ESIMD part.
67+
bool sycl::lowerESIMDConstructs(ModuleDesc &MD, bool OptLevelO0,
68+
bool SplitESIMD) {
69+
// TODO: support options like -debug-pass, -print-[before|after], and others
70+
LoopAnalysisManager LAM;
71+
CGSCCAnalysisManager CGAM;
72+
FunctionAnalysisManager FAM;
73+
ModuleAnalysisManager MAM;
74+
75+
PassBuilder PB;
76+
PB.registerModuleAnalyses(MAM);
77+
PB.registerCGSCCAnalyses(CGAM);
78+
PB.registerFunctionAnalyses(FAM);
79+
PB.registerLoopAnalyses(LAM);
80+
PB.crossRegisterProxies(LAM, FAM, CGAM, MAM);
81+
82+
std::vector<std::string> Names;
83+
MD.saveEntryPointNames(Names);
84+
ModulePassManager MPM = buildESIMDLoweringPipeline(OptLevelO0, SplitESIMD);
85+
PreservedAnalyses Res = MPM.run(MD.getModule(), MAM);
86+
87+
// GenXSPIRVWriterAdaptor pass replaced some functions with "rewritten"
88+
// versions so the entry point table must be rebuilt.
89+
MD.rebuildEntryPoints(Names);
90+
return !Res.areAllPreserved();
91+
}

‎llvm/tools/sycl-post-link/CMakeLists.txt‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,6 @@ set(LLVM_LINK_COMPONENTS
99
TransformUtils
1010
SYCLLowerIR
1111
SYCLPostLink
12-
InstCombine
13-
ScalarOpts
1412
Linker
1513
Passes
1614
Analysis

‎llvm/tools/sycl-post-link/sycl-post-link.cpp‎

Lines changed: 2 additions & 72 deletions
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,6 @@
2121
#include "llvm/Analysis/TargetTransformInfo.h"
2222
#include "llvm/Bitcode/BitcodeWriterPass.h"
2323
#include "llvm/Demangle/Demangle.h"
24-
#include "llvm/GenXIntrinsics/GenXSPIRVWriterAdaptor.h"
2524
#include "llvm/IR/Dominators.h"
2625
#include "llvm/IR/LLVMContext.h"
2726
#include "llvm/IR/Module.h"
@@ -32,14 +31,14 @@
3231
#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
3332
#include "llvm/SYCLLowerIR/DeviceConfigFile.hpp"
3433
#include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
35-
#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
3634
#include "llvm/SYCLLowerIR/HostPipes.h"
3735
#include "llvm/SYCLLowerIR/LowerInvokeSimd.h"
3836
#include "llvm/SYCLLowerIR/SYCLJointMatrixTransform.h"
3937
#include "llvm/SYCLLowerIR/SYCLUtils.h"
4038
#include "llvm/SYCLLowerIR/SpecConstants.h"
4139
#include "llvm/SYCLLowerIR/Support.h"
4240
#include "llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h"
41+
#include "llvm/SYCLPostLink/ESIMDPostSplitProcessing.h"
4342
#include "llvm/SYCLPostLink/ModuleSplitter.h"
4443
#include "llvm/Support/CommandLine.h"
4544
#include "llvm/Support/FileSystem.h"
@@ -49,13 +48,7 @@
4948
#include "llvm/Support/SourceMgr.h"
5049
#include "llvm/Support/SystemUtils.h"
5150
#include "llvm/Support/WithColor.h"
52-
#include "llvm/Transforms/IPO/AlwaysInliner.h"
5351
#include "llvm/Transforms/IPO/StripDeadPrototypes.h"
54-
#include "llvm/Transforms/InstCombine/InstCombine.h"
55-
#include "llvm/Transforms/Scalar.h"
56-
#include "llvm/Transforms/Scalar/DCE.h"
57-
#include "llvm/Transforms/Scalar/EarlyCSE.h"
58-
#include "llvm/Transforms/Scalar/SROA.h"
5952

6053
#include <algorithm>
6154
#include <memory>
@@ -361,69 +354,6 @@ std::string saveModuleSymbolTable(const module_split::ModuleDesc &MD, int I,
361354
return OutFileName;
362355
}
363356

364-
template <class PassClass> bool runModulePass(Module &M) {
365-
ModulePassManager MPM;
366-
ModuleAnalysisManager MAM;
367-
// Register required analysis
368-
MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
369-
MPM.addPass(PassClass{});
370-
PreservedAnalyses Res = MPM.run(M, MAM);
371-
return !Res.areAllPreserved();
372-
}
373-
374-
// When ESIMD code was separated from the regular SYCL code,
375-
// we can safely process ESIMD part.
376-
// TODO: support options like -debug-pass, -print-[before|after], and others
377-
bool lowerEsimdConstructs(module_split::ModuleDesc &MD) {
378-
LoopAnalysisManager LAM;
379-
CGSCCAnalysisManager CGAM;
380-
FunctionAnalysisManager FAM;
381-
ModuleAnalysisManager MAM;
382-
383-
PassBuilder PB;
384-
PB.registerModuleAnalyses(MAM);
385-
PB.registerCGSCCAnalyses(CGAM);
386-
PB.registerFunctionAnalyses(FAM);
387-
PB.registerLoopAnalyses(LAM);
388-
PB.crossRegisterProxies(LAM, FAM, CGAM, MAM);
389-
390-
ModulePassManager MPM;
391-
MPM.addPass(SYCLLowerESIMDPass(!SplitEsimd));
392-
393-
if (!OptLevelO0) {
394-
FunctionPassManager FPM;
395-
FPM.addPass(SROAPass(SROAOptions::ModifyCFG));
396-
MPM.addPass(createModuleToFunctionPassAdaptor(std::move(FPM)));
397-
}
398-
MPM.addPass(ESIMDOptimizeVecArgCallConvPass{});
399-
FunctionPassManager MainFPM;
400-
MainFPM.addPass(ESIMDLowerLoadStorePass{});
401-
402-
if (!OptLevelO0) {
403-
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
404-
MainFPM.addPass(EarlyCSEPass(true));
405-
MainFPM.addPass(InstCombinePass{});
406-
MainFPM.addPass(DCEPass{});
407-
// TODO: maybe remove some passes below that don't affect code quality
408-
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
409-
MainFPM.addPass(EarlyCSEPass(true));
410-
MainFPM.addPass(InstCombinePass{});
411-
MainFPM.addPass(DCEPass{});
412-
}
413-
MPM.addPass(ESIMDLowerSLMReservationCalls{});
414-
MPM.addPass(createModuleToFunctionPassAdaptor(std::move(MainFPM)));
415-
MPM.addPass(GenXSPIRVWriterAdaptor(/*RewriteTypes=*/true,
416-
/*RewriteSingleElementVectorsIn*/ false));
417-
// GenXSPIRVWriterAdaptor pass replaced some functions with "rewritten"
418-
// versions so the entry point table must be rebuilt.
419-
// TODO Change entry point search to analysis?
420-
std::vector<std::string> Names;
421-
MD.saveEntryPointNames(Names);
422-
PreservedAnalyses Res = MPM.run(MD.getModule(), MAM);
423-
MD.rebuildEntryPoints(Names);
424-
return !Res.areAllPreserved();
425-
}
426-
427357
// Compute the filename suffix for the module
428358
StringRef getModuleSuffix(const module_split::ModuleDesc &MD) {
429359
return MD.isESIMD() ? "_esimd" : "";
@@ -610,7 +540,7 @@ handleESIMD(module_split::ModuleDesc &&MDesc, bool &Modified,
610540
for (auto &MD : Result) {
611541
DUMP_ENTRY_POINTS(MD.entries(), MD.Name.c_str(), 3);
612542
if (LowerEsimd && MD.isESIMD())
613-
Modified |= lowerEsimdConstructs(MD);
543+
Modified |= sycl::lowerESIMDConstructs(MD, OptLevelO0, SplitEsimd);
614544
}
615545

616546
if (!SplitEsimd && Result.size() > 1) {

0 commit comments

Comments
 (0)