Skip to content

[SYCL][NFCI] Unify large-grf splitting with per-aspects split #7512

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
merged 8 commits into from
Nov 29, 2022
Merged

[SYCL][NFCI] Unify large-grf splitting with per-aspects split #7512

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
4e1a7a3
Add large-grf attribute into per-aspects splitter
AlexeySachkov Nov 23, 2022
870398f
Stop using largeGRF splitter in sycl-post-link
AlexeySachkov Nov 23, 2022
5699462
Propagate info about large-grf usage from per-aspect splitter
AlexeySachkov Nov 23, 2022
cf6ffa2
Re-arrange tests a bit to match new sycl-post-link output
AlexeySachkov Nov 23, 2022
9722e9f
Remove LargeGRF splitter, because it is not used anymore
AlexeySachkov Nov 23, 2022
15c7398
Remove groupEntryPointsByAttribute because it is not used anymore
AlexeySachkov Nov 23, 2022
4539278
Fix test
AlexeySachkov Nov 23, 2022
722fdd8
[NFC] Use less auto, be more verbose with types
AlexeySachkov Nov 24, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 4 additions & 4 deletions llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-2.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,11 +5,11 @@
; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE
Copy link
Contributor

Choose a reason for hiding this comment

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

Should we add a test with aspects and large GRF to lock down the behavior?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

The thing is that as we add more features, they will affect existing tests anyway, so I would prefer to leave this as-is to have an isolated test for a piece of the implementation

;
; 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 kernel3 --implicit-check-not kernel1 \
; RUN: --implicit-check-not kernel2
;
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M1-SYMS \
; RUN: --implicit-check-not kernel3 --implicit-check-not kernel1 \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 \
; RUN: --implicit-check-not kernel2
;
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M2-SYMS \
Expand All @@ -21,9 +21,9 @@
; CHECK-TABLE-NEXT: _2.sym
; CHECK-TABLE-EMPTY:

; CHECK-M0-SYMS: kernel3
; CHECK-M0-SYMS: kernel0

; CHECK-M1-SYMS: kernel0
; CHECK-M1-SYMS: kernel3

; CHECK-M2-SYMS: kernel1
; CHECK-M2-SYMS: kernel2
Expand Down
10 changes: 5 additions & 5 deletions llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@
; RUN: --implicit-check-not kernel0 --implicit-check-not foo \
; RUN: --implicit-check-not bar
;
; RUN: FileCheck %s -input-file=%t_2.ll --check-prefix CHECK-M2-IR \
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefix CHECK-M1-IR \
; RUN: --implicit-check-not kernel0 --implicit-check-not bar

; We expect to see 3 modules generated:
Expand Down Expand Up @@ -49,14 +49,14 @@
; should also present in a separate device image, because it is an entry point
; with unique set of used aspects.
;
; CHECK-M1-SYMS: foo
; CHECK-M1-SYMS: kernel1
;
; CHECK-M2-SYMS: kernel1
; CHECK-M2-SYMS: foo
;
; @kernel1 uses @foo and therefore @foo should be present in the same module as
; @kernel1 as well
; CHECK-M2-IR-DAG: define spir_func void @foo
; CHECK-M2-IR-DAG: define spir_kernel void @kernel1
; CHECK-M1-IR-DAG: define spir_func void @foo
; CHECK-M1-IR-DAG: define spir_kernel void @kernel1


target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
Expand Down
8 changes: 4 additions & 4 deletions llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,16 +9,16 @@

; 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
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_1.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_1.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM
; RUN: FileCheck %s -input-file=%t_esimd_0.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_large_grf_1.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM

; CHECK: [Code|Properties|Symbols]
; CHECK: {{.*}}esimd_large_grf_0.ll|{{.*}}esimd_large_grf_0.prop|{{.*}}esimd_large_grf_0.sym
; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
; CHECK: {{.*}}esimd_0.ll|{{.*}}esimd_0.prop|{{.*}}esimd_0.sym
; CHECK: {{.*}}esimd_large_grf_1.ll|{{.*}}esimd_large_grf_1.prop|{{.*}}esimd_large_grf_1.sym

; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
; CHECK-ESIMD-LargeGRF-PROP: isLargeGRF=1|1
Expand Down
8 changes: 4 additions & 4 deletions llvm/test/tools/sycl-post-link/sycl-large-grf.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,14 +9,14 @@

; 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_large_grf_1.ll --check-prefixes CHECK-LARGE-GRF-IR
; RUN: FileCheck %s -input-file=%t_large_grf_1.prop --check-prefixes CHECK-LARGE-GRF-PROP
; RUN: FileCheck %s -input-file=%t_0.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_large_grf_1.sym --check-prefixes CHECK-LARGE-GRF-SYM

; CHECK: [Code|Properties|Symbols]
; CHECK: {{.*}}_large_grf_0.ll|{{.*}}_large_grf_0.prop|{{.*}}_large_grf_0.sym
; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
; CHECK: {{.*}}_large_grf_1.ll|{{.*}}_large_grf_1.prop|{{.*}}_large_grf_1.sym

; CHECK-LARGE-GRF-PROP: isLargeGRF=1|1

Expand Down
88 changes: 26 additions & 62 deletions llvm/tools/sycl-post-link/ModuleSplitter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -261,42 +261,6 @@ EntryPointGroupVec groupEntryPointsByScope(ModuleDesc &MD,
return EntryPointGroups;
}

template <class EntryPoinGroupFunc>
EntryPointGroupVec
groupEntryPointsByAttribute(ModuleDesc &MD, StringRef AttrName,
bool EmitOnlyKernelsAsEntryPoints,
EntryPoinGroupFunc F) {
EntryPointGroupVec EntryPointGroups{};
std::map<StringRef, EntryPointSet> EntryPointMap;
Module &M = MD.getModule();

// Only process module entry points:
for (auto &F : M.functions()) {
if (!isEntryPoint(F, EmitOnlyKernelsAsEntryPoints) ||
!MD.isEntryPointCandidate(F)) {
continue;
}
if (F.hasFnAttribute(AttrName)) {
EntryPointMap[AttrName].insert(&F);
} else {
EntryPointMap[""].insert(&F);
}
}
if (!EntryPointMap.empty()) {
EntryPointGroups.reserve(EntryPointMap.size());
for (auto &EPG : EntryPointMap) {
EntryPointGroups.emplace_back(EPG.first, std::move(EPG.second),
MD.getEntryPointGroup().Props);
F(EntryPointGroups.back());
}
} else {
// No entry points met, record this.
EntryPointGroups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{});
F(EntryPointGroups.back());
}
return EntryPointGroups;
}

// Represents a call graph between functions in a module. Nodes are functions,
// edges are "calls" relation.
class CallGraph {
Expand Down Expand Up @@ -741,33 +705,16 @@ void EntryPointGroup::rebuildFromNames(const std::vector<std::string> &Names,
});
}

std::unique_ptr<ModuleSplitterBase>
getLargeGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) {
EntryPointGroupVec Groups = groupEntryPointsByAttribute(
MD, sycl::kernel_props::ATTR_LARGE_GRF, EmitOnlyKernelsAsEntryPoints,
[](EntryPointGroup &G) {
if (G.GroupId == sycl::kernel_props::ATTR_LARGE_GRF) {
G.Props.UsesLargeGRF = true;
}
});
assert(!Groups.empty() && "At least one group is expected");
assert(Groups.size() <= 2 && "At most 2 groups are expected");

if (Groups.size() > 1)
return std::make_unique<ModuleSplitter>(std::move(MD), std::move(Groups));
else
return std::make_unique<ModuleCopier>(std::move(MD), std::move(Groups));
}

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<int, 4> Aspects;
// TODO: extend this further with reqd-sub-group-size, reqd-work-group-size,
// large-grf and other properties
bool UsesLargeGRF = false;
// TODO: extend this further with reqd-sub-group-size, reqd-work-group-size
// and other properties

UsedOptionalFeatures() = default;

Expand All @@ -785,19 +732,27 @@ struct UsedOptionalFeatures {
llvm::sort(Aspects);
}

if (F->hasFnAttribute(sycl::kernel_props::ATTR_LARGE_GRF))
UsesLargeGRF = true;

llvm::hash_code AspectsHash =
llvm::hash_combine_range(Aspects.begin(), Aspects.end());
Hash = static_cast<unsigned>(llvm::hash_combine(AspectsHash));
llvm::hash_code LargeGRFHash = llvm::hash_value(UsesLargeGRF);
Hash = static_cast<unsigned>(llvm::hash_combine(AspectsHash, LargeGRFHash));
}

std::string getName(StringRef BaseName) const {
std::string generateModuleName(StringRef BaseName) const {
if (Aspects.empty())
return BaseName.str() + "-no-aspects";

std::string Ret = BaseName.str() + "-aspects";
for (int A : Aspects) {
Ret += "-" + std::to_string(A);
}

if (UsesLargeGRF)
Ret += "-large-grf";

return Ret;
}

Expand Down Expand Up @@ -833,7 +788,7 @@ struct UsedOptionalFeatures {
return false;
}

return IsEmpty == Other.IsEmpty;
return IsEmpty == Other.IsEmpty && UsesLargeGRF == Other.UsesLargeGRF;
}

unsigned hash() const { return static_cast<unsigned>(Hash); }
Expand Down Expand Up @@ -885,9 +840,18 @@ getSplitterByOptionalFeatures(ModuleDesc &&MD,
Groups.emplace_back(GLOBAL_SCOPE_NAME, EntryPointSet{});
} else {
Groups.reserve(PropertiesToFunctionsMap.size());
for (auto &EPG : PropertiesToFunctionsMap) {
Groups.emplace_back(EPG.first.getName(MD.getEntryPointGroup().GroupId),
std::move(EPG.second), MD.getEntryPointGroup().Props);
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);
}
}

Expand Down
3 changes: 0 additions & 3 deletions llvm/tools/sycl-post-link/ModuleSplitter.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -252,9 +252,6 @@ getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode,
bool AutoSplitIsGlobalScope,
bool EmitOnlyKernelsAsEntryPoints);

std::unique_ptr<ModuleSplitterBase>
getLargeGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints);

std::unique_ptr<ModuleSplitterBase>
getSplitterByOptionalFeatures(ModuleDesc &&MD,
bool EmitOnlyKernelsAsEntryPoints);
Expand Down
Loading