From 3bb0d740ac92610ab45d3d48fb63231265ae0e90 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Feb 2023 07:56:19 -0500 Subject: [PATCH 01/33] WIP --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 102 +++++++++++++++++++ llvm/tools/sycl-post-link/ModuleSplitter.h | 97 ++++++++++++++++++ 2 files changed, 199 insertions(+) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index a517d7639fffc..256d8d8f8b8ee 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -881,5 +881,107 @@ getSplitterByOptionalFeatures(ModuleDesc &&MD, return std::make_unique(std::move(MD), std::move(Groups)); } +void DeviceCodeSplitRulesBuilder::registerRule( + const std::function &Callback) { + Rules.push_back(Rule::get(Callback)); +} + +void DeviceCodeSplitRulesBuilder::registerSimpleStringAttributeRule( + StringRef Attr) { + Rules.push_back(Rule::get(Attr)); +} + +void DeviceCodeSplitRulesBuilder::registerSimpleFlagMetadataRule( + StringRef TrueStr, StringRef FalseStr, StringRef MetadataName) { + Rules.push_back( + Rule::get(TrueStr, FalseStr, MetadataName)); +} + +void DeviceCodeSplitRulesBuilder::registerListOfIntegersInMetadataRule( + StringRef MetadataName) { + Rules.push_back(Rule::get(MetadataName)); +} + +std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { + std::string Result; + for (const auto &R : Rules) { + switch (R.Kind) { + case RuleKind::CALLBACK: + Result += R.getCallbackRuleData().Callback(F); + break; + case RuleKind::SIMPLE_STRING_ATTR: { + auto AttrName = R.getSimpleStringAttrRuleData().Attr; + if (F->hasFnAttribute(AttrName)) { + auto Attr = F->getFnAttribute(AttrName); + assert(Attr.isStringAttribute()); + Result += Attr.getValueAsString(); + } + } break; + case RuleKind::FLAG_METADATA: { + auto Data = R.getFlagMetadataRuleData(); + if (F->hasMetadata(Data.MetadataName)) + Result += Data.TrueStr; + else + Result += Data.FalseStr; + } break; + case RuleKind::INTEGERS_LIST_METADATA: { + auto MetadataName = R.getIntegersListMetadataRuleData().MetadataName; + if (F->hasMetadata(MetadataName)) { + auto *MDN = F->getMetadata(MetadataName); + for (const MDOperand &MDOp : MDN->operands()) + Result += std::to_string( + mdconst::extract(MDOp)->getZExtValue()); + } + } break; + } + } + + return Result; +} + +std::unique_ptr +getSplitterByRules(ModuleDesc &&MD, const DeviceCodeSplitRulesBuilder &Rules, + bool EmitOnlyKernelsAsEntryPoints) { + EntryPointGroupVec Groups; + + StringMap GroupNameToFunctionsMap; + + Module &M = MD.getModule(); + + // Only process module entry points: + for (auto &F : M.functions()) { + if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints) || + !MD.isEntryPointCandidate(F)) { + continue; + } + + auto Key = Rules.executeRules(&F); + GroupNameToFunctionsMap[Key].insert(&F); + } + + if (GroupNameToFunctionsMap.empty()) { + // No entry points met, record this. + Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); + } else { + Groups.reserve(GroupNameToFunctionsMap.size()); + for (auto &It : GroupNameToFunctionsMap) { + auto Name = It.getKey(); + EntryPointSet &EntryPoints = It.getValue(); + + // Start with properties of a source module + EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; + // FIXME: properly set top-level properties + // if (Features.UsesLargeGRF) + // MDProps.UsesLargeGRF = true; + Groups.emplace_back(Name, std::move(EntryPoints), MDProps); + } + } + + if (Groups.size() > 1) + return std::make_unique(std::move(MD), std::move(Groups)); + else + return std::make_unique(std::move(MD), std::move(Groups)); +} + } // namespace module_split } // namespace llvm diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index 037be3f65a891..4c6f1bd7a7169 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -244,6 +244,103 @@ class ModuleSplitterBase { bool hasMoreSplits() const { return remainingSplits() > 0; } }; +class DeviceCodeSplitRulesBuilder { +public: + DeviceCodeSplitRulesBuilder() = default; + + std::string executeRules(Function *) const; + + // Accepts a callback, which should return a string based on provided + // function, which will be used as an entry points group identifier. + void registerRule(const std::function &); + + // Creates a simple rule, which adds a value of a string attribute into a + // resulting identifier. + void registerSimpleStringAttributeRule(StringRef); + + // Creates a simple rule, which adds one or another value to a resulting + // identifier based on a presence of a metadata on a function. + void registerSimpleFlagMetadataRule(StringRef, StringRef, StringRef); + + // Creates a rule, which adds a list of dash-separated integers converted + // into strings listed in a metadata to a resulting identifier. + void registerListOfIntegersInMetadataRule(StringRef); + +private: + enum class RuleKind { + CALLBACK, SIMPLE_STRING_ATTR, FLAG_METADATA, INTEGERS_LIST_METADATA + }; + + struct CallbackRuleData { + constexpr static auto Kind = RuleKind::CALLBACK; + CallbackRuleData() = default; + std::function Callback = nullptr; + }; + + struct SimpleStringAttrRuleData { + constexpr static auto Kind = RuleKind::SIMPLE_STRING_ATTR; + SimpleStringAttrRuleData() = default; + StringRef Attr; + }; + + struct FlagMetadataRuleData { + constexpr static auto Kind = RuleKind::FLAG_METADATA; + FlagMetadataRuleData() = default; + StringRef TrueStr, FalseStr, MetadataName; + }; + + struct IntegersListMetadataRuleData { + constexpr static auto Kind = RuleKind::INTEGERS_LIST_METADATA; + IntegersListMetadataRuleData() = default; + StringRef MetadataName = ""; + }; + + struct Rule { + private: + std::array + Storage; + public: + RuleKind Kind; + + template + static Rule get(Args... args) { + Rule R; + new(R.Storage.data()) T {args...}; + R.Kind = T::Kind; + return R; + } + + CallbackRuleData getCallbackRuleData() const { + assert(Kind == RuleKind::CALLBACK); + return *reinterpret_cast(Storage.data()); + } + + SimpleStringAttrRuleData getSimpleStringAttrRuleData() const { + assert(Kind == RuleKind::SIMPLE_STRING_ATTR); + return *reinterpret_cast(Storage.data()); + } + + FlagMetadataRuleData getFlagMetadataRuleData() const { + assert(Kind == RuleKind::FLAG_METADATA); + return *reinterpret_cast(Storage.data()); + } + + IntegersListMetadataRuleData getIntegersListMetadataRuleData() const { + assert(Kind == RuleKind::INTEGERS_LIST_METADATA); + return *reinterpret_cast(Storage.data()); + } + }; + + std::vector Rules; +}; + +std::unique_ptr +getSplitterByRules(ModuleDesc &&MD, const DeviceCodeSplitRulesBuilder &Rules, + bool EmitOnlyKernelsAsEntryPoints); + std::unique_ptr getSplitterByKernelType(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); From 4ca6479efe9741687cf54ea987a1449e9937d66a Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 27 Mar 2023 10:00:58 -0400 Subject: [PATCH 02/33] start using new splitter --- .../tools/sycl-post-link/assert/property-1.ll | 6 +- .../sycl-post-link/emit_exported_symbols.ll | 10 +-- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 87 +++++++++++++++++++ llvm/tools/sycl-post-link/ModuleSplitter.h | 22 ++++- llvm/tools/sycl-post-link/sycl-post-link.cpp | 64 +++----------- 5 files changed, 127 insertions(+), 62 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/assert/property-1.ll b/llvm/test/tools/sycl-post-link/assert/property-1.ll index 57f76e414b3a8..81e7c674187be 100644 --- a/llvm/test/tools/sycl-post-link/assert/property-1.ll +++ b/llvm/test/tools/sycl-post-link/assert/property-1.ll @@ -12,9 +12,9 @@ ; RUN: FileCheck %s -input-file=%t_0.prop --implicit-check-not TheKernel2 ; ; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table -; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes=CHECK-K1 -; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes=CHECK-K2 -; RUN: FileCheck %s -input-file=%t_2.prop --check-prefixes=CHECK-K3 +; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes=CHECK-K3 +; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes=CHECK-K1 +; RUN: FileCheck %s -input-file=%t_2.prop --check-prefixes=CHECK-K2 ; SYCL source: ; void foo() { diff --git a/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll b/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll index 167b3a8dc6558..7d660fd12bbb6 100644 --- a/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll +++ b/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll @@ -12,11 +12,11 @@ ; ; Per-kernel split ; RUN: sycl-post-link -symbols -split=kernel -emit-exported-symbols -S < %s -o %t.per_kernel.files.table -; RUN: FileCheck %s -input-file=%t.per_kernel.files_0.prop --implicit-check-not="NotExported" --check-prefix=CHECK-PERKERNEL-0-PROP -; RUN: FileCheck %s -input-file=%t.per_kernel.files_1.prop --implicit-check-not="NotExported" --check-prefix=CHECK-PERKERNEL-1-PROP -; RUN: FileCheck %s -input-file=%t.per_kernel.files_2.prop --implicit-check-not="NotExported" --check-prefix=CHECK-PERKERNEL-2-PROP -; RUN: FileCheck %s -input-file=%t.per_kernel.files_3.prop --implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP -; RUN: FileCheck %s -input-file=%t.per_kernel.files_4.prop --implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP +; RUN: FileCheck %s -input-file=%t.per_kernel.files_0.prop --implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP +; RUN: FileCheck %s -input-file=%t.per_kernel.files_1.prop --implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP +; RUN: FileCheck %s -input-file=%t.per_kernel.files_2.prop --implicit-check-not="NotExported" --check-prefix=CHECK-PERKERNEL-0-PROP +; RUN: FileCheck %s -input-file=%t.per_kernel.files_3.prop --implicit-check-not="NotExported" --check-prefix=CHECK-PERKERNEL-1-PROP +; RUN: FileCheck %s -input-file=%t.per_kernel.files_4.prop --implicit-check-not="NotExported" --check-prefix=CHECK-PERKERNEL-2-PROP target triple = "spir64-unknown-unknown" diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 640adc7f671d3..5b4ea3dad666e 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -891,6 +891,12 @@ void DeviceCodeSplitRulesBuilder::registerSimpleStringAttributeRule( Rules.push_back(Rule::get(Attr)); } +void DeviceCodeSplitRulesBuilder::registerSimpleFlagAttributeRule( + StringRef Attr, StringRef TrueStr, StringRef FalseStr) { + Rules.push_back( + Rule::get(TrueStr, FalseStr, Attr)); +} + void DeviceCodeSplitRulesBuilder::registerSimpleFlagMetadataRule( StringRef TrueStr, StringRef FalseStr, StringRef MetadataName) { Rules.push_back( @@ -933,6 +939,14 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { mdconst::extract(MDOp)->getZExtValue()); } } break; + case RuleKind::FLAG_ATTR: { + auto AttrName = R.getFlagAttributeRuleData().AttrName; + if (F->hasFnAttribute(AttrName)) { + Result += R.getFlagAttributeRuleData().TrueStr; + } else { + Result += R.getFlagAttributeRuleData().FalseStr; + } + } break; } } @@ -983,5 +997,78 @@ getSplitterByRules(ModuleDesc &&MD, const DeviceCodeSplitRulesBuilder &Rules, return std::make_unique(std::move(MD), std::move(Groups)); } +std::unique_ptr +getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, + bool EmitOnlyKernelsAsEntryPoints) { + DeviceCodeSplitRulesBuilder RulesBuilder; + + EntryPointsGroupScope Scope = selectDeviceCodeGroupScope( + MD.getModule(), Mode, /* AutoSplitIsGlobalScope */ false); + + if (Scope == Scope_Global) { + // We simply perform entry points filtering, but group all of them together. + RulesBuilder.registerRule([](Function *) -> std::string { + return GLOBAL_SCOPE_NAME; + }); + } else if (Scope == Scope_PerKernel) { + // Per-kernel split is quite simple: every kernel goes into a separate + // module and that's it, no other rules required. + RulesBuilder.registerRule([](Function *F) -> std::string { + return F->getName().str(); + }); + } else if (Scope == Scope_PerModule) { + // The most complex case, because we should account for many other features + // like aspects used in a kernel, large-grf mode, reqd-work-group-size, etc. + + // This is core of per-source device code split + RulesBuilder.registerSimpleStringAttributeRule("sycl-module-id"); + + // Optional features + RulesBuilder.registerSimpleFlagAttributeRule(::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf"); + RulesBuilder.registerListOfIntegersInMetadataRule("sycl_used_aspects"); + RulesBuilder.registerListOfIntegersInMetadataRule("reqd_work_group_size"); + + } else { + llvm_unreachable("Unexpected split scope"); + } + + // std::map is used here to ensure stable ordering of entry point groups, + // which is based on their contents, this greatly helps LIT tests + std::map EntryPointsMap; + + // Only process module entry points: + for (auto &F : MD.getModule().functions()) { + if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints)) + continue; + + std::string Key = RulesBuilder.executeRules(&F); + EntryPointsMap[std::move(Key)].insert(&F); + } + + EntryPointGroupVec Groups; + + if (EntryPointsMap.empty()) { + // No entry points met, record this. + Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); + } else { + Groups.reserve(EntryPointsMap.size()); + for (auto &It : EntryPointsMap) { + EntryPointSet &EntryPoints = It.second; + + // Start with properties of a source module + EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; + // FIXME: Propagate LargeGRF flag to entry points group + // if (Features.UsesLargeGRF) + // MDProps.UsesLargeGRF = true; + Groups.emplace_back(It.first, std::move(EntryPoints), MDProps); + } + } + + if (Groups.size() > 1) + return std::make_unique(std::move(MD), std::move(Groups)); + else + return std::make_unique(std::move(MD), std::move(Groups)); +} + } // namespace module_split } // namespace llvm diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index 27092deb126be..1d0333560e77c 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -258,6 +258,10 @@ class DeviceCodeSplitRulesBuilder { // resulting identifier. void registerSimpleStringAttributeRule(StringRef); + // Creates a simple rule, which adds one or another value to a resulting + // identifier based on a presence of a metadata on a function. + void registerSimpleFlagAttributeRule(StringRef, StringRef, StringRef = ""); + // Creates a simple rule, which adds one or another value to a resulting // identifier based on a presence of a metadata on a function. void registerSimpleFlagMetadataRule(StringRef, StringRef, StringRef); @@ -268,7 +272,7 @@ class DeviceCodeSplitRulesBuilder { private: enum class RuleKind { - CALLBACK, SIMPLE_STRING_ATTR, FLAG_METADATA, INTEGERS_LIST_METADATA + CALLBACK, SIMPLE_STRING_ATTR, FLAG_METADATA, INTEGERS_LIST_METADATA, FLAG_ATTR }; struct CallbackRuleData { @@ -289,6 +293,12 @@ class DeviceCodeSplitRulesBuilder { StringRef TrueStr, FalseStr, MetadataName; }; + struct FlagAttributeRuleData { + constexpr static auto Kind = RuleKind::FLAG_ATTR; + FlagAttributeRuleData() = default; + StringRef TrueStr, FalseStr, AttrName; + }; + struct IntegersListMetadataRuleData { constexpr static auto Kind = RuleKind::INTEGERS_LIST_METADATA; IntegersListMetadataRuleData() = default; @@ -300,6 +310,7 @@ class DeviceCodeSplitRulesBuilder { std::array Storage; public: @@ -332,6 +343,11 @@ class DeviceCodeSplitRulesBuilder { assert(Kind == RuleKind::INTEGERS_LIST_METADATA); return *reinterpret_cast(Storage.data()); } + + FlagAttributeRuleData getFlagAttributeRuleData() const { + assert(Kind == RuleKind::FLAG_ATTR); + return *reinterpret_cast(Storage.data()); + } }; std::vector Rules; @@ -341,6 +357,10 @@ std::unique_ptr getSplitterByRules(ModuleDesc &&MD, const DeviceCodeSplitRulesBuilder &Rules, bool EmitOnlyKernelsAsEntryPoints); +std::unique_ptr +getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, + bool EmitOnlyKernelsAsEntryPoints); + std::unique_ptr getSplitterByKernelType(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index f9110752bc331..108cd95525b4c 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -759,60 +759,19 @@ processInputModule(std::unique_ptr M) { (SplitMode == module_split::SPLIT_AUTO)) && "invalid split mode for IR-only output"); - // Top-level per-kernel/per-source splitter. SYCL/ESIMD splitting is applied - // to modules resulting from all other kinds of splitting. - std::unique_ptr ScopedSplitter = - module_split::getSplitterByMode(module_split::ModuleDesc{std::move(M)}, - SplitMode, IROutputOnly, - EmitOnlyKernelsAsEntryPoints); - - SmallVector TopLevelModules; + // FIXME: handle IROutputOnly + std::unique_ptr Splitter = + module_split::getDeviceCodeSplitter(module_split::ModuleDesc{std::move(M)}, + SplitMode, EmitOnlyKernelsAsEntryPoints); + const bool Split = Splitter->remainingSplits() > 1; + Modified |= Split; // FIXME: this check should be performed on all split levels if (DeviceGlobals) - ScopedSplitter->verifyNoCrossModuleDeviceGlobalUsage(); - - const bool SplitByScope = ScopedSplitter->remainingSplits() > 1; - bool SplitByOptionalFeatures = false; - - while (ScopedSplitter->hasMoreSplits()) { - module_split::ModuleDesc MD = ScopedSplitter->nextSplit(); - - if (IROutputOnly || SplitMode == module_split::SPLIT_NONE) { - // We can't perform any kind of split. - TopLevelModules.emplace_back(std::move(MD)); - continue; - } - - std::unique_ptr OptionalFeaturesSplitter = - module_split::getSplitterByOptionalFeatures( - std::move(MD), EmitOnlyKernelsAsEntryPoints); - - // Here we perform second-level splitting based on device-specific - // features used/declared in entry points. - // This step is mandatory, because it is required for functional - // correctness, i.e. to prevent speculative compilation of kernels that use - // optional features on a HW which doesn't support them. - SplitByOptionalFeatures |= OptionalFeaturesSplitter->remainingSplits() > 1; - - while (OptionalFeaturesSplitter->hasMoreSplits()) { - TopLevelModules.emplace_back(OptionalFeaturesSplitter->nextSplit()); - } - } - - Modified |= SplitByScope; - Modified |= SplitByOptionalFeatures; - - // TODO this nested splitting scheme will not scale well when other split - // "dimensions" will be added. Some infra/"split manager" needs to be - // implemented in this case - e.g. all needed splitters are registered, then - // split manager applies them in the order added and runs needed tforms on the - // "leaf" ModuleDesc's resulted from splitting. Some bookkeeping is needed for - // ESIMD splitter to link back needed modules. + Splitter->verifyNoCrossModuleDeviceGlobalUsage(); - // Based on results from the top-level splitting, we perform some lower-level - // splitting for various unique features. - for (module_split::ModuleDesc &MDesc : TopLevelModules) { + while (Splitter->hasMoreSplits()) { + module_split::ModuleDesc MDesc = Splitter->nextSplit(); DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1); MDesc.fixupLinkageOfDirectInvokeSimdTargets(); @@ -829,7 +788,7 @@ processInputModule(std::unique_ptr M) { const bool SplitByESIMD = ESIMDSplitter->remainingSplits() > 1; Modified |= SplitByESIMD; - if (SplitByESIMD && SplitByScope && + if (SplitByESIMD && Split && (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 @@ -874,8 +833,7 @@ processInputModule(std::unique_ptr M) { Modified = true; } - bool SplitOccurred = - SplitByScope || SplitByESIMD || SplitByOptionalFeatures; + bool SplitOccurred = Split || SplitByESIMD; if (IROutputOnly) { if (SplitOccurred) { From acc36d0dc7d78241a3e2dcbef82a27343ed8c591 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 06:02:57 -0400 Subject: [PATCH 03/33] fixes --- .../split-with-kernel-declarations.ll | 6 ++--- ...bal_variable_many_kernels_in_one_module.ll | 4 ++-- ...bal_variable_many_modules_no_dev_global.ll | 6 ++--- ..._variable_many_modules_no_dev_img_scope.ll | 6 ++--- ...lobal_variable_many_modules_two_vars_ok.ll | 4 ++-- .../device-requirements/aspects.ll | 4 ++-- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 21 +++++++++++++++--- llvm/tools/sycl-post-link/ModuleSplitter.h | 5 ++++- llvm/tools/sycl-post-link/sycl-post-link.cpp | 22 ++++++++++++++----- 9 files changed, 54 insertions(+), 24 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-code-split/split-with-kernel-declarations.ll b/llvm/test/tools/sycl-post-link/device-code-split/split-with-kernel-declarations.ll index 9b56190a71d06..595427a786e7b 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/split-with-kernel-declarations.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/split-with-kernel-declarations.ll @@ -8,9 +8,9 @@ ; ; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t1.table ; RUN: FileCheck %s -input-file=%t1.table --check-prefix CHECK-PER-KERNEL-TABLE -; RUN: FileCheck %s -input-file=%t1_0.sym --check-prefix CHECK-PER-KERNEL-SYM0 -; RUN: FileCheck %s -input-file=%t1_1.sym --check-prefix CHECK-PER-KERNEL-SYM1 -; RUN: FileCheck %s -input-file=%t1_2.sym --check-prefix CHECK-PER-KERNEL-SYM2 +; RUN: FileCheck %s -input-file=%t1_0.sym --check-prefix CHECK-PER-KERNEL-SYM1 +; RUN: FileCheck %s -input-file=%t1_1.sym --check-prefix CHECK-PER-KERNEL-SYM2 +; RUN: FileCheck %s -input-file=%t1_2.sym --check-prefix CHECK-PER-KERNEL-SYM0 ; With per-source split, there should be two device images ; CHECK-PER-SOURCE-TABLE: [Code|Properties|Symbols] diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll index 02bc8a42def25..03ccc2f5f696b 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll @@ -1,6 +1,6 @@ ; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table -; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0 -; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD1 +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 ; This test is intended to check that sycl-post-link generates no errors ; when a device global variable with the 'device_image_scope' property diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll index fd141fbaead7a..3b94a48cb3e90 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll @@ -1,7 +1,7 @@ ; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table -; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0 -; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 -; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD2 +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD2 +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 +; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD1 ; This test is intended to check that sycl-post-link generates no error if the ; 'device_image_scope' property is attached to not a device global variable. diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll index 5162e0f0456d3..847ecbc2b102d 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll @@ -1,7 +1,7 @@ ; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table -; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0 -; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 -; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD2 +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD2 +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 +; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD1 ; ModuleID = 'llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll' source_filename = "llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll" diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll index b06ad62371ae4..7e63a72c8a358 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll @@ -1,6 +1,6 @@ ; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table -; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0 -; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD1 +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 ; This test is intended to check that sycl-post-link generates no errors ; when each device global variable with the 'device_image_scope' property diff --git a/llvm/test/tools/sycl-post-link/device-requirements/aspects.ll b/llvm/test/tools/sycl-post-link/device-requirements/aspects.ll index 730b97e139748..98e92650e516f 100644 --- a/llvm/test/tools/sycl-post-link/device-requirements/aspects.ll +++ b/llvm/test/tools/sycl-post-link/device-requirements/aspects.ll @@ -17,8 +17,8 @@ ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP-AUTO-SPLIT ; RUN: sycl-post-link -split=kernel < %s -o %t.files.table -; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-0 -; RUN: FileCheck %s -input-file=%t.files_1.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-1 +; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-1 +; RUN: FileCheck %s -input-file=%t.files_1.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-0 ; CHECK-PROP-AUTO-SPLIT: [SYCL/device requirements] ; CHECK-PROP-AUTO-SPLIT-NEXT: aspects=2|gCAAAAAAAAAAAAAABAAAAYAAAAQCAAAAMAAAAA diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 5b4ea3dad666e..5591cefd37c0f 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -998,12 +998,18 @@ getSplitterByRules(ModuleDesc &&MD, const DeviceCodeSplitRulesBuilder &Rules, } std::unique_ptr -getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, +getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, bool EmitOnlyKernelsAsEntryPoints) { + if (IROutputOnly && SPLIT_NONE == Mode) { + EntryPointGroupVec Groups; + Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); + return std::make_unique(std::move(MD), std::move(Groups)); + } + DeviceCodeSplitRulesBuilder RulesBuilder; EntryPointsGroupScope Scope = selectDeviceCodeGroupScope( - MD.getModule(), Mode, /* AutoSplitIsGlobalScope */ false); + MD.getModule(), Mode, IROutputOnly); if (Scope == Scope_Global) { // We simply perform entry points filtering, but group all of them together. @@ -1063,12 +1069,21 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, Groups.emplace_back(It.first, std::move(EntryPoints), MDProps); } } + bool DoSplit = (Mode != SPLIT_NONE && + (Groups.size() > 1 || !Groups.cbegin()->Functions.empty())); - if (Groups.size() > 1) + if (DoSplit) return std::make_unique(std::move(MD), std::move(Groups)); else return std::make_unique(std::move(MD), std::move(Groups)); } +std::unique_ptr +getModuleCopier(ModuleDesc &&MD) { + EntryPointGroupVec Groups; + Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); + return std::make_unique(std::move(MD), std::move(Groups)); +} + } // namespace module_split } // namespace llvm diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index 1d0333560e77c..3223b2b762f31 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -358,7 +358,7 @@ getSplitterByRules(ModuleDesc &&MD, const DeviceCodeSplitRulesBuilder &Rules, bool EmitOnlyKernelsAsEntryPoints); std::unique_ptr -getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, +getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, bool EmitOnlyKernelsAsEntryPoints); std::unique_ptr @@ -373,6 +373,9 @@ std::unique_ptr getSplitterByOptionalFeatures(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); +std::unique_ptr +getModuleCopier(ModuleDesc &&MD); + #ifndef NDEBUG void dumpEntryPoints(const EntryPointSet &C, const char *msg = "", int Tab = 0); void dumpEntryPoints(const Module &M, bool OnlyKernelsAreEntryPoints = false, diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 108cd95525b4c..282e17e95eff7 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -451,8 +451,20 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, if (MD.isESIMD()) { PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isEsimdImage", true}); } - if (MD.isLargeGRF()) - PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true}); + + { + // check for large GRF property + bool HasLargeGRF = false; + for (const auto *F : MD.entries()) { + if (F->hasFnAttribute(::sycl::kernel_props::ATTR_LARGE_GRF)) { + HasLargeGRF = true; + break; + } + } + + if (HasLargeGRF) + PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true}); + } { std::vector FuncNames = getKernelNamesUsingAssert(M); for (const StringRef &FName : FuncNames) @@ -759,10 +771,10 @@ processInputModule(std::unique_ptr M) { (SplitMode == module_split::SPLIT_AUTO)) && "invalid split mode for IR-only output"); - // FIXME: handle IROutputOnly std::unique_ptr Splitter = - module_split::getDeviceCodeSplitter(module_split::ModuleDesc{std::move(M)}, - SplitMode, EmitOnlyKernelsAsEntryPoints); + module_split::getDeviceCodeSplitter( + module_split::ModuleDesc{std::move(M)}, SplitMode, IROutputOnly, + EmitOnlyKernelsAsEntryPoints); const bool Split = Splitter->remainingSplits() > 1; Modified |= Split; From 89212149f088910601f95a4458d84523f2eb7434 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 06:47:07 -0400 Subject: [PATCH 04/33] remaining fixes --- .../device-code-split/per-aspect-split-2.ll | 8 +++---- .../per-reqd-wg-size-split-1.ll | 16 ++++++------- .../per-reqd-wg-size-split-2.ll | 18 +++++++-------- ...lobal_variable_many_modules_two_vars_ok.ll | 2 +- .../sycl-post-link/emit_exported_symbols.ll | 4 ++-- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 22 +++++++++++++++++- llvm/tools/sycl-post-link/ModuleSplitter.h | 23 ++++++++++++++++++- 7 files changed, 67 insertions(+), 26 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll index 83d32c1688608..773424fa91fcb 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll @@ -21,12 +21,12 @@ ; CHECK-TABLE-NEXT: _2.sym ; CHECK-TABLE-EMPTY: -; CHECK-M0-SYMS: kernel0 +; CHECK-M0-SYMS: kernel3 -; CHECK-M1-SYMS: kernel3 +; CHECK-M1-SYMS: kernel1 +; CHECK-M1-SYMS: kernel2 -; CHECK-M2-SYMS: kernel1 -; CHECK-M2-SYMS: kernel2 +; CHECK-M2-SYMS: kernel0 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" target triple = "spir64-unknown-linux" diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-1.ll index c9fe37276faae..fa5ffe782a7db 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-1.ll @@ -10,29 +10,29 @@ ; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M1-IR \ +; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \ +; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M1-SYMS \ +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2 -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M2-SYMS \ +; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 ; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M1-IR \ +; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2 -; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \ +; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M1-SYMS \ +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2 -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M2-SYMS \ +; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 ; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll index c724ca3284909..cb38a596a7ba9 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-reqd-wg-size-split-2.ll @@ -5,15 +5,15 @@ ; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE ; ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \ -; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 \ -; RUN: --implicit-check-not kernel2 +; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2 ; -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M2-SYMS \ -; RUN: --implicit-check-not kernel0 --implicit-check-not kernel3 +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M1-SYMS \ +; RUN: --implicit-check-not kernel1 --implicit-check-not kernel3 \ +; RUN: --implicit-check-not kernel2 ; -; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M1-SYMS \ +; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 \ -; RUN: --implicit-check-not kernel3 +; RUN: --implicit-check-not kernel0 ; CHECK-TABLE: Code ; CHECK-TABLE-NEXT: _0.sym @@ -21,12 +21,12 @@ ; CHECK-TABLE-NEXT: _2.sym ; CHECK-TABLE-EMPTY: -; CHECK-M0-SYMS: kernel3 +; CHECK-M0-SYMS: kernel1 +; CHECK-M0-SYMS: kernel2 ; CHECK-M1-SYMS: kernel0 -; CHECK-M2-SYMS: kernel1 -; CHECK-M2-SYMS: kernel2 +; CHECK-M2-SYMS: kernel3 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" target triple = "spir64-unknown-linux" diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll index 7e63a72c8a358..9afd208726d79 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll @@ -1,6 +1,6 @@ ; RUN: sycl-post-link --device-globals --split=source -S < %s -o %t.files.table -; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD1 ; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD0 +; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD1 ; This test is intended to check that sycl-post-link generates no errors ; when each device global variable with the 'device_image_scope' property diff --git a/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll b/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll index 7d660fd12bbb6..49a94eefa2575 100644 --- a/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll +++ b/llvm/test/tools/sycl-post-link/emit_exported_symbols.ll @@ -6,8 +6,8 @@ ; ; Per-module split ; RUN: sycl-post-link -symbols -split=source -emit-exported-symbols -S < %s -o %t.per_module.files.table -; RUN: FileCheck %s -input-file=%t.per_module.files_0.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-0-PROP -; RUN: FileCheck %s -input-file=%t.per_module.files_1.prop -implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP +; RUN: FileCheck %s -input-file=%t.per_module.files_0.prop -implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP +; RUN: FileCheck %s -input-file=%t.per_module.files_1.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-0-PROP ; RUN: FileCheck %s -input-file=%t.per_module.files_2.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-2-PROP ; ; Per-kernel split diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 5591cefd37c0f..8ab40d1f87c81 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -908,6 +908,11 @@ void DeviceCodeSplitRulesBuilder::registerListOfIntegersInMetadataRule( Rules.push_back(Rule::get(MetadataName)); } +void DeviceCodeSplitRulesBuilder::registerListOfIntegersInMetadataSortedRule( + StringRef MetadataName) { + Rules.push_back(Rule::get(MetadataName)); +} + std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { std::string Result; for (const auto &R : Rules) { @@ -939,6 +944,21 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { mdconst::extract(MDOp)->getZExtValue()); } } break; + case RuleKind::SORTED_INTEGERS_LIST_METADATA: { + auto MetadataName = R.getSortedIntegersListMetadataRuleData().MetadataName; + if (F->hasMetadata(MetadataName)) { + auto *MDN = F->getMetadata(MetadataName); + + SmallVector Values; + for (const MDOperand &MDOp : MDN->operands()) + Values.push_back(mdconst::extract(MDOp)->getZExtValue()); + + llvm::sort(Values); + + for (unsigned V : Values) + Result += std::to_string(V); + } + } break; case RuleKind::FLAG_ATTR: { auto AttrName = R.getFlagAttributeRuleData().AttrName; if (F->hasFnAttribute(AttrName)) { @@ -1031,7 +1051,7 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, // Optional features RulesBuilder.registerSimpleFlagAttributeRule(::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf"); - RulesBuilder.registerListOfIntegersInMetadataRule("sycl_used_aspects"); + RulesBuilder.registerListOfIntegersInMetadataSortedRule("sycl_used_aspects"); RulesBuilder.registerListOfIntegersInMetadataRule("reqd_work_group_size"); } else { diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index 3223b2b762f31..f5b2e9ad4c992 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -270,9 +270,18 @@ class DeviceCodeSplitRulesBuilder { // into strings listed in a metadata to a resulting identifier. void registerListOfIntegersInMetadataRule(StringRef); + // Creates a rule, which adds a list of sorted dash-separated integers + // converted into strings listed in a metadata to a resulting identifier. + void registerListOfIntegersInMetadataSortedRule(StringRef); + private: enum class RuleKind { - CALLBACK, SIMPLE_STRING_ATTR, FLAG_METADATA, INTEGERS_LIST_METADATA, FLAG_ATTR + CALLBACK, + SIMPLE_STRING_ATTR, + FLAG_METADATA, + INTEGERS_LIST_METADATA, + FLAG_ATTR, + SORTED_INTEGERS_LIST_METADATA }; struct CallbackRuleData { @@ -305,12 +314,19 @@ class DeviceCodeSplitRulesBuilder { StringRef MetadataName = ""; }; + struct SortedIntegersListMetadataRuleData { + constexpr static auto Kind = RuleKind::SORTED_INTEGERS_LIST_METADATA; + SortedIntegersListMetadataRuleData() = default; + StringRef MetadataName = ""; + }; + struct Rule { private: std::array Storage; public: @@ -339,6 +355,11 @@ class DeviceCodeSplitRulesBuilder { return *reinterpret_cast(Storage.data()); } + SortedIntegersListMetadataRuleData getSortedIntegersListMetadataRuleData() const { + assert(Kind == RuleKind::SORTED_INTEGERS_LIST_METADATA); + return *reinterpret_cast(Storage.data()); + } + IntegersListMetadataRuleData getIntegersListMetadataRuleData() const { assert(Kind == RuleKind::INTEGERS_LIST_METADATA); return *reinterpret_cast(Storage.data()); From 74577a3504e32843d31d3fd761b981ad847b5b58 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 06:50:47 -0400 Subject: [PATCH 05/33] remove dead code --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 263 ------------------- llvm/tools/sycl-post-link/ModuleSplitter.h | 12 - 2 files changed, 275 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 8ab40d1f87c81..36c96304c6b3f 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -199,68 +199,6 @@ groupEntryPointsByKernelType(ModuleDesc &MD, return EntryPointGroups; } -// This function decides how entry points of the input module M will be -// distributed ("split") into multiple modules based on the command options and -// IR attributes. The decision is recorded in the output vector EntryPointGroups -// which contains pairs of group id and entry points for that group. Each such -// group along with IR it depends on (globals, functions from its call graph, -// ...) will constitute a separate module. -EntryPointGroupVec groupEntryPointsByScope(ModuleDesc &MD, - EntryPointsGroupScope EntryScope, - bool EmitOnlyKernelsAsEntryPoints) { - EntryPointGroupVec EntryPointGroups{}; - // Use MapVector for deterministic order of traversal (helps tests). - MapVector EntryPointMap; - Module &M = MD.getModule(); - - // Only process module entry points: - for (Function &F : M.functions()) { - if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints) || - !MD.isEntryPointCandidate(F)) - continue; - - switch (EntryScope) { - case Scope_PerKernel: - EntryPointMap[F.getName()].insert(&F); - break; - - case Scope_PerModule: { - if (!llvm::sycl::utils::isSYCLExternalFunction(&F)) - // TODO It may make sense to group all entry points w/o the attribute - // into a separate module rather than issuing an error. Should probably - // be controlled by an option. - error("no '" + Twine(llvm::sycl::utils::ATTR_SYCL_MODULE_ID) + - "' attribute for entry point '" + F.getName() + - "', per-module split is not possible"); - - Attribute Id = F.getFnAttribute(llvm::sycl::utils::ATTR_SYCL_MODULE_ID); - StringRef Val = Id.getValueAsString(); - EntryPointMap[Val].insert(&F); - break; - } - - case Scope_Global: - // the map key is not significant here - EntryPointMap[GLOBAL_SCOPE_NAME].insert(&F); - break; - } - } - - if (!EntryPointMap.empty()) { - EntryPointGroups.reserve(EntryPointMap.size()); - for (auto &EPG : EntryPointMap) { - EntryPointGroups.emplace_back(EPG.first, std::move(EPG.second), - MD.getEntryPointGroup().Props); - EntryPointGroup &G = EntryPointGroups.back(); - G.Props.Scope = EntryScope; - } - } else { - // No entry points met, record this. - EntryPointGroups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); - } - return EntryPointGroups; -} - // Represents a call graph between functions in a module. Nodes are functions, // edges are "calls" relation. class CallGraph { @@ -423,24 +361,6 @@ getSplitterByKernelType(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) { return std::make_unique(std::move(MD), std::move(Groups)); } -std::unique_ptr -getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode, - bool AutoSplitIsGlobalScope, - bool EmitOnlyKernelsAsEntryPoints) { - EntryPointsGroupScope Scope = - selectDeviceCodeGroupScope(MD.getModule(), Mode, AutoSplitIsGlobalScope); - EntryPointGroupVec Groups = - groupEntryPointsByScope(MD, Scope, EmitOnlyKernelsAsEntryPoints); - assert(!Groups.empty() && "At least one group is expected"); - bool DoSplit = (Mode != SPLIT_NONE && - (Groups.size() > 1 || !Groups.cbegin()->Functions.empty())); - - if (DoSplit) - return std::make_unique(std::move(MD), std::move(Groups)); - else - return std::make_unique(std::move(MD), std::move(Groups)); -} - void ModuleSplitterBase::verifyNoCrossModuleDeviceGlobalUsage() { const Module &M = getInputModule(); // Early exit if there is only one group @@ -705,182 +625,6 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, }); } -namespace { -// Data structure, which represent a combination of all possible optional -// features used in a function. -// -// It has extra methods to be useable as a key in llvm::DenseMap. -struct UsedOptionalFeatures { - SmallVector Aspects; - bool UsesLargeGRF = false; - SmallVector ReqdWorkGroupSize; - // TODO: extend this further with reqd-sub-group-size and other properties - - UsedOptionalFeatures() = default; - - UsedOptionalFeatures(const Function *F) { - if (const MDNode *MDN = F->getMetadata("sycl_used_aspects")) { - auto ExtractIntegerFromMDNodeOperand = [=](const MDOperand &N) { - Constant *C = cast(N.get())->getValue(); - return C->getUniqueInteger().getSExtValue(); - }; - - // !sycl_used_aspects is supposed to contain unique values, no duplicates - // are expected here - llvm::transform(MDN->operands(), std::back_inserter(Aspects), - ExtractIntegerFromMDNodeOperand); - llvm::sort(Aspects); - } - - if (F->hasFnAttribute(::sycl::kernel_props::ATTR_LARGE_GRF)) - UsesLargeGRF = true; - - if (const MDNode *MDN = F->getMetadata("reqd_work_group_size")) { - size_t NumOperands = MDN->getNumOperands(); - assert(NumOperands >= 1 && NumOperands <= 3 && - "reqd_work_group_size does not have between 1 and 3 operands."); - ReqdWorkGroupSize.reserve(NumOperands); - for (const MDOperand &MDOp : MDN->operands()) - ReqdWorkGroupSize.push_back( - mdconst::extract(MDOp)->getZExtValue()); - } - - llvm::hash_code AspectsHash = - llvm::hash_combine_range(Aspects.begin(), Aspects.end()); - llvm::hash_code LargeGRFHash = llvm::hash_value(UsesLargeGRF); - llvm::hash_code ReqdWorkGroupSizeHash = llvm::hash_combine_range( - ReqdWorkGroupSize.begin(), ReqdWorkGroupSize.end()); - Hash = static_cast( - llvm::hash_combine(AspectsHash, LargeGRFHash, ReqdWorkGroupSizeHash)); - } - - std::string generateModuleName(StringRef BaseName) const { - std::string Ret = BaseName.str(); - if (!ReqdWorkGroupSize.empty()) { - Ret += "-reqd-wg-size"; - for (int V : ReqdWorkGroupSize) - Ret += "-" + std::to_string(V); - } - - if (Aspects.empty()) - return Ret + "-no-aspects"; - - Ret += "-aspects"; - for (int A : Aspects) { - Ret += "-" + std::to_string(A); - } - - if (UsesLargeGRF) - Ret += "-large-grf"; - - return Ret; - } - - static UsedOptionalFeatures getTombstone() { - UsedOptionalFeatures Ret; - Ret.IsTombstoneKey = true; - return Ret; - } - - static UsedOptionalFeatures getEmpty() { - UsedOptionalFeatures Ret; - Ret.IsEmpty = true; - return Ret; - } - -private: - // For DenseMap: - llvm::hash_code Hash = {}; - bool IsTombstoneKey = false; - bool IsEmpty = false; - -public: - bool operator==(const UsedOptionalFeatures &Other) const { - // Tombstone does not compare equal to any other item - if (IsTombstoneKey || Other.IsTombstoneKey) - return false; - - if (Aspects.size() != Other.Aspects.size()) - return false; - - for (size_t I = 0, E = Aspects.size(); I != E; ++I) { - if (Aspects[I] != Other.Aspects[I]) - return false; - } - - return IsEmpty == Other.IsEmpty && UsesLargeGRF == Other.UsesLargeGRF; - } - - unsigned hash() const { return static_cast(Hash); } -}; - -struct UsedOptionalFeaturesAsKeyInfo { - static inline UsedOptionalFeatures getEmptyKey() { - return UsedOptionalFeatures::getEmpty(); - } - - static inline UsedOptionalFeatures getTombstoneKey() { - return UsedOptionalFeatures::getTombstone(); - } - - static unsigned getHashValue(const UsedOptionalFeatures &Value) { - return Value.hash(); - } - - static bool isEqual(const UsedOptionalFeatures &LHS, - const UsedOptionalFeatures &RHS) { - return LHS == RHS; - } -}; -} // namespace - -std::unique_ptr -getSplitterByOptionalFeatures(ModuleDesc &&MD, - bool EmitOnlyKernelsAsEntryPoints) { - EntryPointGroupVec Groups; - - DenseMap - PropertiesToFunctionsMap; - - Module &M = MD.getModule(); - - // Only process module entry points: - for (auto &F : M.functions()) { - if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints) || - !MD.isEntryPointCandidate(F)) { - continue; - } - - auto Key = UsedOptionalFeatures(&F); - PropertiesToFunctionsMap[std::move(Key)].insert(&F); - } - - if (PropertiesToFunctionsMap.empty()) { - // No entry points met, record this. - Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); - } else { - Groups.reserve(PropertiesToFunctionsMap.size()); - for (auto &It : PropertiesToFunctionsMap) { - const UsedOptionalFeatures &Features = It.first; - EntryPointSet &EntryPoints = It.second; - - // Start with properties of a source module - EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; - // Propagate LargeGRF flag to entry points group - if (Features.UsesLargeGRF) - MDProps.UsesLargeGRF = true; - Groups.emplace_back( - Features.generateModuleName(MD.getEntryPointGroup().GroupId), - std::move(EntryPoints), MDProps); - } - } - - if (Groups.size() > 1) - return std::make_unique(std::move(MD), std::move(Groups)); - else - return std::make_unique(std::move(MD), std::move(Groups)); -} - void DeviceCodeSplitRulesBuilder::registerRule( const std::function &Callback) { Rules.push_back(Rule::get(Callback)); @@ -1098,12 +842,5 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, return std::make_unique(std::move(MD), std::move(Groups)); } -std::unique_ptr -getModuleCopier(ModuleDesc &&MD) { - EntryPointGroupVec Groups; - Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); - return std::make_unique(std::move(MD), std::move(Groups)); -} - } // namespace module_split } // namespace llvm diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index f5b2e9ad4c992..c01dbb721914b 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -385,18 +385,6 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, std::unique_ptr getSplitterByKernelType(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); -std::unique_ptr -getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode, - bool AutoSplitIsGlobalScope, - bool EmitOnlyKernelsAsEntryPoints); - -std::unique_ptr -getSplitterByOptionalFeatures(ModuleDesc &&MD, - bool EmitOnlyKernelsAsEntryPoints); - -std::unique_ptr -getModuleCopier(ModuleDesc &&MD); - #ifndef NDEBUG void dumpEntryPoints(const EntryPointSet &C, const char *msg = "", int Tab = 0); void dumpEntryPoints(const Module &M, bool OnlyKernelsAreEntryPoints = false, From 0108a36b0290051415ba51379217a6cf43324e03 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 07:10:11 -0400 Subject: [PATCH 06/33] move some code around --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 174 ++++++++++++++----- llvm/tools/sycl-post-link/ModuleSplitter.h | 134 -------------- 2 files changed, 130 insertions(+), 178 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 36c96304c6b3f..7632ef67c015c 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -625,6 +625,136 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, }); } +class DeviceCodeSplitRulesBuilder { +public: + DeviceCodeSplitRulesBuilder() = default; + + std::string executeRules(Function *) const; + + // Accepts a callback, which should return a string based on provided + // function, which will be used as an entry points group identifier. + void registerRule(const std::function &); + + // Creates a simple rule, which adds a value of a string attribute into a + // resulting identifier. + void registerSimpleStringAttributeRule(StringRef); + + // Creates a simple rule, which adds one or another value to a resulting + // identifier based on a presence of a metadata on a function. + void registerSimpleFlagAttributeRule(StringRef, StringRef, StringRef = ""); + + // Creates a simple rule, which adds one or another value to a resulting + // identifier based on a presence of a metadata on a function. + void registerSimpleFlagMetadataRule(StringRef, StringRef, StringRef); + + // Creates a rule, which adds a list of dash-separated integers converted + // into strings listed in a metadata to a resulting identifier. + void registerListOfIntegersInMetadataRule(StringRef); + + // Creates a rule, which adds a list of sorted dash-separated integers + // converted into strings listed in a metadata to a resulting identifier. + void registerListOfIntegersInMetadataSortedRule(StringRef); + +private: + enum class RuleKind { + CALLBACK, + SIMPLE_STRING_ATTR, + FLAG_METADATA, + INTEGERS_LIST_METADATA, + FLAG_ATTR, + SORTED_INTEGERS_LIST_METADATA + }; + + struct CallbackRuleData { + constexpr static auto Kind = RuleKind::CALLBACK; + CallbackRuleData() = default; + std::function Callback = nullptr; + }; + + struct SimpleStringAttrRuleData { + constexpr static auto Kind = RuleKind::SIMPLE_STRING_ATTR; + SimpleStringAttrRuleData() = default; + StringRef Attr; + }; + + struct FlagMetadataRuleData { + constexpr static auto Kind = RuleKind::FLAG_METADATA; + FlagMetadataRuleData() = default; + StringRef TrueStr, FalseStr, MetadataName; + }; + + struct FlagAttributeRuleData { + constexpr static auto Kind = RuleKind::FLAG_ATTR; + FlagAttributeRuleData() = default; + StringRef TrueStr, FalseStr, AttrName; + }; + + struct IntegersListMetadataRuleData { + constexpr static auto Kind = RuleKind::INTEGERS_LIST_METADATA; + IntegersListMetadataRuleData() = default; + StringRef MetadataName = ""; + }; + + struct SortedIntegersListMetadataRuleData { + constexpr static auto Kind = RuleKind::SORTED_INTEGERS_LIST_METADATA; + SortedIntegersListMetadataRuleData() = default; + StringRef MetadataName = ""; + }; + + struct Rule { + private: + std::array + Storage; + public: + RuleKind Kind; + + template + static Rule get(Args... args) { + Rule R; + new(R.Storage.data()) T {args...}; + R.Kind = T::Kind; + return R; + } + + CallbackRuleData getCallbackRuleData() const { + assert(Kind == RuleKind::CALLBACK); + return *reinterpret_cast(Storage.data()); + } + + SimpleStringAttrRuleData getSimpleStringAttrRuleData() const { + assert(Kind == RuleKind::SIMPLE_STRING_ATTR); + return *reinterpret_cast(Storage.data()); + } + + FlagMetadataRuleData getFlagMetadataRuleData() const { + assert(Kind == RuleKind::FLAG_METADATA); + return *reinterpret_cast(Storage.data()); + } + + SortedIntegersListMetadataRuleData getSortedIntegersListMetadataRuleData() const { + assert(Kind == RuleKind::SORTED_INTEGERS_LIST_METADATA); + return *reinterpret_cast(Storage.data()); + } + + IntegersListMetadataRuleData getIntegersListMetadataRuleData() const { + assert(Kind == RuleKind::INTEGERS_LIST_METADATA); + return *reinterpret_cast(Storage.data()); + } + + FlagAttributeRuleData getFlagAttributeRuleData() const { + assert(Kind == RuleKind::FLAG_ATTR); + return *reinterpret_cast(Storage.data()); + } + }; + + std::vector Rules; +}; + void DeviceCodeSplitRulesBuilder::registerRule( const std::function &Callback) { Rules.push_back(Rule::get(Callback)); @@ -717,50 +847,6 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { return Result; } -std::unique_ptr -getSplitterByRules(ModuleDesc &&MD, const DeviceCodeSplitRulesBuilder &Rules, - bool EmitOnlyKernelsAsEntryPoints) { - EntryPointGroupVec Groups; - - StringMap GroupNameToFunctionsMap; - - Module &M = MD.getModule(); - - // Only process module entry points: - for (auto &F : M.functions()) { - if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints) || - !MD.isEntryPointCandidate(F)) { - continue; - } - - auto Key = Rules.executeRules(&F); - GroupNameToFunctionsMap[Key].insert(&F); - } - - if (GroupNameToFunctionsMap.empty()) { - // No entry points met, record this. - Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); - } else { - Groups.reserve(GroupNameToFunctionsMap.size()); - for (auto &It : GroupNameToFunctionsMap) { - auto Name = It.getKey(); - EntryPointSet &EntryPoints = It.getValue(); - - // Start with properties of a source module - EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; - // FIXME: properly set top-level properties - // if (Features.UsesLargeGRF) - // MDProps.UsesLargeGRF = true; - Groups.emplace_back(Name, std::move(EntryPoints), MDProps); - } - } - - if (Groups.size() > 1) - return std::make_unique(std::move(MD), std::move(Groups)); - else - return std::make_unique(std::move(MD), std::move(Groups)); -} - std::unique_ptr getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, bool EmitOnlyKernelsAsEntryPoints) { diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index c01dbb721914b..84590e669671d 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -244,140 +244,6 @@ class ModuleSplitterBase { bool hasMoreSplits() const { return remainingSplits() > 0; } }; -class DeviceCodeSplitRulesBuilder { -public: - DeviceCodeSplitRulesBuilder() = default; - - std::string executeRules(Function *) const; - - // Accepts a callback, which should return a string based on provided - // function, which will be used as an entry points group identifier. - void registerRule(const std::function &); - - // Creates a simple rule, which adds a value of a string attribute into a - // resulting identifier. - void registerSimpleStringAttributeRule(StringRef); - - // Creates a simple rule, which adds one or another value to a resulting - // identifier based on a presence of a metadata on a function. - void registerSimpleFlagAttributeRule(StringRef, StringRef, StringRef = ""); - - // Creates a simple rule, which adds one or another value to a resulting - // identifier based on a presence of a metadata on a function. - void registerSimpleFlagMetadataRule(StringRef, StringRef, StringRef); - - // Creates a rule, which adds a list of dash-separated integers converted - // into strings listed in a metadata to a resulting identifier. - void registerListOfIntegersInMetadataRule(StringRef); - - // Creates a rule, which adds a list of sorted dash-separated integers - // converted into strings listed in a metadata to a resulting identifier. - void registerListOfIntegersInMetadataSortedRule(StringRef); - -private: - enum class RuleKind { - CALLBACK, - SIMPLE_STRING_ATTR, - FLAG_METADATA, - INTEGERS_LIST_METADATA, - FLAG_ATTR, - SORTED_INTEGERS_LIST_METADATA - }; - - struct CallbackRuleData { - constexpr static auto Kind = RuleKind::CALLBACK; - CallbackRuleData() = default; - std::function Callback = nullptr; - }; - - struct SimpleStringAttrRuleData { - constexpr static auto Kind = RuleKind::SIMPLE_STRING_ATTR; - SimpleStringAttrRuleData() = default; - StringRef Attr; - }; - - struct FlagMetadataRuleData { - constexpr static auto Kind = RuleKind::FLAG_METADATA; - FlagMetadataRuleData() = default; - StringRef TrueStr, FalseStr, MetadataName; - }; - - struct FlagAttributeRuleData { - constexpr static auto Kind = RuleKind::FLAG_ATTR; - FlagAttributeRuleData() = default; - StringRef TrueStr, FalseStr, AttrName; - }; - - struct IntegersListMetadataRuleData { - constexpr static auto Kind = RuleKind::INTEGERS_LIST_METADATA; - IntegersListMetadataRuleData() = default; - StringRef MetadataName = ""; - }; - - struct SortedIntegersListMetadataRuleData { - constexpr static auto Kind = RuleKind::SORTED_INTEGERS_LIST_METADATA; - SortedIntegersListMetadataRuleData() = default; - StringRef MetadataName = ""; - }; - - struct Rule { - private: - std::array - Storage; - public: - RuleKind Kind; - - template - static Rule get(Args... args) { - Rule R; - new(R.Storage.data()) T {args...}; - R.Kind = T::Kind; - return R; - } - - CallbackRuleData getCallbackRuleData() const { - assert(Kind == RuleKind::CALLBACK); - return *reinterpret_cast(Storage.data()); - } - - SimpleStringAttrRuleData getSimpleStringAttrRuleData() const { - assert(Kind == RuleKind::SIMPLE_STRING_ATTR); - return *reinterpret_cast(Storage.data()); - } - - FlagMetadataRuleData getFlagMetadataRuleData() const { - assert(Kind == RuleKind::FLAG_METADATA); - return *reinterpret_cast(Storage.data()); - } - - SortedIntegersListMetadataRuleData getSortedIntegersListMetadataRuleData() const { - assert(Kind == RuleKind::SORTED_INTEGERS_LIST_METADATA); - return *reinterpret_cast(Storage.data()); - } - - IntegersListMetadataRuleData getIntegersListMetadataRuleData() const { - assert(Kind == RuleKind::INTEGERS_LIST_METADATA); - return *reinterpret_cast(Storage.data()); - } - - FlagAttributeRuleData getFlagAttributeRuleData() const { - assert(Kind == RuleKind::FLAG_ATTR); - return *reinterpret_cast(Storage.data()); - } - }; - - std::vector Rules; -}; - -std::unique_ptr -getSplitterByRules(ModuleDesc &&MD, const DeviceCodeSplitRulesBuilder &Rules, - bool EmitOnlyKernelsAsEntryPoints); - std::unique_ptr getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, bool EmitOnlyKernelsAsEntryPoints); From 0887e06be45813c605cfcd02888849d2c67618e4 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 09:27:55 -0400 Subject: [PATCH 07/33] refactor using std::variant --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 176 +++++++------------ 1 file changed, 64 insertions(+), 112 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 7632ef67c015c..73f267a940281 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -31,6 +31,7 @@ #include #include #include +#include using namespace llvm; using namespace llvm::module_split; @@ -656,100 +657,51 @@ class DeviceCodeSplitRulesBuilder { void registerListOfIntegersInMetadataSortedRule(StringRef); private: - enum class RuleKind { - CALLBACK, - SIMPLE_STRING_ATTR, - FLAG_METADATA, - INTEGERS_LIST_METADATA, - FLAG_ATTR, - SORTED_INTEGERS_LIST_METADATA - }; - - struct CallbackRuleData { - constexpr static auto Kind = RuleKind::CALLBACK; - CallbackRuleData() = default; - std::function Callback = nullptr; - }; - - struct SimpleStringAttrRuleData { - constexpr static auto Kind = RuleKind::SIMPLE_STRING_ATTR; - SimpleStringAttrRuleData() = default; - StringRef Attr; - }; - - struct FlagMetadataRuleData { - constexpr static auto Kind = RuleKind::FLAG_METADATA; - FlagMetadataRuleData() = default; - StringRef TrueStr, FalseStr, MetadataName; - }; - - struct FlagAttributeRuleData { - constexpr static auto Kind = RuleKind::FLAG_ATTR; - FlagAttributeRuleData() = default; - StringRef TrueStr, FalseStr, AttrName; - }; - - struct IntegersListMetadataRuleData { - constexpr static auto Kind = RuleKind::INTEGERS_LIST_METADATA; - IntegersListMetadataRuleData() = default; - StringRef MetadataName = ""; - }; - - struct SortedIntegersListMetadataRuleData { - constexpr static auto Kind = RuleKind::SORTED_INTEGERS_LIST_METADATA; - SortedIntegersListMetadataRuleData() = default; - StringRef MetadataName = ""; - }; - struct Rule { - private: - std::array + private: + std::variant, + std::function> Storage; - public: - RuleKind Kind; - - template - static Rule get(Args... args) { - Rule R; - new(R.Storage.data()) T {args...}; - R.Kind = T::Kind; - return R; - } - - CallbackRuleData getCallbackRuleData() const { - assert(Kind == RuleKind::CALLBACK); - return *reinterpret_cast(Storage.data()); - } - - SimpleStringAttrRuleData getSimpleStringAttrRuleData() const { - assert(Kind == RuleKind::SIMPLE_STRING_ATTR); - return *reinterpret_cast(Storage.data()); - } - - FlagMetadataRuleData getFlagMetadataRuleData() const { - assert(Kind == RuleKind::FLAG_METADATA); - return *reinterpret_cast(Storage.data()); + public: + enum class RKind { + K_Callback, + K_SimpleStringAttribute, + K_FlagMetadata, + K_FlagAttribute, + K_IntegersListMetadata, + K_SortedIntegersListMetadata + }; + RKind Kind; + + // Returns an index into std::variant<...> Storage defined above, which + // corresponds to the specified rule Kind. + constexpr static std::size_t storage_index(RKind K) { + switch (K) { + case RKind::K_SimpleStringAttribute: + case RKind::K_IntegersListMetadata: + case RKind::K_SortedIntegersListMetadata: + return 0; + case RKind::K_Callback: + return 2; + case RKind::K_FlagMetadata: + case RKind::K_FlagAttribute: + return 1; + } + // can't use llvm_unreachable in constexpr context + return std::variant_npos; } - SortedIntegersListMetadataRuleData getSortedIntegersListMetadataRuleData() const { - assert(Kind == RuleKind::SORTED_INTEGERS_LIST_METADATA); - return *reinterpret_cast(Storage.data()); + template + auto getStorage() const { + return std::get(Storage); } - IntegersListMetadataRuleData getIntegersListMetadataRuleData() const { - assert(Kind == RuleKind::INTEGERS_LIST_METADATA); - return *reinterpret_cast(Storage.data()); + template + Rule(RKind K, Args... args) : Storage(args...), Kind(K) { + assert(storage_index(K) == Storage.index()); } - FlagAttributeRuleData getFlagAttributeRuleData() const { - assert(Kind == RuleKind::FLAG_ATTR); - return *reinterpret_cast(Storage.data()); - } + Rule(Rule&& Other) = default; }; std::vector Rules; @@ -757,60 +709,60 @@ class DeviceCodeSplitRulesBuilder { void DeviceCodeSplitRulesBuilder::registerRule( const std::function &Callback) { - Rules.push_back(Rule::get(Callback)); + Rules.emplace_back(Rule::RKind::K_Callback, Callback); } void DeviceCodeSplitRulesBuilder::registerSimpleStringAttributeRule( StringRef Attr) { - Rules.push_back(Rule::get(Attr)); + Rules.emplace_back(Rule::RKind::K_SimpleStringAttribute, Attr); } void DeviceCodeSplitRulesBuilder::registerSimpleFlagAttributeRule( StringRef Attr, StringRef TrueStr, StringRef FalseStr) { - Rules.push_back( - Rule::get(TrueStr, FalseStr, Attr)); + Rules.emplace_back( + Rule::RKind::K_FlagAttribute, std::tuple{Attr, TrueStr, FalseStr}); } void DeviceCodeSplitRulesBuilder::registerSimpleFlagMetadataRule( StringRef TrueStr, StringRef FalseStr, StringRef MetadataName) { - Rules.push_back( - Rule::get(TrueStr, FalseStr, MetadataName)); + Rules.emplace_back(Rule::RKind::K_FlagMetadata, + std::tuple{MetadataName, TrueStr, FalseStr}); } void DeviceCodeSplitRulesBuilder::registerListOfIntegersInMetadataRule( StringRef MetadataName) { - Rules.push_back(Rule::get(MetadataName)); + Rules.emplace_back(Rule::RKind::K_IntegersListMetadata, MetadataName); } void DeviceCodeSplitRulesBuilder::registerListOfIntegersInMetadataSortedRule( StringRef MetadataName) { - Rules.push_back(Rule::get(MetadataName)); + Rules.emplace_back(Rule::RKind::K_SortedIntegersListMetadata, MetadataName); } std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { std::string Result; for (const auto &R : Rules) { switch (R.Kind) { - case RuleKind::CALLBACK: - Result += R.getCallbackRuleData().Callback(F); + case Rule::RKind::K_Callback: + Result += R.getStorage()(F); break; - case RuleKind::SIMPLE_STRING_ATTR: { - auto AttrName = R.getSimpleStringAttrRuleData().Attr; + case Rule::RKind::K_SimpleStringAttribute: { + auto AttrName = R.getStorage(); if (F->hasFnAttribute(AttrName)) { auto Attr = F->getFnAttribute(AttrName); assert(Attr.isStringAttribute()); Result += Attr.getValueAsString(); } } break; - case RuleKind::FLAG_METADATA: { - auto Data = R.getFlagMetadataRuleData(); - if (F->hasMetadata(Data.MetadataName)) - Result += Data.TrueStr; + case Rule::RKind::K_FlagMetadata: { + auto Data = R.getStorage(); + if (F->hasMetadata(std::get<0>(Data))) + Result += std::get<1>(Data); else - Result += Data.FalseStr; + Result += std::get<2>(Data); } break; - case RuleKind::INTEGERS_LIST_METADATA: { - auto MetadataName = R.getIntegersListMetadataRuleData().MetadataName; + case Rule::RKind::K_IntegersListMetadata: { + auto MetadataName = R.getStorage(); if (F->hasMetadata(MetadataName)) { auto *MDN = F->getMetadata(MetadataName); for (const MDOperand &MDOp : MDN->operands()) @@ -818,8 +770,8 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { mdconst::extract(MDOp)->getZExtValue()); } } break; - case RuleKind::SORTED_INTEGERS_LIST_METADATA: { - auto MetadataName = R.getSortedIntegersListMetadataRuleData().MetadataName; + case Rule::RKind::K_SortedIntegersListMetadata: { + auto MetadataName = R.getStorage(); if (F->hasMetadata(MetadataName)) { auto *MDN = F->getMetadata(MetadataName); @@ -833,12 +785,12 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { Result += std::to_string(V); } } break; - case RuleKind::FLAG_ATTR: { - auto AttrName = R.getFlagAttributeRuleData().AttrName; - if (F->hasFnAttribute(AttrName)) { - Result += R.getFlagAttributeRuleData().TrueStr; + case Rule::RKind::K_FlagAttribute: { + auto Data = R.getStorage(); + if (F->hasFnAttribute(std::get<0>(Data))) { + Result += std::get<1>(Data); } else { - Result += R.getFlagAttributeRuleData().FalseStr; + Result += std::get<2>(Data); } } break; } From a39bba1537a689b9a43cb47655d31ab8bbcaf8cf Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 09:33:45 -0400 Subject: [PATCH 08/33] move some method definitions into class definition --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 62 ++++++++------------ 1 file changed, 24 insertions(+), 38 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 73f267a940281..5803db96e243f 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -634,27 +634,45 @@ class DeviceCodeSplitRulesBuilder { // Accepts a callback, which should return a string based on provided // function, which will be used as an entry points group identifier. - void registerRule(const std::function &); + void registerRule(const std::function &Callback) { + Rules.emplace_back(Rule::RKind::K_Callback, Callback); + } // Creates a simple rule, which adds a value of a string attribute into a // resulting identifier. - void registerSimpleStringAttributeRule(StringRef); + void registerSimpleStringAttributeRule(StringRef AttrName) { + Rules.emplace_back(Rule::RKind::K_SimpleStringAttribute, AttrName); + } // Creates a simple rule, which adds one or another value to a resulting // identifier based on a presence of a metadata on a function. - void registerSimpleFlagAttributeRule(StringRef, StringRef, StringRef = ""); + void registerSimpleFlagAttributeRule(StringRef AttrName, + StringRef IfPresentStr, + StringRef IfAbsentStr = "") { + Rules.emplace_back(Rule::RKind::K_FlagAttribute, + std::tuple{AttrName, IfPresentStr, IfAbsentStr}); + } // Creates a simple rule, which adds one or another value to a resulting // identifier based on a presence of a metadata on a function. - void registerSimpleFlagMetadataRule(StringRef, StringRef, StringRef); + void registerSimpleFlagMetadataRule(StringRef MetadataName, + StringRef IfPresentStr, + StringRef IfAbsentStr = "") { + Rules.emplace_back(Rule::RKind::K_FlagMetadata, + std::tuple{MetadataName, IfPresentStr, IfAbsentStr}); + } // Creates a rule, which adds a list of dash-separated integers converted // into strings listed in a metadata to a resulting identifier. - void registerListOfIntegersInMetadataRule(StringRef); + void registerListOfIntegersInMetadataRule(StringRef MetadataName) { + Rules.emplace_back(Rule::RKind::K_IntegersListMetadata, MetadataName); + } // Creates a rule, which adds a list of sorted dash-separated integers // converted into strings listed in a metadata to a resulting identifier. - void registerListOfIntegersInMetadataSortedRule(StringRef); + void registerListOfIntegersInMetadataSortedRule(StringRef MetadataName) { + Rules.emplace_back(Rule::RKind::K_SortedIntegersListMetadata, MetadataName); + } private: struct Rule { @@ -707,38 +725,6 @@ class DeviceCodeSplitRulesBuilder { std::vector Rules; }; -void DeviceCodeSplitRulesBuilder::registerRule( - const std::function &Callback) { - Rules.emplace_back(Rule::RKind::K_Callback, Callback); -} - -void DeviceCodeSplitRulesBuilder::registerSimpleStringAttributeRule( - StringRef Attr) { - Rules.emplace_back(Rule::RKind::K_SimpleStringAttribute, Attr); -} - -void DeviceCodeSplitRulesBuilder::registerSimpleFlagAttributeRule( - StringRef Attr, StringRef TrueStr, StringRef FalseStr) { - Rules.emplace_back( - Rule::RKind::K_FlagAttribute, std::tuple{Attr, TrueStr, FalseStr}); -} - -void DeviceCodeSplitRulesBuilder::registerSimpleFlagMetadataRule( - StringRef TrueStr, StringRef FalseStr, StringRef MetadataName) { - Rules.emplace_back(Rule::RKind::K_FlagMetadata, - std::tuple{MetadataName, TrueStr, FalseStr}); -} - -void DeviceCodeSplitRulesBuilder::registerListOfIntegersInMetadataRule( - StringRef MetadataName) { - Rules.emplace_back(Rule::RKind::K_IntegersListMetadata, MetadataName); -} - -void DeviceCodeSplitRulesBuilder::registerListOfIntegersInMetadataSortedRule( - StringRef MetadataName) { - Rules.emplace_back(Rule::RKind::K_SortedIntegersListMetadata, MetadataName); -} - std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { std::string Result; for (const auto &R : Rules) { From eab7e02686163c5b1d691eabe6ab8a8310185d3c Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 09:51:41 -0400 Subject: [PATCH 09/33] use SmallString + tiny refactorings --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 33 ++++++++++++-------- 1 file changed, 20 insertions(+), 13 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 5803db96e243f..cfb22d72faf8f 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -726,29 +726,33 @@ class DeviceCodeSplitRulesBuilder { }; std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { - std::string Result; + SmallString<256> Result; for (const auto &R : Rules) { switch (R.Kind) { case Rule::RKind::K_Callback: Result += R.getStorage()(F); break; + case Rule::RKind::K_SimpleStringAttribute: { - auto AttrName = R.getStorage(); + StringRef AttrName = R.getStorage(); if (F->hasFnAttribute(AttrName)) { - auto Attr = F->getFnAttribute(AttrName); - assert(Attr.isStringAttribute()); + Attribute Attr = F->getFnAttribute(AttrName); Result += Attr.getValueAsString(); } } break; + case Rule::RKind::K_FlagMetadata: { - auto Data = R.getStorage(); + std::tuple Data = + R.getStorage(); if (F->hasMetadata(std::get<0>(Data))) Result += std::get<1>(Data); else Result += std::get<2>(Data); } break; + case Rule::RKind::K_IntegersListMetadata: { - auto MetadataName = R.getStorage(); + StringRef MetadataName = + R.getStorage(); if (F->hasMetadata(MetadataName)) { auto *MDN = F->getMetadata(MetadataName); for (const MDOperand &MDOp : MDN->operands()) @@ -756,10 +760,12 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { mdconst::extract(MDOp)->getZExtValue()); } } break; + case Rule::RKind::K_SortedIntegersListMetadata: { - auto MetadataName = R.getStorage(); + StringRef MetadataName = + R.getStorage(); if (F->hasMetadata(MetadataName)) { - auto *MDN = F->getMetadata(MetadataName); + MDNode *MDN = F->getMetadata(MetadataName); SmallVector Values; for (const MDOperand &MDOp : MDN->operands()) @@ -771,18 +777,19 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { Result += std::to_string(V); } } break; + case Rule::RKind::K_FlagAttribute: { - auto Data = R.getStorage(); - if (F->hasFnAttribute(std::get<0>(Data))) { + std::tuple Data = + R.getStorage(); + if (F->hasFnAttribute(std::get<0>(Data))) Result += std::get<1>(Data); - } else { + else Result += std::get<2>(Data); - } } break; } } - return Result; + return (std::string)Result; } std::unique_ptr From f802bc2a2db1e04f9cbf25f3f25c89b417f95614 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 09:59:28 -0400 Subject: [PATCH 10/33] Add some comments --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index cfb22d72faf8f..d8931de6f9267 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -626,6 +626,19 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, }); } +// This is a helper class, which allows to group/categorize function based on +// provided rules. It is intended to be used in device code split +// implementation. +// +// "Rule" is a simple routine, which returns a string for an llvm::Function +// passed to it. There could be more than one rule and they are applied in order +// of their registration. Results obtained from those rules are concatenated +// together to produce the final result. +// +// There are some predefined rules for the most popular use-cases, like grouping +// functions together based on an attribute value or presence of a metadata. +// However, there is also a possibility to register a custom callback function +// as a rule, to implement custom/more complex logic. class DeviceCodeSplitRulesBuilder { public: DeviceCodeSplitRulesBuilder() = default; From 2c11625e67471c76d828c947237b8d26582f7fbd Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 10:01:45 -0400 Subject: [PATCH 11/33] cleanup some large-grf-related dead code --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 6 +----- llvm/tools/sycl-post-link/ModuleSplitter.h | 6 ------ llvm/tools/sycl-post-link/sycl-post-link.cpp | 3 --- 3 files changed, 1 insertion(+), 14 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index d8931de6f9267..d7f504f0ceb13 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -594,8 +594,7 @@ void ModuleDesc::dump() const { llvm::errs() << "split_module::ModuleDesc[" << Name << "] {\n"; llvm::errs() << " ESIMD:" << toString(EntryPoints.Props.HasESIMD) << ", SpecConstMet:" << (Props.SpecConstsMet ? "YES" : "NO") - << ", LargeGRF:" - << (EntryPoints.Props.UsesLargeGRF ? "YES" : "NO") << "\n"; + << "\n"; dumpEntryPoints(entries(), EntryPoints.GroupId.c_str(), 1); llvm::errs() << "}\n"; } @@ -871,9 +870,6 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, // Start with properties of a source module EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; - // FIXME: Propagate LargeGRF flag to entry points group - // if (Features.UsesLargeGRF) - // MDProps.UsesLargeGRF = true; Groups.emplace_back(It.first, std::move(EntryPoints), MDProps); } } diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index 84590e669671d..3652f756ce1e1 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -55,8 +55,6 @@ struct EntryPointGroup { struct Properties { // Whether all EPs are ESIMD, SYCL or there are both kinds. SyclEsimdSplitStatus HasESIMD = SyclEsimdSplitStatus::SYCL_AND_ESIMD; - // Whether any of the EPs use large GRF mode. - bool UsesLargeGRF = false; // Scope represented by EPs in a group EntryPointsGroupScope Scope = Scope_Global; @@ -65,7 +63,6 @@ struct EntryPointGroup { Res.HasESIMD = HasESIMD == Other.HasESIMD ? HasESIMD : SyclEsimdSplitStatus::SYCL_AND_ESIMD; - Res.UsesLargeGRF = UsesLargeGRF || Other.UsesLargeGRF; // Scope remains global return Res; } @@ -90,8 +87,6 @@ struct EntryPointGroup { bool isSycl() const { return Props.HasESIMD == SyclEsimdSplitStatus::SYCL_ONLY; } - // Tells if some entry points use large GRF mode. - bool isLargeGRF() const { return Props.UsesLargeGRF; } void saveNames(std::vector &Dest) const; void rebuildFromNames(const std::vector &Names, const Module &M); @@ -146,7 +141,6 @@ class ModuleDesc { bool isESIMD() const { return EntryPoints.isEsimd(); } bool isSYCL() const { return EntryPoints.isSycl(); } - bool isLargeGRF() const { return EntryPoints.isLargeGRF(); } const EntryPointSet &entries() const { return EntryPoints.Functions; } const EntryPointGroup &getEntryPointGroup() const { return EntryPoints; } diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 282e17e95eff7..118877e87ba81 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -581,9 +581,6 @@ bool lowerEsimdConstructs(module_split::ModuleDesc &MD) { // Compute the filename suffix for the module StringRef getModuleSuffix(const module_split::ModuleDesc &MD) { - if (MD.isLargeGRF()) { - return MD.isESIMD() ? "_esimd_large_grf" : "_large_grf"; - } return MD.isESIMD() ? "_esimd" : ""; } From 471967c58f1fbc6bce57dfc8c4f09e29ac918e8b Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 28 Mar 2023 10:02:38 -0400 Subject: [PATCH 12/33] a bit of clang-format --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index d7f504f0ceb13..f662aa23f5d4a 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -815,20 +815,18 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, DeviceCodeSplitRulesBuilder RulesBuilder; - EntryPointsGroupScope Scope = selectDeviceCodeGroupScope( - MD.getModule(), Mode, IROutputOnly); + EntryPointsGroupScope Scope = + selectDeviceCodeGroupScope(MD.getModule(), Mode, IROutputOnly); if (Scope == Scope_Global) { // We simply perform entry points filtering, but group all of them together. - RulesBuilder.registerRule([](Function *) -> std::string { - return GLOBAL_SCOPE_NAME; - }); + RulesBuilder.registerRule( + [](Function *) -> std::string { return GLOBAL_SCOPE_NAME; }); } else if (Scope == Scope_PerKernel) { // Per-kernel split is quite simple: every kernel goes into a separate // module and that's it, no other rules required. - RulesBuilder.registerRule([](Function *F) -> std::string { - return F->getName().str(); - }); + RulesBuilder.registerRule( + [](Function *F) -> std::string { return F->getName().str(); }); } else if (Scope == Scope_PerModule) { // The most complex case, because we should account for many other features // like aspects used in a kernel, large-grf mode, reqd-work-group-size, etc. @@ -837,8 +835,10 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, RulesBuilder.registerSimpleStringAttributeRule("sycl-module-id"); // Optional features - RulesBuilder.registerSimpleFlagAttributeRule(::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf"); - RulesBuilder.registerListOfIntegersInMetadataSortedRule("sycl_used_aspects"); + RulesBuilder.registerSimpleFlagAttributeRule( + ::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf"); + RulesBuilder.registerListOfIntegersInMetadataSortedRule( + "sycl_used_aspects"); RulesBuilder.registerListOfIntegersInMetadataRule("reqd_work_group_size"); } else { From 466e2df0593b9535246b5ae1e7326caa1a94a437 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 29 Mar 2023 08:21:24 -0400 Subject: [PATCH 13/33] Fixes to -ir-output-only flow --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 6 ------ 1 file changed, 6 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index f662aa23f5d4a..fcdc899c4fe2f 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -807,12 +807,6 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { std::unique_ptr getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, bool EmitOnlyKernelsAsEntryPoints) { - if (IROutputOnly && SPLIT_NONE == Mode) { - EntryPointGroupVec Groups; - Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); - return std::make_unique(std::move(MD), std::move(Groups)); - } - DeviceCodeSplitRulesBuilder RulesBuilder; EntryPointsGroupScope Scope = From a76e8d2ec1d4bf8109b0246fdee12098cd45497d Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 31 Mar 2023 05:42:49 -0400 Subject: [PATCH 14/33] tiny sycl-post-link cleanup --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 16 +++++++--------- 1 file changed, 7 insertions(+), 9 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 6a9cbc984de04..f009e03aa3364 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -778,10 +778,10 @@ processInputModule(std::unique_ptr M) { module_split::getDeviceCodeSplitter( module_split::ModuleDesc{std::move(M)}, SplitMode, IROutputOnly, EmitOnlyKernelsAsEntryPoints); - const bool Split = Splitter->remainingSplits() > 1; - Modified |= Split; + bool SplitOccurred = Splitter->remainingSplits() > 1; + Modified |= SplitOccurred; - // FIXME: this check should be performed on all split levels + // FIXME: this check is not performed for ESIMD splits if (DeviceGlobals) Splitter->verifyNoCrossModuleDeviceGlobalUsage(); @@ -800,11 +800,11 @@ processInputModule(std::unique_ptr M) { std::unique_ptr ESIMDSplitter = module_split::getSplitterByKernelType(std::move(MDesc), EmitOnlyKernelsAsEntryPoints); - const bool SplitByESIMD = ESIMDSplitter->remainingSplits() > 1; - Modified |= SplitByESIMD; + SplitOccurred |= ESIMDSplitter->remainingSplits() > 1; + Modified |= SplitOccurred; - if (SplitByESIMD && Split && - (SplitMode == module_split::SPLIT_PER_KERNEL) && !SplitEsimd) { + if (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. @@ -848,8 +848,6 @@ processInputModule(std::unique_ptr M) { Modified = true; } - bool SplitOccurred = Split || SplitByESIMD; - if (IROutputOnly) { if (SplitOccurred) { error("some modules had to be split, '-" + IROutputOnly.ArgStr + From 7384fe4cbae714692d127ccd21450f418686c421 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 31 Mar 2023 05:47:21 -0400 Subject: [PATCH 15/33] Some renamings --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 24 ++++++++++---------- 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index fcdc899c4fe2f..9fdb80745a0ef 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -638,11 +638,11 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, // functions together based on an attribute value or presence of a metadata. // However, there is also a possibility to register a custom callback function // as a rule, to implement custom/more complex logic. -class DeviceCodeSplitRulesBuilder { +class FunctionsCategorizer { public: - DeviceCodeSplitRulesBuilder() = default; + FunctionsCategorizer() = default; - std::string executeRules(Function *) const; + std::string computeCategoryFor(Function *) const; // Accepts a callback, which should return a string based on provided // function, which will be used as an entry points group identifier. @@ -737,7 +737,7 @@ class DeviceCodeSplitRulesBuilder { std::vector Rules; }; -std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { +std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { SmallString<256> Result; for (const auto &R : Rules) { switch (R.Kind) { @@ -807,33 +807,33 @@ std::string DeviceCodeSplitRulesBuilder::executeRules(Function *F) const { std::unique_ptr getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, bool EmitOnlyKernelsAsEntryPoints) { - DeviceCodeSplitRulesBuilder RulesBuilder; + FunctionsCategorizer Categorizer; EntryPointsGroupScope Scope = selectDeviceCodeGroupScope(MD.getModule(), Mode, IROutputOnly); if (Scope == Scope_Global) { // We simply perform entry points filtering, but group all of them together. - RulesBuilder.registerRule( + Categorizer.registerRule( [](Function *) -> std::string { return GLOBAL_SCOPE_NAME; }); } else if (Scope == Scope_PerKernel) { // Per-kernel split is quite simple: every kernel goes into a separate // module and that's it, no other rules required. - RulesBuilder.registerRule( + Categorizer.registerRule( [](Function *F) -> std::string { return F->getName().str(); }); } else if (Scope == Scope_PerModule) { // The most complex case, because we should account for many other features // like aspects used in a kernel, large-grf mode, reqd-work-group-size, etc. // This is core of per-source device code split - RulesBuilder.registerSimpleStringAttributeRule("sycl-module-id"); + Categorizer.registerSimpleStringAttributeRule("sycl-module-id"); // Optional features - RulesBuilder.registerSimpleFlagAttributeRule( + Categorizer.registerSimpleFlagAttributeRule( ::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf"); - RulesBuilder.registerListOfIntegersInMetadataSortedRule( + Categorizer.registerListOfIntegersInMetadataSortedRule( "sycl_used_aspects"); - RulesBuilder.registerListOfIntegersInMetadataRule("reqd_work_group_size"); + Categorizer.registerListOfIntegersInMetadataRule("reqd_work_group_size"); } else { llvm_unreachable("Unexpected split scope"); @@ -848,7 +848,7 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints)) continue; - std::string Key = RulesBuilder.executeRules(&F); + std::string Key = Categorizer.computeCategoryFor(&F); EntryPointsMap[std::move(Key)].insert(&F); } From 2a367012d3ea0b86fe5b41431ca1a1bddf959875 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 31 Mar 2023 05:55:16 -0400 Subject: [PATCH 16/33] Refactoring --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 31 ++++++++++---------- 1 file changed, 16 insertions(+), 15 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 9fdb80745a0ef..b07f0a329c6bd 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -812,30 +812,33 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, EntryPointsGroupScope Scope = selectDeviceCodeGroupScope(MD.getModule(), Mode, IROutputOnly); - if (Scope == Scope_Global) { + switch (Scope) { + case Scope_Global: // We simply perform entry points filtering, but group all of them together. Categorizer.registerRule( [](Function *) -> std::string { return GLOBAL_SCOPE_NAME; }); - } else if (Scope == Scope_PerKernel) { + break; + case Scope_PerKernel: // Per-kernel split is quite simple: every kernel goes into a separate // module and that's it, no other rules required. Categorizer.registerRule( [](Function *F) -> std::string { return F->getName().str(); }); - } else if (Scope == Scope_PerModule) { + break; + case Scope_PerModule: // The most complex case, because we should account for many other features // like aspects used in a kernel, large-grf mode, reqd-work-group-size, etc. // This is core of per-source device code split - Categorizer.registerSimpleStringAttributeRule("sycl-module-id"); + Categorizer.registerSimpleStringAttributeRule( + sycl::utils::ATTR_SYCL_MODULE_ID); // Optional features Categorizer.registerSimpleFlagAttributeRule( ::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf"); - Categorizer.registerListOfIntegersInMetadataSortedRule( - "sycl_used_aspects"); + Categorizer.registerListOfIntegersInMetadataSortedRule("sycl_used_aspects"); Categorizer.registerListOfIntegersInMetadataRule("reqd_work_group_size"); - - } else { + break; + default: llvm_unreachable("Unexpected split scope"); } @@ -859,14 +862,12 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{}); } else { Groups.reserve(EntryPointsMap.size()); - for (auto &It : EntryPointsMap) { - EntryPointSet &EntryPoints = It.second; - - // Start with properties of a source module - EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; - Groups.emplace_back(It.first, std::move(EntryPoints), MDProps); - } + // Start with properties of a source module + EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; + for (auto &It : EntryPointsMap) + Groups.emplace_back(It.first, std::move(It.second), MDProps); } + bool DoSplit = (Mode != SPLIT_NONE && (Groups.size() > 1 || !Groups.cbegin()->Functions.empty())); From 1528921a614bb6ec590f51199fa52a3e726acf74 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 31 Mar 2023 06:03:56 -0400 Subject: [PATCH 17/33] clang-format --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 26 ++++++++++---------- 1 file changed, 13 insertions(+), 13 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index b07f0a329c6bd..350f5d60724ff 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -692,6 +692,7 @@ class FunctionsCategorizer { std::variant, std::function> Storage; + public: enum class RKind { K_Callback, @@ -707,31 +708,30 @@ class FunctionsCategorizer { // corresponds to the specified rule Kind. constexpr static std::size_t storage_index(RKind K) { switch (K) { - case RKind::K_SimpleStringAttribute: - case RKind::K_IntegersListMetadata: - case RKind::K_SortedIntegersListMetadata: - return 0; - case RKind::K_Callback: - return 2; - case RKind::K_FlagMetadata: - case RKind::K_FlagAttribute: - return 1; + case RKind::K_SimpleStringAttribute: + case RKind::K_IntegersListMetadata: + case RKind::K_SortedIntegersListMetadata: + return 0; + case RKind::K_Callback: + return 2; + case RKind::K_FlagMetadata: + case RKind::K_FlagAttribute: + return 1; } // can't use llvm_unreachable in constexpr context return std::variant_npos; } - template - auto getStorage() const { + template auto getStorage() const { return std::get(Storage); } - template + template Rule(RKind K, Args... args) : Storage(args...), Kind(K) { assert(storage_index(K) == Storage.index()); } - Rule(Rule&& Other) = default; + Rule(Rule &&Other) = default; }; std::vector Rules; From aaf1e92b82415656faec197d9afa340fe54627bd Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 3 Apr 2023 09:41:43 -0400 Subject: [PATCH 18/33] apply comments --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 4 ++-- llvm/tools/sycl-post-link/sycl-post-link.cpp | 10 +++------- 2 files changed, 5 insertions(+), 9 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 350f5d60724ff..aa276dae4e69d 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -873,8 +873,8 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, if (DoSplit) return std::make_unique(std::move(MD), std::move(Groups)); - else - return std::make_unique(std::move(MD), std::move(Groups)); + + return std::make_unique(std::move(MD), std::move(Groups)); } } // namespace module_split diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index f009e03aa3364..e346e98deb914 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -455,13 +455,9 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, { // check for large GRF property - bool HasLargeGRF = false; - for (const auto *F : MD.entries()) { - if (F->hasFnAttribute(::sycl::kernel_props::ATTR_LARGE_GRF)) { - HasLargeGRF = true; - break; - } - } + bool HasLargeGRF = llvm::any_of(MD.entries(), [](const Function *F) { + return F->hasFnAttribute(::sycl::kernel_props::ATTR_LARGE_GRF); + }); if (HasLargeGRF) PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true}); From b4edc982f09a368d0b50a13ac1ac0dd7ba5c5e13 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 3 Apr 2023 09:45:16 -0400 Subject: [PATCH 19/33] An attempt to fix LIT tests --- llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll | 4 ++-- llvm/test/tools/sycl-post-link/sycl-large-grf.ll | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll b/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll index 919eabccfa3c4..a8bf294179d35 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll @@ -17,8 +17,8 @@ ; CHECK: [Code|Properties|Symbols] ; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym -; CHECK: {{.*}}esimd-large-grf.ll.tmp_1.ll|{{.*}}esimd-large-grf.ll.tmp_1.prop|{{.*}}esimd-large-grf.ll.tmp_1.sym -; CHECK: {{.*}}esimd-large-grf.ll.tmp_esimd_1.ll|{{.*}}esimd-large-grf.ll.tmp_esimd_1.prop|{{.*}}esimd-large-grf.ll.tmp_esimd_1.sym +; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym +; CHECK: {{.*}}_esimd_1.ll|{{.*}}_esimd_1.prop|{{.*}}_esimd_1.sym ; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1 ; CHECK-ESIMD-LargeGRF-PROP: isLargeGRF=1|1 diff --git a/llvm/test/tools/sycl-post-link/sycl-large-grf.ll b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll index 26340ddad59e9..34b98c1e5e5d9 100644 --- a/llvm/test/tools/sycl-post-link/sycl-large-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll @@ -16,7 +16,7 @@ ; CHECK: [Code|Properties|Symbols] ; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym -; CHECK: {{.*}}-large-grf.ll.tmp_1.ll|{{.*}}-large-grf.ll.tmp_1.prop|{{.*}}-large-grf.ll.tmp_1.sym +; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym ; CHECK-LARGE-GRF-PROP: isLargeGRF=1|1 From 2332b57b9eb8b5c81eabb2119f76a1c0a9c20c07 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 3 Apr 2023 12:41:14 -0400 Subject: [PATCH 20/33] Actual fix for tests --- llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll | 6 +++--- llvm/test/tools/sycl-post-link/sycl-large-grf.ll | 6 +++--- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll b/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll index a8bf294179d35..6bae075384405 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll @@ -9,11 +9,11 @@ ; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table -; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR --implicit-check-not='RegisterAllocMode' -; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP +; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR --implicit-check-not='RegisterAllocMode' +; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM ; RUN: FileCheck %s -input-file=%t_esimd_1.sym --check-prefixes CHECK-ESIMD-SYM -; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM +; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM ; CHECK: [Code|Properties|Symbols] ; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym diff --git a/llvm/test/tools/sycl-post-link/sycl-large-grf.ll b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll index 34b98c1e5e5d9..105926200b0e9 100644 --- a/llvm/test/tools/sycl-post-link/sycl-large-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll @@ -9,10 +9,10 @@ ; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t.table -; RUN: FileCheck %s -input-file=%t_large_grf_0.ll --check-prefixes CHECK-LARGE-GRF-IR -; RUN: FileCheck %s -input-file=%t_large_grf_0.prop --check-prefixes CHECK-LARGE-GRF-PROP +; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-LARGE-GRF-IR +; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-LARGE-GRF-PROP ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM -; RUN: FileCheck %s -input-file=%t_large_grf_0.sym --check-prefixes CHECK-LARGE-GRF-SYM +; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-LARGE-GRF-SYM ; CHECK: [Code|Properties|Symbols] ; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym From 13169d46b86d40196060de52bfe1dff86d75409f Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 4 Apr 2023 14:00:32 -0400 Subject: [PATCH 21/33] Add brief Rule::RKind descriptions --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index aa276dae4e69d..3a34f7e508391 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -695,11 +695,17 @@ class FunctionsCategorizer { public: enum class RKind { + // Custom callback function K_Callback, + // Copy value of the specified attribute, if present K_SimpleStringAttribute, + // Use one or another string based on the specified metadata presence K_FlagMetadata, + // Use one or another string based on the specified attribute presence K_FlagAttribute, + // Concatenate and use list of integers from the specified metadata K_IntegersListMetadata, + // Sort, concatenate and use list of integers from the specified metadata K_SortedIntegersListMetadata }; RKind Kind; From 67082dda9a95d3dabaa3c4c9158e86f8b1b5ba27 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 4 Apr 2023 14:09:45 -0400 Subject: [PATCH 22/33] Add delimeter to distinguish between strings produced by different rules --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 3a34f7e508391..1d0d97a891e50 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -774,7 +774,7 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { if (F->hasMetadata(MetadataName)) { auto *MDN = F->getMetadata(MetadataName); for (const MDOperand &MDOp : MDN->operands()) - Result += std::to_string( + Result += "-" + std::to_string( mdconst::extract(MDOp)->getZExtValue()); } } break; @@ -792,7 +792,7 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { llvm::sort(Values); for (unsigned V : Values) - Result += std::to_string(V); + Result += "-" + std::to_string(V); } } break; @@ -805,6 +805,8 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { Result += std::get<2>(Data); } break; } + + Result += "-"; } return (std::string)Result; From 54ffaeb5fe632090b9e607e60afda00f2c6c9c74 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 4 Apr 2023 14:25:21 -0400 Subject: [PATCH 23/33] clang-format --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 1d0d97a891e50..cd7208198c7d9 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -774,8 +774,9 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { if (F->hasMetadata(MetadataName)) { auto *MDN = F->getMetadata(MetadataName); for (const MDOperand &MDOp : MDN->operands()) - Result += "-" + std::to_string( - mdconst::extract(MDOp)->getZExtValue()); + Result += + "-" + std::to_string( + mdconst::extract(MDOp)->getZExtValue()); } } break; From cf3a36576a91129a36b8f3086bfc7f063855a815 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 6 Apr 2023 09:33:11 -0400 Subject: [PATCH 24/33] [NFC] A few more changes Added some comments here and there. Wrapped some code in `ModuleSplitter.cpp` into `namespace { ... }`. Removed `default` case from a fully-coverage switch: we have downstream environments where this is considered an error. --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 6 ++++-- llvm/tools/sycl-post-link/sycl-post-link.cpp | 3 +++ 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index cd7208198c7d9..0d19a6cb0523a 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -625,6 +625,7 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, }); } +namespace { // This is a helper class, which allows to group/categorize function based on // provided rules. It is intended to be used in device code split // implementation. @@ -812,6 +813,7 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { return (std::string)Result; } +} // namespace std::unique_ptr getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, @@ -842,13 +844,13 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, sycl::utils::ATTR_SYCL_MODULE_ID); // Optional features + // Note: Add more rules at the end of the list to avoid chaning orders of + // output files in existing tests. Categorizer.registerSimpleFlagAttributeRule( ::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf"); Categorizer.registerListOfIntegersInMetadataSortedRule("sycl_used_aspects"); Categorizer.registerListOfIntegersInMetadataRule("reqd_work_group_size"); break; - default: - llvm_unreachable("Unexpected split scope"); } // std::map is used here to ensure stable ordering of entry point groups, diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index e346e98deb914..d95b15d50c151 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -781,6 +781,9 @@ processInputModule(std::unique_ptr M) { if (DeviceGlobals) Splitter->verifyNoCrossModuleDeviceGlobalUsage(); + // 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. while (Splitter->hasMoreSplits()) { module_split::ModuleDesc MDesc = Splitter->nextSplit(); DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1); From 863abed65c5951fff893438947177befa43cf156 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 12 Apr 2023 06:03:29 -0400 Subject: [PATCH 25/33] Fix sycl-post-link emitting warning incorrectly --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index d95b15d50c151..a8027e6625db1 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -799,11 +799,10 @@ processInputModule(std::unique_ptr M) { std::unique_ptr ESIMDSplitter = module_split::getSplitterByKernelType(std::move(MDesc), EmitOnlyKernelsAsEntryPoints); - SplitOccurred |= ESIMDSplitter->remainingSplits() > 1; - Modified |= SplitOccurred; + bool ESIMDSplitOccurred = ESIMDSplitter->remainingSplits() > 1; - if (SplitOccurred && (SplitMode == module_split::SPLIT_PER_KERNEL) && - !SplitEsimd) { + if (ESIMDSplitOccurred && 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. @@ -812,6 +811,8 @@ processInputModule(std::unique_ptr M) { SplitEsimd.ValueStr + " must also be specified"); } SmallVector MMs; + SplitOccurred |= ESIMDSplitOccurred; + Modified |= SplitOccurred; while (ESIMDSplitter->hasMoreSplits()) { module_split::ModuleDesc MDesc2 = ESIMDSplitter->nextSplit(); From adadb0d7e99b322530c9643ac33da8827cd628a4 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 20 Apr 2023 10:08:59 -0400 Subject: [PATCH 26/33] Restore sycl-opt-level attribute functionality after merge --- llvm/include/llvm/SYCLLowerIR/SYCLUtils.h | 1 + .../tools/sycl-post-link/sycl-opt-level.ll | 9 ++++++-- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 5 +++-- llvm/tools/sycl-post-link/ModuleSplitter.h | 8 ------- llvm/tools/sycl-post-link/sycl-post-link.cpp | 21 ++++++++++++++++--- 5 files changed, 29 insertions(+), 15 deletions(-) diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h index 241859a47426d..c9ebcdae53f4b 100644 --- a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h @@ -22,6 +22,7 @@ namespace llvm { namespace sycl { namespace utils { constexpr char ATTR_SYCL_MODULE_ID[] = "sycl-module-id"; +constexpr char ATTR_SYCL_OPTLEVEL[] = "sycl-optlevel"; using CallGraphNodeAction = ::std::function; using CallGraphFunctionFilter = diff --git a/llvm/test/tools/sycl-post-link/sycl-opt-level.ll b/llvm/test/tools/sycl-post-link/sycl-opt-level.ll index 349e6e445027d..8967d3bc02621 100644 --- a/llvm/test/tools/sycl-post-link/sycl-opt-level.ll +++ b/llvm/test/tools/sycl-post-link/sycl-opt-level.ll @@ -8,13 +8,18 @@ ; RUN: FileCheck %s -input-file=%t.table ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-OPT-LEVEL-PROP-0 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-OPT-LEVEL-PROP-1 +; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM-0 +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYM-1 ; CHECK: [Code|Properties|Symbols] ; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym ; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym +; CHECK-EMPTY: -; CHECK-OPT-LEVEL-PROP-0: optLevel=1|0 -; CHECK-OPT-LEVEL-PROP-1: optLevel=1|2 +; CHECK-OPT-LEVEL-PROP-0: optLevel=1|2 +; CHECK-OPT-LEVEL-PROP-1: optLevel=1|0 +; CHECK-SYM-0: _Z3fooii +; CHECK-SYM-1: _Z3booii target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" target triple = "spir64-unknown-unknown" diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index a2d6fcd0748a7..1bf3dd2ac4d98 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -43,7 +43,6 @@ constexpr char GLOBAL_SCOPE_NAME[] = ""; constexpr char SYCL_SCOPE_NAME[] = ""; constexpr char ESIMD_SCOPE_NAME[] = ""; constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd"; -constexpr char ATTR_OPT_LEVEL[] = "sycl-optlevel"; bool hasIndirectFunctionsOrCalls(const Module &M) { for (const auto &F : M.functions()) { @@ -595,7 +594,7 @@ void ModuleDesc::dump() const { llvm::errs() << "split_module::ModuleDesc[" << Name << "] {\n"; llvm::errs() << " ESIMD:" << toString(EntryPoints.Props.HasESIMD) << ", SpecConstMet:" << (Props.SpecConstsMet ? "YES" : "NO") - << ", OptLevel:" << EntryPoints.getOptLevel() << "\n"; + << "\n"; dumpEntryPoints(entries(), EntryPoints.GroupId.c_str(), 1); llvm::errs() << "}\n"; } @@ -851,6 +850,8 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, ::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf"); Categorizer.registerListOfIntegersInMetadataSortedRule("sycl_used_aspects"); Categorizer.registerListOfIntegersInMetadataRule("reqd_work_group_size"); + Categorizer.registerSimpleStringAttributeRule( + sycl::utils::ATTR_SYCL_OPTLEVEL); break; } diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index e53bbdfcac2a1..3652f756ce1e1 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -55,8 +55,6 @@ struct EntryPointGroup { struct Properties { // Whether all EPs are ESIMD, SYCL or there are both kinds. SyclEsimdSplitStatus HasESIMD = SyclEsimdSplitStatus::SYCL_AND_ESIMD; - // front-end opt level for kernel compilation - int OptLevel = -1; // Scope represented by EPs in a group EntryPointsGroupScope Scope = Scope_Global; @@ -66,8 +64,6 @@ struct EntryPointGroup { ? HasESIMD : SyclEsimdSplitStatus::SYCL_AND_ESIMD; // Scope remains global - // OptLevel is expected to be the same for both merging EPGs - assert(OptLevel == Other.OptLevel && "OptLevels are not same"); return Res; } }; @@ -92,9 +88,6 @@ struct EntryPointGroup { return Props.HasESIMD == SyclEsimdSplitStatus::SYCL_ONLY; } - // Returns opt level - int getOptLevel() const { return Props.OptLevel; } - void saveNames(std::vector &Dest) const; void rebuildFromNames(const std::vector &Names, const Module &M); void rebuild(const Module &M); @@ -148,7 +141,6 @@ class ModuleDesc { bool isESIMD() const { return EntryPoints.isEsimd(); } bool isSYCL() const { return EntryPoints.isSycl(); } - int getOptLevel() const { return EntryPoints.getOptLevel(); } const EntryPointSet &entries() const { return EntryPoints.Functions; } const EntryPointGroup &getEntryPointGroup() const { return EntryPoints; } diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index e2384d16c2665..97548f450a5b1 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -40,6 +40,7 @@ #include "llvm/SYCLLowerIR/HostPipes.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" #include "llvm/SYCLLowerIR/LowerKernelProps.h" +#include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/InitLLVM.h" @@ -462,9 +463,23 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, if (HasLargeGRF) PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true}); } - if (MD.getOptLevel() != -1) - PropSet[PropSetRegTy::SYCL_MISC_PROP].insert( - {"optLevel", MD.getOptLevel()}); + { + // Handle sycl-optlevel property + int OptLevel = -1; + for (const Function &F : M.functions()) { + if (!F.hasFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL)) + continue; + + // getAsInteger returns true on error + if (!F.getFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL) + .getValueAsString() + .getAsInteger(10, OptLevel)) + break; + } + + if (OptLevel != -1) + PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"optLevel", OptLevel}); + } { std::vector FuncNames = getKernelNamesUsingAssert(M); for (const StringRef &FName : FuncNames) From 9158ca89088e7508c81b9f56e79f2bebf1e91dbd Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 20 Apr 2023 10:17:08 -0400 Subject: [PATCH 27/33] Fix test after merge --- .../device-code-split/per-aspect-split-1.ll | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll index 687e50e0f6ffd..faec71a602ffd 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-1.ll @@ -10,29 +10,29 @@ ; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M1-IR \ +; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \ +; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M1-SYMS \ +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M2-SYMS \ +; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M1-IR \ +; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \ +; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M1-SYMS \ +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 -; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M2-SYMS \ +; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \ ; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 ; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table From 7ff7531215698fee717b3dadcb0987b3af1dee27 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 20 Apr 2023 10:25:29 -0400 Subject: [PATCH 28/33] Apply comments --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 1bf3dd2ac4d98..d977c45ba231c 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -689,8 +689,10 @@ class FunctionsCategorizer { private: struct Rule { + using TupleOfThreeStringRef = std::tuple; + private: - std::variant, + std::variant> Storage; @@ -761,7 +763,7 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { } break; case Rule::RKind::K_FlagMetadata: { - std::tuple Data = + Rule::TupleOfThreeStringRef Data = R.getStorage(); if (F->hasMetadata(std::get<0>(Data))) Result += std::get<1>(Data); @@ -799,7 +801,7 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { } break; case Rule::RKind::K_FlagAttribute: { - std::tuple Data = + Rule::TupleOfThreeStringRef Data = R.getStorage(); if (F->hasFnAttribute(std::get<0>(Data))) Result += std::get<1>(Data); From e4b452e6142f5d1175758213a42f9821740dc85e Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 20 Apr 2023 11:03:59 -0400 Subject: [PATCH 29/33] Better solution for review comments --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 31 ++++++++++---------- 1 file changed, 16 insertions(+), 15 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index d977c45ba231c..75ef12999409b 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -663,7 +663,7 @@ class FunctionsCategorizer { StringRef IfPresentStr, StringRef IfAbsentStr = "") { Rules.emplace_back(Rule::RKind::K_FlagAttribute, - std::tuple{AttrName, IfPresentStr, IfAbsentStr}); + Rule::FlagRuleData{AttrName, IfPresentStr, IfAbsentStr}); } // Creates a simple rule, which adds one or another value to a resulting @@ -671,8 +671,9 @@ class FunctionsCategorizer { void registerSimpleFlagMetadataRule(StringRef MetadataName, StringRef IfPresentStr, StringRef IfAbsentStr = "") { - Rules.emplace_back(Rule::RKind::K_FlagMetadata, - std::tuple{MetadataName, IfPresentStr, IfAbsentStr}); + Rules.emplace_back( + Rule::RKind::K_FlagMetadata, + Rule::FlagRuleData{MetadataName, IfPresentStr, IfAbsentStr}); } // Creates a rule, which adds a list of dash-separated integers converted @@ -689,10 +690,12 @@ class FunctionsCategorizer { private: struct Rule { - using TupleOfThreeStringRef = std::tuple; + struct FlagRuleData { + StringRef Name, IfPresentStr, IfAbsentStr; + }; private: - std::variant> Storage; @@ -763,12 +766,11 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { } break; case Rule::RKind::K_FlagMetadata: { - Rule::TupleOfThreeStringRef Data = - R.getStorage(); - if (F->hasMetadata(std::get<0>(Data))) - Result += std::get<1>(Data); + Rule::FlagRuleData Data = R.getStorage(); + if (F->hasMetadata(Data.Name)) + Result += Data.IfPresentStr; else - Result += std::get<2>(Data); + Result += Data.IfAbsentStr; } break; case Rule::RKind::K_IntegersListMetadata: { @@ -801,12 +803,11 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { } break; case Rule::RKind::K_FlagAttribute: { - Rule::TupleOfThreeStringRef Data = - R.getStorage(); - if (F->hasFnAttribute(std::get<0>(Data))) - Result += std::get<1>(Data); + Rule::FlagRuleData Data = R.getStorage(); + if (F->hasFnAttribute(Data.Name)) + Result += Data.IfPresentStr; else - Result += std::get<2>(Data); + Result += Data.IfAbsentStr; } break; } From d61e93e7fac6c82f3bfdfcb5ede68ea5d96d555c Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 21 Apr 2023 03:41:57 -0400 Subject: [PATCH 30/33] Apply comments --- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 12 ++++++------ llvm/tools/sycl-post-link/sycl-post-link.cpp | 9 +++++++-- 2 files changed, 13 insertions(+), 8 deletions(-) diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 75ef12999409b..b23bb7e4c0e1d 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -658,7 +658,7 @@ class FunctionsCategorizer { } // Creates a simple rule, which adds one or another value to a resulting - // identifier based on a presence of a metadata on a function. + // identifier based on the presence of a metadata on a function. void registerSimpleFlagAttributeRule(StringRef AttrName, StringRef IfPresentStr, StringRef IfAbsentStr = "") { @@ -667,7 +667,7 @@ class FunctionsCategorizer { } // Creates a simple rule, which adds one or another value to a resulting - // identifier based on a presence of a metadata on a function. + // identifier based on the presence of a metadata on a function. void registerSimpleFlagMetadataRule(StringRef MetadataName, StringRef IfPresentStr, StringRef IfAbsentStr = "") { @@ -791,13 +791,13 @@ std::string FunctionsCategorizer::computeCategoryFor(Function *F) const { if (F->hasMetadata(MetadataName)) { MDNode *MDN = F->getMetadata(MetadataName); - SmallVector Values; + SmallVector Values; for (const MDOperand &MDOp : MDN->operands()) Values.push_back(mdconst::extract(MDOp)->getZExtValue()); llvm::sort(Values); - for (unsigned V : Values) + for (std::uint64_t V : Values) Result += "-" + std::to_string(V); } } break; @@ -880,8 +880,8 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly, Groups.reserve(EntryPointsMap.size()); // Start with properties of a source module EntryPointGroup::Properties MDProps = MD.getEntryPointGroup().Props; - for (auto &It : EntryPointsMap) - Groups.emplace_back(It.first, std::move(It.second), MDProps); + for (auto &[Key, EntryPoints] : EntryPointsMap) + Groups.emplace_back(Key, std::move(EntryPoints), MDProps); } bool DoSplit = (Mode != SPLIT_NONE && diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 97548f450a5b1..25e084812e5c7 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -473,8 +473,13 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, // getAsInteger returns true on error if (!F.getFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL) .getValueAsString() - .getAsInteger(10, OptLevel)) - break; + .getAsInteger(10, OptLevel)) { + // It is expected that device-code split has separated kernels with + // different values of sycl-optlevel attribute. Therefore, it is enough + // to only look at the first function with such attribute to compute + // the property for the whole device image. + break; + } } if (OptLevel != -1) From ec3783179a3c609033d4c202df5ef6915a588fdf Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 21 Apr 2023 03:58:07 -0400 Subject: [PATCH 31/33] Apply clang-format --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 25e084812e5c7..d4f796883fde2 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -478,7 +478,7 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, // different values of sycl-optlevel attribute. Therefore, it is enough // to only look at the first function with such attribute to compute // the property for the whole device image. - break; + break; } } From d674eb283b869c71e3858ec98f67239362961b5c Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 27 Apr 2023 05:39:49 -0400 Subject: [PATCH 32/33] Consider only entry points when emitting optLevel property --- .../sycl-opt-level-external-funcs.ll | 68 +++++++++++++++++++ llvm/tools/sycl-post-link/sycl-post-link.cpp | 6 +- 2 files changed, 71 insertions(+), 3 deletions(-) create mode 100644 llvm/test/tools/sycl-post-link/sycl-opt-level-external-funcs.ll diff --git a/llvm/test/tools/sycl-post-link/sycl-opt-level-external-funcs.ll b/llvm/test/tools/sycl-post-link/sycl-opt-level-external-funcs.ll new file mode 100644 index 0000000000000..053760fa3b86f --- /dev/null +++ b/llvm/test/tools/sycl-post-link/sycl-opt-level-external-funcs.ll @@ -0,0 +1,68 @@ +; This test checks, that 'optLevel' property is only emitted based on the module +; entry points +; +; In this test we have functions 'foo' and 'boo' defined in different +; translation units. They are both entry points and 'foo' calls 'boo'. +; As a result, we expect two modules: +; - module with 'foo' (as entry point) and 'bar' (included as dependency) with +; 'optLevel' set to 1 (taken from 'foo') +; - module with 'bar' (as entry point) with 'optLevel' set to 2 (taken from +; 'bar') + +; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table +; RUN: FileCheck %s -input-file=%t.table +; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-OPT-LEVEL-PROP-0 +; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-OPT-LEVEL-PROP-1 +; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM-0 +; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYM-1 +; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK-IR-0 +; RUN: FileCheck %s -input-file=%t_1.ll --check-prefix CHECK-IR-1 + +; CHECK: [Code|Properties|Symbols] +; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym +; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym +; CHECK-EMPTY: + +; CHECK-OPT-LEVEL-PROP-0: optLevel=1|1 +; CHECK-OPT-LEVEL-PROP-1: optLevel=1|2 +; CHECK-SYM-0: _Z3fooii +; CHECK-SYM-0-EMPTY: +; CHECK-SYM-1: _Z3barii +; +; CHECK-IR-0-DAG: define {{.*}} @_Z3fooii{{.*}} #[[#ATTR0:]] +; CHECK-IR-0-DAG: define {{.*}} @_Z3barii{{.*}} #[[#ATTR1:]] +; CHECK-IR-0-DAG: attributes #[[#ATTR0]] = { {{.*}} "sycl-optlevel"="1" } +; CHECK-IR-0-DAG: attributes #[[#ATTR1]] = { {{.*}} "sycl-optlevel"="2" } +; +; CHECK-IR-1: define {{.*}} @_Z3barii{{.*}} #[[#ATTR0:]] +; CHECK-IR-1: attributes #[[#ATTR0]] = { {{.*}} "sycl-optlevel"="2" } + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" +target triple = "spir64-unknown-unknown" + +define dso_local spir_func noundef i32 @_Z3fooii(i32 noundef %a, i32 noundef %b) local_unnamed_addr #0 { +entry: + %call = call i32 @_Z3barii(i32 %a, i32 %b) + %sub = sub nsw i32 %a, %call + ret i32 %sub +} + +define dso_local spir_func noundef i32 @_Z3barii(i32 noundef %a, i32 noundef %b) #1 { +entry: + %retval = alloca i32, align 4 + %a.addr = alloca i32, align 4 + %b.addr = alloca i32, align 4 + %retval.ascast = addrspacecast i32* %retval to i32 addrspace(4)* + %a.addr.ascast = addrspacecast i32* %a.addr to i32 addrspace(4)* + %b.addr.ascast = addrspacecast i32* %b.addr to i32 addrspace(4)* + store i32 %a, i32 addrspace(4)* %a.addr.ascast, align 4 + store i32 %b, i32 addrspace(4)* %b.addr.ascast, align 4 + %0 = load i32, i32 addrspace(4)* %a.addr.ascast, align 4 + %1 = load i32, i32 addrspace(4)* %b.addr.ascast, align 4 + %add = add nsw i32 %0, %1 + ret i32 %add +} + +attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "sycl-module-id"="test3.cpp" "sycl-optlevel"="1" } +attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "sycl-module-id"="test2.cpp" "sycl-optlevel"="2" } + diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index d4f796883fde2..d3c90adbd3767 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -466,12 +466,12 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, { // Handle sycl-optlevel property int OptLevel = -1; - for (const Function &F : M.functions()) { - if (!F.hasFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL)) + for (const Function *F : MD.entries()) { + if (!F->hasFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL)) continue; // getAsInteger returns true on error - if (!F.getFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL) + if (!F->getFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL) .getValueAsString() .getAsInteger(10, OptLevel)) { // It is expected that device-code split has separated kernels with From 0f49bfbf756c1abb041b13126ee515f5db5914c9 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 27 Apr 2023 06:12:27 -0400 Subject: [PATCH 33/33] This (rather) hacky change should help fix regressions --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index d3c90adbd3767..a676cab38fc59 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -463,7 +463,18 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, if (HasLargeGRF) PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true}); } - { + // FIXME: Remove 'if' below when possible + // GPU backend has a problem with accepting optimization level options in form + // described by Level Zero specification (-ze-opt-level=1) when 'invoke_simd' + // functionality is involved. JIT compilation results in the following error: + // error: VLD: Failed to compile SPIR-V with following error: + // invalid api option: -ze-opt-level=O1 + // -11 (PI_ERROR_BUILD_PROGRAM_FAILURE) + // 'if' below essentially preserves the behavior (presumably mistakenly) + // implemented in intel/llvm#8763: ignore 'optLevel' property for images which + // were produced my merge after ESIMD split + if (MD.getEntryPointGroup().Props.HasESIMD != + module_split::SyclEsimdSplitStatus::SYCL_AND_ESIMD) { // Handle sycl-optlevel property int OptLevel = -1; for (const Function *F : MD.entries()) {