diff --git a/llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h b/llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h
index a69e7d645abda..b0364f79397a9 100644
--- a/llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h
+++ b/llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h
@@ -9,16 +9,56 @@
// required optimimizations.
//===----------------------------------------------------------------------===//
+#ifndef LLVM_SYCL_POST_LINK_ESIMD_POST_SPLIT_PROCESSING_H
+#define LLVM_SYCL_POST_LINK_ESIMD_POST_SPLIT_PROCESSING_H
+
#include "llvm/SYCLPostLink/ModuleSplitter.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/Support/Error.h"
+
namespace llvm {
namespace sycl {
+struct ESIMDProcessingOptions {
+ llvm::module_split::IRSplitMode SplitMode =
+ llvm::module_split::IRSplitMode::SPLIT_NONE;
+ bool EmitOnlyKernelsAsEntryPoints = false;
+ bool AllowDeviceImageDependencies = false;
+ bool LowerESIMD = false;
+ bool SplitESIMD = false;
+ unsigned OptLevel = 0;
+};
+
/// Lowers ESIMD constructs after separation from regular SYCL code.
-/// \SplitESIMD identifies that ESIMD splitting is requested in the compilation.
-/// Returns true if the given \MD has been modified.
-bool lowerESIMDConstructs(llvm::module_split::ModuleDesc &MD, bool OptLevelO0,
- bool SplitESIMD);
+/// \p Options.SplitESIMD identifies that ESIMD splitting is requested in the
+/// compilation. Returns true if the given \p MD has been modified.
+bool lowerESIMDConstructs(llvm::module_split::ModuleDesc &MD,
+ const ESIMDProcessingOptions &Options);
+
+/// Performs ESIMD processing that happens in the following steps:
+/// 1) Separate ESIMD Module from SYCL code.
+/// \p Options.EmitOnlyKernelsAsEntryPoints and
+/// \p Options.AllowDeviceImageDependencies are being used in the splitting.
+/// 2) If \p Options.LowerESIMD is true then ESIMD lowering pipeline is applied
+/// to the ESIMD Module.
+/// If \p Options.OptLevel is not O0 then ESIMD Module is being optimized
+/// after the lowering.
+/// 3.1) If \p Options.SplitESIMD is true then both ESIMD and non-ESIMD modules
+/// are returned.
+/// 3.2) Otherwise, two Modules are being linked into one Module which is
+/// returned. After the linking graphs become disjoint because functions
+/// shared between graphs are cloned and renamed.
+///
+/// \p Modified value indicates whether the Module has been modified.
+/// \p SplitOccurred value indicates whether split has occurred before or during
+/// function's invocation.
+Expected<SmallVector<module_split::ModuleDesc, 2>>
+handleESIMD(llvm::module_split::ModuleDesc MDesc,
+ const ESIMDProcessingOptions &Options, bool &Modified,
+ bool &SplitOccurred);
} // namespace sycl
} // namespace llvm
+
+#endif // LLVM_SYCL_POST_LINK_ESIMD_POST_SPLIT_PROCESSING_H
diff --git a/llvm/lib/SYCLPostLink/CMakeLists.txt b/llvm/lib/SYCLPostLink/CMakeLists.txt
index ffe061acf1b66..131b249bc10b5 100644
--- a/llvm/lib/SYCLPostLink/CMakeLists.txt
+++ b/llvm/lib/SYCLPostLink/CMakeLists.txt
@@ -24,6 +24,7 @@ add_llvm_component_library(LLVMSYCLPostLink
Demangle
InstCombine
IRPrinter
+ Linker
Passes
ScalarOpts
Support
diff --git a/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp b/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp
index f2ae96e1c5e6b..4da338b244ad9 100644
--- a/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp
+++ b/llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp
@@ -11,9 +11,11 @@
#include "llvm/SYCLPostLink/ESIMDPostSplitProcessing.h"
#include "llvm/GenXIntrinsics/GenXSPIRVWriterAdaptor.h"
+#include "llvm/Linker/Linker.h"
#include "llvm/Passes/PassBuilder.h"
#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
#include "llvm/SYCLPostLink/ModuleSplitter.h"
+#include "llvm/Support/FormatVariadic.h"
#include "llvm/Transforms/IPO/AlwaysInliner.h"
#include "llvm/Transforms/IPO/StripDeadPrototypes.h"
#include "llvm/Transforms/InstCombine/InstCombine.h"
@@ -30,9 +32,10 @@ using namespace llvm::module_split;
namespace {
-ModulePassManager buildESIMDLoweringPipeline(bool OptLevelO0, bool SplitESIMD) {
+ModulePassManager
+buildESIMDLoweringPipeline(const sycl::ESIMDProcessingOptions &Options) {
ModulePassManager MPM;
- MPM.addPass(SYCLLowerESIMDPass(!SplitESIMD));
+ MPM.addPass(SYCLLowerESIMDPass(!Options.SplitESIMD));
FunctionPassManager FPM;
FPM.addPass(SROAPass(SROAOptions::ModifyCFG));
@@ -56,12 +59,29 @@ ModulePassManager buildESIMDLoweringPipeline(bool OptLevelO0, bool SplitESIMD) {
return MPM;
}
+Expected<ModuleDesc> linkModules(ModuleDesc MD1, ModuleDesc MD2) {
+ std::vector<std::string> Names;
+ MD1.saveEntryPointNames(Names);
+ MD2.saveEntryPointNames(Names);
+ bool LinkError =
+ llvm::Linker::linkModules(MD1.getModule(), MD2.releaseModulePtr());
+
+ if (LinkError)
+ return createStringError(
+ formatv("link failed. Module names: {0}, {1}", MD1.Name, MD2.Name));
+
+ ModuleDesc Res(MD1.releaseModulePtr(), std::move(Names));
+ Res.assignMergedProperties(MD1, MD2);
+ Res.Name = (Twine("linked[") + MD1.Name + "," + MD2.Name + "]").str();
+ return std::move(Res);
+}
+
} // anonymous namespace
// When ESIMD code was separated from the regular SYCL code,
// we can safely process ESIMD part.
-bool sycl::lowerESIMDConstructs(ModuleDesc &MD, bool OptLevelO0,
- bool SplitESIMD) {
+bool sycl::lowerESIMDConstructs(ModuleDesc &MD,
+ const sycl::ESIMDProcessingOptions &Options) {
// TODO: support options like -debug-pass, -print-[before|after], and others
LoopAnalysisManager LAM;
CGSCCAnalysisManager CGAM;
@@ -77,7 +97,7 @@ bool sycl::lowerESIMDConstructs(ModuleDesc &MD, bool OptLevelO0,
std::vector<std::string> Names;
MD.saveEntryPointNames(Names);
- ModulePassManager MPM = buildESIMDLoweringPipeline(OptLevelO0, SplitESIMD);
+ ModulePassManager MPM = buildESIMDLoweringPipeline(Options);
PreservedAnalyses Res = MPM.run(MD.getModule(), MAM);
// GenXSPIRVWriterAdaptor pass replaced some functions with "rewritten"
@@ -85,3 +105,56 @@ bool sycl::lowerESIMDConstructs(ModuleDesc &MD, bool OptLevelO0,
MD.rebuildEntryPoints(Names);
return !Res.areAllPreserved();
}
+
+Expected<SmallVector<ModuleDesc, 2>>
+llvm::sycl::handleESIMD(ModuleDesc MDesc,
+ const sycl::ESIMDProcessingOptions &Options,
+ bool &Modified, bool &SplitOccurred) {
+ SmallVector<ModuleDesc, 2> Result =
+ splitByESIMD(std::move(MDesc), Options.EmitOnlyKernelsAsEntryPoints,
+ Options.AllowDeviceImageDependencies);
+
+ assert(Result.size() <= 2 &&
+ "Split modules aren't expected to be more than 2.");
+
+ SplitOccurred |= Result.size() > 1;
+
+ for (ModuleDesc &MD : Result) {
+#ifdef LLVM_ENABLE_DUMP
+ dumpEntryPoints(MD.entries(), MD.Name.c_str(), 4);
+#endif // LLVM_ENABLE_DUMP
+ if (Options.LowerESIMD && MD.isESIMD())
+ Modified |= lowerESIMDConstructs(MD, Options);
+ }
+
+ if (Options.SplitESIMD || Result.size() == 1)
+ return std::move(Result);
+
+ // SYCL/ESIMD splitting is not requested, link back into single module.
+ int ESIMDInd = Result[0].isESIMD() ? 0 : 1;
+ int SYCLInd = 1 - ESIMDInd;
+ assert(Result[SYCLInd].isSYCL() &&
+ "Result[SYCLInd].isSYCL() expected to be true.");
+
+ // Make sure that no link conflicts occur.
+ Result[ESIMDInd].renameDuplicatesOf(Result[SYCLInd].getModule(), ".esimd");
+ auto LinkedOrErr = linkModules(std::move(Result[0]), std::move(Result[1]));
+ if (!LinkedOrErr)
+ return LinkedOrErr.takeError();
+
+ ModuleDesc &Linked = *LinkedOrErr;
+ Linked.restoreLinkageOfDirectInvokeSimdTargets();
+ std::vector<std::string> Names;
+ Linked.saveEntryPointNames(Names);
+ // Cleanup may remove some entry points, need to save/rebuild.
+ Linked.cleanup(Options.AllowDeviceImageDependencies);
+ Linked.rebuildEntryPoints(Names);
+ Result.clear();
+ Result.emplace_back(std::move(Linked));
+#ifdef LLVM_ENABLE_DUMP
+ dumpEntryPoints(Result.back().entries(), Result.back().Name.c_str(), 4);
+#endif // LLVM_ENABLE_DUMP
+ Modified = true;
+
+ return std::move(Result);
+}
diff --git a/llvm/tools/sycl-post-link/CMakeLists.txt b/llvm/tools/sycl-post-link/CMakeLists.txt
index ace9e66e44535..d229f80767fd3 100644
--- a/llvm/tools/sycl-post-link/CMakeLists.txt
+++ b/llvm/tools/sycl-post-link/CMakeLists.txt
@@ -9,7 +9,6 @@ set(LLVM_LINK_COMPONENTS
TransformUtils
SYCLLowerIR
SYCLPostLink
- Linker
Passes
Analysis
)
diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp
index 512f7a2614b31..0c28ba1f96447 100644
--- a/llvm/tools/sycl-post-link/sycl-post-link.cpp
+++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp
@@ -59,8 +59,6 @@
using namespace llvm;
using namespace llvm::sycl;
-using string_vector = std::vector<std::string>;
-
namespace {
#ifdef NDEBUG
@@ -259,6 +257,19 @@ struct IrPropSymFilenameTriple {
std::string Sym;
};
+unsigned getOptLevel() {
+ if (OptLevelO3)
+ return 3;
+ if (OptLevelO2 || OptLevelOs || OptLevelOz)
+ return 2;
+ if (OptLevelO1)
+ return 1;
+ if (OptLevelO0)
+ return 0;
+
+ return 2; // default value
+}
+
static void writeToFile(const StringRef Filename, const StringRef Content) {
std::error_code EC;
raw_fd_ostream OS{Filename, EC, sys::fs::OpenFlags::OF_None};
@@ -414,23 +425,6 @@ void saveDeviceLibModule(
saveModule(OutTables, DeviceLibMD, I, OutputPrefix, "");
}
-module_split::ModuleDesc link(module_split::ModuleDesc &&MD1,
- module_split::ModuleDesc &&MD2) {
- std::vector<std::string> Names;
- MD1.saveEntryPointNames(Names);
- MD2.saveEntryPointNames(Names);
- bool LinkError =
- llvm::Linker::linkModules(MD1.getModule(), MD2.releaseModulePtr());
-
- if (LinkError) {
- error(" error when linking SYCL and ESIMD modules");
- }
- module_split::ModuleDesc Res(MD1.releaseModulePtr(), std::move(Names));
- Res.assignMergedProperties(MD1, MD2);
- Res.Name = "linked[" + MD1.Name + "," + MD2.Name + "]";
- return Res;
-}
-
bool processSpecConstants(module_split::ModuleDesc &MD) {
MD.Props.SpecConstsMet = false;
@@ -496,64 +490,6 @@ void addTableRow(util::SimpleTable &Table,
Table.addRow(Row);
}
-SmallVector<module_split::ModuleDesc, 2>
-handleESIMD(module_split::ModuleDesc &&MDesc, bool &Modified,
- bool &SplitOccurred) {
- // Do SYCL/ESIMD splitting. It happens always, as ESIMD and SYCL must
- // undergo different set of LLVMIR passes. After this they are linked back
- // together to form single module with disjoint SYCL and ESIMD call graphs
- // unless -split-esimd option is specified. The graphs become disjoint
- // when linked back because functions shared between graphs are cloned and
- // renamed.
- SmallVector<module_split::ModuleDesc, 2> Result =
- module_split::splitByESIMD(std::move(MDesc), EmitOnlyKernelsAsEntryPoints,
- AllowDeviceImageDependencies);
-
- if (Result.size() > 1 && SplitOccurred &&
- (SplitMode == module_split::SPLIT_PER_KERNEL) && !SplitEsimd) {
- // Controversial state reached - SYCL and ESIMD entry points resulting
- // from SYCL/ESIMD split (which is done always) are linked back, since
- // -split-esimd is not specified, but per-kernel split is requested.
- warning("SYCL and ESIMD entry points detected and split mode is "
- "per-kernel, so " +
- SplitEsimd.ValueStr + " must also be specified");
- }
- SplitOccurred |= Result.size() > 1;
-
- for (auto &MD : Result) {
- DUMP_ENTRY_POINTS(MD.entries(), MD.Name.c_str(), 3);
- if (LowerEsimd && MD.isESIMD())
- Modified |= sycl::lowerESIMDConstructs(MD, OptLevelO0, SplitEsimd);
- }
-
- if (!SplitEsimd && Result.size() > 1) {
- // SYCL/ESIMD splitting is not requested, link back into single module.
- assert(Result.size() == 2 &&
- "Unexpected number of modules as results of ESIMD split");
- int ESIMDInd = Result[0].isESIMD() ? 0 : 1;
- int SYCLInd = 1 - ESIMDInd;
- assert(Result[SYCLInd].isSYCL() &&
- "no non-ESIMD module as a result ESIMD split?");
-
- // ... but before that, make sure no link conflicts will occur.
- Result[ESIMDInd].renameDuplicatesOf(Result[SYCLInd].getModule(), ".esimd");
- module_split::ModuleDesc Linked =
- link(std::move(Result[0]), std::move(Result[1]));
- Linked.restoreLinkageOfDirectInvokeSimdTargets();
- string_vector Names;
- Linked.saveEntryPointNames(Names);
- // cleanup may remove some entry points, need to save/rebuild
- Linked.cleanup(AllowDeviceImageDependencies);
- Linked.rebuildEntryPoints(Names);
- Result.clear();
- Result.emplace_back(std::move(Linked));
- DUMP_ENTRY_POINTS(Result.back().entries(), Result.back().Name.c_str(), 3);
- Modified = true;
- }
-
- return Result;
-}
-
// Checks if the given target and module are compatible.
// A target and module are compatible if all the optional kernel features
// the module uses are supported by that target (i.e. that module can be
@@ -671,10 +607,17 @@ processInputModule(std::unique_ptr<Module> M, const StringRef OutputPrefix) {
MDesc.fixupLinkageOfDirectInvokeSimdTargets();
- SmallVector<module_split::ModuleDesc, 2> MMs =
- handleESIMD(std::move(MDesc), Modified, SplitOccurred);
+ ESIMDProcessingOptions Options = {SplitMode,
+ EmitOnlyKernelsAsEntryPoints,
+ AllowDeviceImageDependencies,
+ LowerEsimd,
+ SplitEsimd,
+ getOptLevel()};
+ auto ModulesOrErr =
+ handleESIMD(std::move(MDesc), Options, Modified, SplitOccurred);
+ CHECK_AND_EXIT(ModulesOrErr.takeError());
+ SmallVector<module_split::ModuleDesc, 2> &MMs = *ModulesOrErr;
assert(MMs.size() && "at least one module is expected after ESIMD split");
-
SmallVector<module_split::ModuleDesc, 2> MMsWithDefaultSpecConsts;
for (size_t I = 0; I != MMs.size(); ++I) {
if (GenerateDeviceImageWithDefaultSpecConsts) {