Skip to content

[SYCL][NFC] Extract specialization constant's processing from sycl-post-link to SYCLPostLink library. #19022

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

Merged
merged 6 commits into from
Oct 9, 2025
Merged

[SYCL][NFC] Extract specialization constant's processing from sycl-post-link to SYCLPostLink library. #19022

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
d0de9a5
[SYCL][NFC] Extract specialization constant's processing from sycl-po…
maksimsab Jun 17, 2025
0325133
Merge branch 'sycl' into move_spec_const
maksimsab Jun 27, 2025
91fd691
address feedback from CR
maksimsab Jun 27, 2025
0a5a20e
fix comments
maksimsab Sep 29, 2025
861cefc
Merge branch 'sycl' into move_spec_const
maksimsab Sep 29, 2025
2ec7bc5
simplify function's interface
maksimsab Oct 6, 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
42 changes: 42 additions & 0 deletions llvm/include/llvm/SYCLPostLink/SpecializationConstants.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,42 @@
//= SpecializationConstants.h - Processing of SYCL Specialization Constants ==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// Specialization constants processing consists of lowering and generation
// of new module with spec consts replaced by default values.
//===----------------------------------------------------------------------===//

#ifndef LLVM_SYCL_POST_LINK_SPECIALIZATION_CONSTANTS_H
#define LLVM_SYCL_POST_LINK_SPECIALIZATION_CONSTANTS_H

#include "llvm/ADT/SmallVector.h"
#include "llvm/SYCLLowerIR/SpecConstants.h"
#include "llvm/SYCLPostLink/ModuleSplitter.h"

#include <optional>

namespace llvm {
namespace sycl {

/// Handling consists of SpecConsts's lowering depending on the given
/// \p Mode. If \p Mode is std::nullopt, then no lowering happens.
/// If \p GenerateModuleDescWithDefaultSpecConsts is true, then a generation
/// of new modules with specialization constants replaced by default values
/// happens and the result is written in \p NewModuleDescs.
/// Otherwise, \p NewModuleDescs is expected to be nullptr.
///
/// \returns Boolean value indicating whether the lowering has changed the input
/// modules.
bool handleSpecializationConstants(
llvm::SmallVectorImpl<module_split::ModuleDesc> &MDs,
std::optional<SpecConstantsPass::HandlingMode> Mode,
bool GenerateModuleDescWithDefaultSpecConsts = false,
llvm::SmallVectorImpl<module_split::ModuleDesc> *NewModuleDescs = nullptr);

} // namespace sycl
} // namespace llvm

#endif // LLVM_SYCL_POST_LINK_SPECIALIZATION_CONSTANTS_H
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 @@ -2,6 +2,7 @@ add_llvm_component_library(LLVMSYCLPostLink
ComputeModuleRuntimeInfo.cpp
ESIMDPostSplitProcessing.cpp
ModuleSplitter.cpp
SpecializationConstants.cpp

ADDITIONAL_HEADER_DIRS
${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLPostLink
Expand Down
100 changes: 100 additions & 0 deletions llvm/lib/SYCLPostLink/SpecializationConstants.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,100 @@
//= SpecializationConstants.cpp - Processing of SYCL Specialization Constants //
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// See comments in the header.
//===----------------------------------------------------------------------===//

#include "llvm/SYCLPostLink/SpecializationConstants.h"

#include "llvm/ADT/SmallVector.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/Module.h"
#include "llvm/IR/PassManager.h"
#include "llvm/Passes/PassBuilder.h"
#include "llvm/SYCLLowerIR/SpecConstants.h"
#include "llvm/SYCLPostLink/ModuleSplitter.h"
#include "llvm/Transforms/IPO/StripDeadPrototypes.h"

#include <optional>

using namespace llvm;

namespace {

bool lowerSpecConstants(module_split::ModuleDesc &MD,
SpecConstantsPass::HandlingMode Mode) {
ModulePassManager RunSpecConst;
ModuleAnalysisManager MAM;
SpecConstantsPass SCP(Mode);
// Register required analysis.
MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
RunSpecConst.addPass(std::move(SCP));

// Perform the spec constant intrinsics transformation on resulting module.
PreservedAnalyses Res = RunSpecConst.run(MD.getModule(), MAM);
MD.Props.SpecConstsMet = !Res.areAllPreserved();
return MD.Props.SpecConstsMet;
}

/// Function generates the copy of the given \p MD where all uses of
/// Specialization Constants are replaced by corresponding default values.
/// If the Module in \p MD doesn't contain specialization constants then
/// std::nullopt is returned.
std::optional<module_split::ModuleDesc>
cloneModuleWithSpecConstsReplacedByDefaultValues(
const module_split::ModuleDesc &MD) {
std::optional<module_split::ModuleDesc> NewMD;
if (!checkModuleContainsSpecConsts(MD.getModule()))
return NewMD;

NewMD = MD.clone();
NewMD->setSpecConstantDefault(true);

ModulePassManager MPM;
ModuleAnalysisManager MAM;
SpecConstantsPass SCP(SpecConstantsPass::HandlingMode::default_values);
MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
MPM.addPass(std::move(SCP));
MPM.addPass(StripDeadPrototypesPass());

PreservedAnalyses Res = MPM.run(NewMD->getModule(), MAM);
NewMD->Props.SpecConstsMet = !Res.areAllPreserved();
assert(NewMD->Props.SpecConstsMet &&
"This property should be true since the presence of SpecConsts "
"has been checked before the run of the pass");
NewMD->rebuildEntryPoints();
return NewMD;
}

} // namespace

bool llvm::sycl::handleSpecializationConstants(
SmallVectorImpl<module_split::ModuleDesc> &MDs,
std::optional<SpecConstantsPass::HandlingMode> Mode,
bool GenerateModuleDescWithDefaultSpecConsts,
SmallVectorImpl<module_split::ModuleDesc> *NewModuleDescs) {
[[maybe_unused]] bool AreArgumentsCompatible =
(GenerateModuleDescWithDefaultSpecConsts && NewModuleDescs) ||
(!GenerateModuleDescWithDefaultSpecConsts && !NewModuleDescs);
assert(AreArgumentsCompatible &&
"NewModuleDescs pointer is nullptr iff "
"GenerateModuleDescWithDefaultSpecConsts is false.");

bool Modified = false;
for (module_split::ModuleDesc &MD : MDs) {
if (GenerateModuleDescWithDefaultSpecConsts)
if (std::optional<module_split::ModuleDesc> NewMD =
cloneModuleWithSpecConstsReplacedByDefaultValues(MD);
NewMD)
NewModuleDescs->push_back(std::move(*NewMD));

if (Mode)
Modified |= lowerSpecConstants(MD, *Mode);
}

return Modified;
}
72 changes: 11 additions & 61 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 @@ -40,6 +40,7 @@
#include "llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h"
#include "llvm/SYCLPostLink/ESIMDPostSplitProcessing.h"
#include "llvm/SYCLPostLink/ModuleSplitter.h"
#include "llvm/SYCLPostLink/SpecializationConstants.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/FileSystem.h"
#include "llvm/Support/InitLLVM.h"
Expand All @@ -48,7 +49,6 @@
#include "llvm/Support/SourceMgr.h"
#include "llvm/Support/SystemUtils.h"
#include "llvm/Support/WithColor.h"
#include "llvm/Transforms/IPO/StripDeadPrototypes.h"

#include <algorithm>
#include <memory>
Expand Down Expand Up @@ -451,56 +451,6 @@ module_split::ModuleDesc link(module_split::ModuleDesc &&MD1,
return Res;
}

bool processSpecConstants(module_split::ModuleDesc &MD) {
MD.Props.SpecConstsMet = false;

if (SpecConstLower.getNumOccurrences() == 0)
return false;

ModulePassManager RunSpecConst;
ModuleAnalysisManager MAM;
SpecConstantsPass SCP(SpecConstLower == SC_NATIVE_MODE
? SpecConstantsPass::HandlingMode::native
: SpecConstantsPass::HandlingMode::emulation);
// Register required analysis
MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
RunSpecConst.addPass(std::move(SCP));

// Perform the spec constant intrinsics transformation on resulting module
PreservedAnalyses Res = RunSpecConst.run(MD.getModule(), MAM);
MD.Props.SpecConstsMet = !Res.areAllPreserved();
return MD.Props.SpecConstsMet;
}

/// Function generates the copy of the given ModuleDesc where all uses of
/// Specialization Constants are replaced by corresponding default values.
/// If the Module in MD doesn't contain specialization constants then
/// std::nullopt is returned.
std::optional<module_split::ModuleDesc>
processSpecConstantsWithDefaultValues(const module_split::ModuleDesc &MD) {
std::optional<module_split::ModuleDesc> NewModuleDesc;
if (!checkModuleContainsSpecConsts(MD.getModule()))
return NewModuleDesc;

NewModuleDesc = MD.clone();
NewModuleDesc->setSpecConstantDefault(true);

ModulePassManager MPM;
ModuleAnalysisManager MAM;
SpecConstantsPass SCP(SpecConstantsPass::HandlingMode::default_values);
MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
MPM.addPass(std::move(SCP));
MPM.addPass(StripDeadPrototypesPass());

PreservedAnalyses Res = MPM.run(NewModuleDesc->getModule(), MAM);
NewModuleDesc->Props.SpecConstsMet = !Res.areAllPreserved();
assert(NewModuleDesc->Props.SpecConstsMet &&
"This property should be true since the presence of SpecConsts "
"has been checked before the run of the pass");
NewModuleDesc->rebuildEntryPoints();
return NewModuleDesc;
}

constexpr int MAX_COLUMNS_IN_FILE_TABLE = 3;

void addTableRow(util::SimpleTable &Table,
Expand Down Expand Up @@ -682,6 +632,12 @@ processInputModule(std::unique_ptr<Module> M) {
error(toString(std::move(E)));
}

std::optional<SpecConstantsPass::HandlingMode> SCMode;
if (SpecConstLower.getNumOccurrences() > 0)
SCMode = SpecConstLower == SC_NATIVE_MODE
Copy link
Contributor

Choose a reason for hiding this comment

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

is there a default value we can use instead of having it being optional? maybe like not_specified or something?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Would it be better to have not_specified value? We will have to check that value and return errors deeper in lowering.

Copy link
Contributor

Choose a reason for hiding this comment

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

personally i would prefer that

Copy link
Contributor Author

@maksimsab maksimsab Oct 6, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

What is the expected behaviour in the following snippet then would be?

  ModulePassManager MPM;
  ModuleAnalysisManager MAM;
  SpecConstantsPass SCP(SpecConstantsPass::HandlingMode::not_specified);
  MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
  MPM.addPass(std::move(SCP));

  PreservedAnalyses Res = MPM.run(NewMD->getModule(), MAM);

Copy link
Contributor

Choose a reason for hiding this comment

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

@sarnex , @maksimsab , I guess the answer to Nick's question in this thread is: there is no default value currently that can be used instead of having SCMode to be optional.

This PR is about extracting existing spec const processing to a library. Updating SpecConstantsPass and introducing new HandlingMode is out of scope of this PR. We can consider it as a future refactoring, but it should not hold this PR from merging.

Copy link
Contributor

Choose a reason for hiding this comment

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

it's not worth blocking the PR over, but i'm not sure i understand, the old code does not seem to use optional for the handling mode, so im wondering why we need it in this NFC change

Copy link
Contributor

Choose a reason for hiding this comment

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

That is a good question! The reason is how logic was split between sycl-post-link and SYCLPostLink library.

The info about spec const mode comes from cl::opt option in sycl-post-link tool, which is:
cl::opt<SpecConstLowerMode> SpecConstLower and is cl::Optional.

Since in the old code SpecConstLower is global for that file and the logic which depends on its optionality is in the same file, we did not need additional optional variable. So there is code like:

if (SpecConstLower.getNumOccurrences() == 0)
    return false;

So, semantically, old code used optional, just not std::optional.

In the new code, we need to keep optional behavior and pass this information somehow to the library function. It seems using std::optional is the closest way to keep the original semantics.

The alternative (to keep the code as close to the original as possible) that I see is to pass cl::opt<SpecConstLowerMode> SpecConstLower as parameter, instead of std::optional<SpecConstantsPass::HandlingMode> Mode.
That doesn't look better to me...

What do you think?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Updating SpecConstantsPass and introducing new HandlingMode is out of scope of this PR.

That is correct. Introducing default value would lead to many changes in SpecConstsPass which worth a separate patch.

Copy link
Contributor

@sarnex sarnex Oct 7, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

i agree that calling getNumOccurences() == 0 is similar to checking if the optional is empty. so yeah feel free to merge the pr and address this in a follow up

? SpecConstantsPass::HandlingMode::native
: SpecConstantsPass::HandlingMode::emulation;

// It is important that we *DO NOT* preserve all the splits in memory at the
// same time, because it leads to a huge RAM consumption by the tool on bigger
// inputs.
Expand All @@ -696,16 +652,10 @@ processInputModule(std::unique_ptr<Module> M) {
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) {
std::optional<module_split::ModuleDesc> NewMD =
processSpecConstantsWithDefaultValues(MMs[I]);
if (NewMD)
MMsWithDefaultSpecConsts.push_back(std::move(*NewMD));
}

Modified |= processSpecConstants(MMs[I]);
}
Modified |= handleSpecializationConstants(
MMs, SCMode, GenerateDeviceImageWithDefaultSpecConsts,
GenerateDeviceImageWithDefaultSpecConsts ? &MMsWithDefaultSpecConsts
: nullptr);

if (IROutputOnly) {
if (SplitOccurred) {
Expand Down
Loading