Skip to content

[ESIMD][NFC] Extract ESIMD handling from sycl-post-link to library. #18684

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 3 commits into
base: sycl
Choose a base branch
from
Open

[ESIMD][NFC] Extract ESIMD handling from sycl-post-link to library. #18684

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
cb17d33
[ESIMD][NFC] Extract ESIMD handling from sycl-post-link to library.
maksimsab May 27, 2025
c490309
Merge branch 'sycl' into move_esimd_split
maksimsab Jun 13, 2025
8293ede
address CR feedback
maksimsab Jun 13, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
43 changes: 39 additions & 4 deletions llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,14 +11,49 @@

#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
1 change: 1 addition & 0 deletions llvm/lib/SYCLPostLink/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,7 @@ add_llvm_component_library(LLVMSYCLPostLink
Demangle
InstCombine
IRPrinter
Linker
Passes
ScalarOpts
Support
Expand Down
93 changes: 86 additions & 7 deletions llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -30,11 +32,12 @@ 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));

if (!OptLevelO0) {
if (Options.OptLevel != 0) {
FunctionPassManager FPM;
FPM.addPass(SROAPass(SROAOptions::ModifyCFG));
MPM.addPass(createModuleToFunctionPassAdaptor(std::move(FPM)));
Expand All @@ -43,7 +46,7 @@ ModulePassManager buildESIMDLoweringPipeline(bool OptLevelO0, bool SplitESIMD) {
FunctionPassManager MainFPM;
MainFPM.addPass(ESIMDLowerLoadStorePass{});

if (!OptLevelO0) {
if (Options.OptLevel != 0) {
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
MainFPM.addPass(EarlyCSEPass(true));
MainFPM.addPass(InstCombinePass{});
Expand All @@ -60,12 +63,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 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;
Expand All @@ -81,11 +101,70 @@ 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"
// versions so the entry point table must be rebuilt.
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.");
if (Result.size() == 2 && SplitOccurred &&
Options.SplitMode == module_split::SPLIT_PER_KERNEL &&
!Options.SplitESIMD)
return createStringError("SYCL and ESIMD entry points detected with "
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm not sure I understand this error message. Does it mean that if the modules contain both SYCL and ESIMD entry points, then split-mode=per-kernel and split-esimd=false are incompatible?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

size == 2 means that we have at least 1 ESIMD entry point and at least 1 non-ESIMD entry point.
-split-esimd=false would lead to linking 2 entry points into common module.
SPLIT_PER_KERNEL dictates that one module should contain only 1 entry point.
This is a controversial state that has been initially mentioned in the comment. However, there is a aspect that it has been initially a warning while my PR makes this an error which is not NFC. I could consider to leave it as a warning.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm not against the error, but the wording is quite confusing to me. I'd go for something like "--split-esimd=false is incompatible with split-mode=per-kernel".

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In fact, -split-esimd=false is compatible with split by kernel in terms of input arguments. The problem in that place is a combination of factors: Result.size() == 2, -split-esimd=false and split-mode=per-kernel.
The following combination is alright: Result.size() == 1, split-esimd=false, split-mode=per-kernel.

"-split-mode=per-kernel and -split-esimd=false. "
"So -split-esimd=true is mandatory.");

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 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 Result;
}
1 change: 0 additions & 1 deletion llvm/tools/sycl-post-link/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,6 @@ set(LLVM_LINK_COMPONENTS
TransformUtils
SYCLLowerIR
SYCLPostLink
Linker
Passes
Analysis
)
Expand Down
103 changes: 23 additions & 80 deletions llvm/tools/sycl-post-link/sycl-post-link.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,8 +59,6 @@
using namespace llvm;
using namespace llvm::sycl;

using string_vector = std::vector<std::string>;

namespace {

#ifdef NDEBUG
Expand Down Expand Up @@ -256,6 +254,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
}

void writeToFile(const std::string &Filename, const std::string &Content) {
std::error_code EC;
raw_fd_ostream OS{Filename, EC, sys::fs::OpenFlags::OF_None};
Expand Down Expand Up @@ -431,23 +442,6 @@ void saveDeviceLibModule(
saveModule(OutTables, DeviceLibMD, I, DeviceLibFileName);
}

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;

Expand Down Expand Up @@ -513,64 +507,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
Expand Down Expand Up @@ -688,10 +624,17 @@ processInputModule(std::unique_ptr<Module> M) {

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) {
Expand Down
Loading