Skip to content

Commit 6127715

Browse files
[SYCL] Implement device image properties for virtual functions (#14875)
Implementation design explaining those changes in a bigger picture can be found in #10540 Key things implemented here: - device code split to outline virtual functions into separate device images - emission of new properties for virtual functions - generation of `calls-indirectly` LLVM IR attribute for kernels that construct objects with virtual functions, but don't do calls - device image manipulations to cleanup or preserve virtual functions depending on a device image Even though those pieces are technically independent from each other, it is hard to split them apart into separate PRs, because they all have to be either present or absent for existing E2E tests for virtual functions to work.
1 parent 429b01d commit 6127715

15 files changed

+676
-19
lines changed

llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -82,6 +82,11 @@ struct EntryPointGroup {
8282
// Scope remains global
8383
return Res;
8484
}
85+
86+
// Indicates that this group holds definitions of virtual functions - they
87+
// are outlined into separate device images and should be removed from all
88+
// other modules. The flag is used in ModuleDesc::cleanup
89+
bool HasVirtualFunctionDefinitions = false;
8590
};
8691

8792
std::string GroupId;

llvm/include/llvm/Support/PropertySetIO.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -209,6 +209,7 @@ class PropertySetRegistry {
209209
static constexpr char SYCL_DEVICE_GLOBALS[] = "SYCL/device globals";
210210
static constexpr char SYCL_DEVICE_REQUIREMENTS[] = "SYCL/device requirements";
211211
static constexpr char SYCL_HOST_PIPES[] = "SYCL/host pipes";
212+
static constexpr char SYCL_VIRTUAL_FUNCTIONS[] = "SYCL/virtual functions";
212213

213214
/// Function for bulk addition of an entire property set in the given
214215
/// \p Category .

llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp

Lines changed: 63 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,9 @@
88
// See comments in the header.
99
//===----------------------------------------------------------------------===//
1010
#include "llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h"
11+
#include "llvm/ADT/SmallString.h"
12+
#include "llvm/ADT/SmallVector.h"
13+
#include "llvm/ADT/StringSet.h"
1114
#include "llvm/Demangle/Demangle.h"
1215
#include "llvm/IR/PassInstrumentation.h"
1316
#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
@@ -188,6 +191,10 @@ PropSetRegTy computeModuleProperties(const Module &M,
188191
if (GlobProps.EmitExportedSymbols) {
189192
// extract exported functions if any and save them into property set
190193
for (const auto *F : EntryPoints) {
194+
// Virtual functions use a different mechanism of dynamic linking, they
195+
// should not be registered here.
196+
if (F->hasFnAttribute("indirectly-callable"))
197+
continue;
191198
// TODO FIXME some of SYCL/ESIMD functions maybe marked with __regcall CC,
192199
// so they won't make it into the export list. Should the check be
193200
// F->getCallingConv() != CallingConv::SPIR_KERNEL?
@@ -201,11 +208,19 @@ PropSetRegTy computeModuleProperties(const Module &M,
201208
if (GlobProps.EmitImportedSymbols) {
202209
// record imported functions in the property set
203210
for (const auto &F : M) {
204-
if ( // A function that can be imported may still be defined in one split
205-
// image. Only add import property if this is not the image where the
206-
// function is defined.
207-
F.isDeclaration() && module_split::canBeImportedFunction(F)) {
211+
// A function that can be imported may still be defined in one split
212+
// image. Only add import property if this is not the image where the
213+
// function is defined.
214+
if (!F.isDeclaration())
215+
continue;
208216

217+
// Even though virtual functions are considered to be imported by the
218+
// function below, we shouldn't list them in the property because they
219+
// use different mechanism for dynamic linking.
220+
if (F.hasFnAttribute("indirectly-callable"))
221+
continue;
222+
223+
if (module_split::canBeImportedFunction(F)) {
209224
// StripDeadPrototypes is called during module splitting
210225
// cleanup. At this point all function decls should have uses.
211226
assert(!F.use_empty() && "Function F has no uses");
@@ -354,6 +369,50 @@ PropSetRegTy computeModuleProperties(const Module &M,
354369
PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "specConstsReplacedWithDefault",
355370
1);
356371

372+
{ // Properties related to virtual functions
373+
StringSet<> UsedVFSets;
374+
bool AddedVFSetProperty = false;
375+
for (const Function &F : M) {
376+
if (F.isDeclaration())
377+
continue;
378+
379+
if (F.hasFnAttribute("indirectly-callable")) {
380+
PropSet.add(PropSetRegTy::SYCL_VIRTUAL_FUNCTIONS,
381+
"virtual-functions-set",
382+
F.getFnAttribute("indirectly-callable").getValueAsString());
383+
AddedVFSetProperty = true;
384+
// Device code split should ensure that virtual functions that belong
385+
// to different sets are split into separate device images and hence
386+
// there is no need to scan other functions.
387+
break;
388+
}
389+
390+
if (F.hasFnAttribute("calls-indirectly")) {
391+
SmallVector<StringRef, 4> Sets;
392+
F.getFnAttribute("calls-indirectly")
393+
.getValueAsString()
394+
.split(Sets, ',', /* MaxSplits */ -1, /* KeepEmpty */ false);
395+
for (auto Set : Sets)
396+
UsedVFSets.insert(Set);
397+
}
398+
}
399+
400+
if (!UsedVFSets.empty()) {
401+
assert(!AddedVFSetProperty &&
402+
"device image cannot have both virtual-functions-set and "
403+
"uses-virtual-functions-set property");
404+
SmallString<128> AllSets;
405+
for (auto &It : UsedVFSets) {
406+
if (!AllSets.empty())
407+
AllSets += ',';
408+
AllSets += It.getKey();
409+
}
410+
411+
PropSet.add(PropSetRegTy::SYCL_VIRTUAL_FUNCTIONS,
412+
"uses-virtual-functions-set", AllSets);
413+
}
414+
}
415+
357416
return PropSet;
358417
}
359418
std::string computeModuleSymbolTable(const Module &M,

llvm/lib/SYCLLowerIR/ModuleSplitter.cpp

Lines changed: 44 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -136,7 +136,10 @@ bool isEntryPoint(const Function &F, bool EmitOnlyKernelsAsEntryPoints) {
136136
!isGenericBuiltin(F.getName());
137137
}
138138

139-
return false;
139+
// Even if we are emitting only kernels as entry points, virtual functions
140+
// should still be treated as entry points, because they are going to be
141+
// outlined into separate device images and linked in later.
142+
return F.hasFnAttribute("indirectly-callable");
140143
}
141144

142145
// Represents "dependency" or "use" graph of global objects (functions and
@@ -668,6 +671,22 @@ bool mustPreserveGV(const GlobalValue &GV) {
668671
// TODO: try to move all passes (cleanup, spec consts, compile time properties)
669672
// in one place and execute MPM.run() only once.
670673
void ModuleDesc::cleanup() {
674+
// Any definitions of virtual functions should be removed and turned into
675+
// declarations, they are supposed to be provided by a different module.
676+
if (!EntryPoints.Props.HasVirtualFunctionDefinitions) {
677+
for (Function &F : *M)
678+
if (F.hasFnAttribute("indirectly-callable")) {
679+
F.deleteBody();
680+
if (F.hasComdat())
681+
F.setComdat(nullptr);
682+
}
683+
} else {
684+
// Otherwise externalize them so they are not dropped by GlobalDCE
685+
for (Function &F : *M)
686+
if (F.hasFnAttribute("indirectly-callable"))
687+
F.setLinkage(GlobalValue::LinkageTypes::ExternalLinkage);
688+
}
689+
671690
ModuleAnalysisManager MAM;
672691
MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
673692
ModulePassManager MPM;
@@ -1057,6 +1076,17 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
10571076
Categorizer.registerSimpleStringAttributeRule(
10581077
sycl::utils::ATTR_SYCL_MODULE_ID);
10591078

1079+
// This attribute marks virtual functions and effectively dictates how they
1080+
// should be groupped together. By design we won't split those groups of
1081+
// virtual functions further even if functions from the same group use
1082+
// different optional features and therefore this rule is put here.
1083+
// Strictly speaking, we don't even care about module-id splitting for
1084+
// those, but to avoid that we need to refactor the whole categorizer.
1085+
// However, this is good enough as it is for an initial version.
1086+
// TODO: for AOT use case we shouldn't be outlining those and instead should
1087+
// only select those functions which are compatible with the target device
1088+
Categorizer.registerSimpleStringAttributeRule("indirectly-callable");
1089+
10601090
// Optional features
10611091
// Note: Add more rules at the end of the list to avoid chaning orders of
10621092
// output files in existing tests.
@@ -1096,8 +1126,19 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
10961126
Groups.reserve(EntryPointsMap.size());
10971127
// Start with properties of a source module
10981128
EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props;
1099-
for (auto &[Key, EntryPoints] : EntryPointsMap)
1100-
Groups.emplace_back(Key, std::move(EntryPoints), MDProps);
1129+
for (auto &[Key, EntryPoints] : EntryPointsMap) {
1130+
bool HasVirtualFunctions = false;
1131+
for (auto *F : EntryPoints) {
1132+
if (F->hasFnAttribute("indirectly-callable")) {
1133+
HasVirtualFunctions = true;
1134+
break;
1135+
}
1136+
}
1137+
1138+
auto PropsCopy = MDProps;
1139+
PropsCopy.HasVirtualFunctionDefinitions = HasVirtualFunctions;
1140+
Groups.emplace_back(Key, std::move(EntryPoints), PropsCopy);
1141+
}
11011142
}
11021143

11031144
bool DoSplit = (Mode != SPLIT_NONE &&

0 commit comments

Comments
 (0)