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 ff9aa0f29376e..3175872ee5a86 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 @@ -5,11 +5,11 @@ ; 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 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 \ @@ -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 diff --git a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll index 5fa587abca234..6b97e1d99d6c6 100644 --- a/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll +++ b/llvm/test/tools/sycl-post-link/device-code-split/per-aspect-split-3.ll @@ -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: @@ -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" 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 ccc61ca20195b..db33ae76bea0b 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,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 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 a4ae724faef15..67e74d53d88f9 100644 --- a/llvm/test/tools/sycl-post-link/sycl-large-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll @@ -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 diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 3c4daefe6b24d..c3d75c19f8678 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -261,42 +261,6 @@ EntryPointGroupVec groupEntryPointsByScope(ModuleDesc &MD, return EntryPointGroups; } -template -EntryPointGroupVec -groupEntryPointsByAttribute(ModuleDesc &MD, StringRef AttrName, - bool EmitOnlyKernelsAsEntryPoints, - EntryPoinGroupFunc F) { - EntryPointGroupVec EntryPointGroups{}; - std::map 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 { @@ -741,24 +705,6 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, }); } -std::unique_ptr -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(std::move(MD), std::move(Groups)); - else - return std::make_unique(std::move(MD), std::move(Groups)); -} - namespace { // Data structure, which represent a combination of all possible optional // features used in a function. @@ -766,8 +712,9 @@ namespace { // It has extra methods to be useable as a key in llvm::DenseMap. struct UsedOptionalFeatures { SmallVector 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; @@ -785,12 +732,16 @@ 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(llvm::hash_combine(AspectsHash)); + llvm::hash_code LargeGRFHash = llvm::hash_value(UsesLargeGRF); + Hash = static_cast(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"; @@ -798,6 +749,10 @@ struct UsedOptionalFeatures { for (int A : Aspects) { Ret += "-" + std::to_string(A); } + + if (UsesLargeGRF) + Ret += "-large-grf"; + return Ret; } @@ -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(Hash); } @@ -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); } } diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index 5c023dd21ef31..037be3f65a891 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -252,9 +252,6 @@ getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode, bool AutoSplitIsGlobalScope, bool EmitOnlyKernelsAsEntryPoints); -std::unique_ptr -getLargeGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); - std::unique_ptr getSplitterByOptionalFeatures(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 7d98af370b706..bbdb336871d59 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -806,108 +806,96 @@ processInputModule(std::unique_ptr M) { for (module_split::ModuleDesc &MDesc : TopLevelModules) { DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1); - // FIXME: large grf should be handled by properties splitter above - std::unique_ptr LargeGRFSplitter = - module_split::getLargeGRFSplitter(std::move(MDesc), - EmitOnlyKernelsAsEntryPoints); - const bool SplitByLargeGRF = LargeGRFSplitter->remainingSplits() > 1; - Modified |= SplitByLargeGRF; - - // Now split further by "large-grf" attribute. - while (LargeGRFSplitter->hasMoreSplits()) { - module_split::ModuleDesc MDesc1 = LargeGRFSplitter->nextSplit(); - DUMP_ENTRY_POINTS(MDesc1.entries(), MDesc1.Name.c_str(), 2); - MDesc1.fixupLinkageOfDirectInvokeSimdTargets(); - - // Do SYCL/ESIMD splitting. It happens always, as ESIMD and SYCL must - // undergo different set of LLVMIR passes. After this they are linked back - // together to form single module with disjoint SYCL and ESIMD call graphs - // unless -split-esimd option is specified. The graphs become disjoint - // when linked back because functions shared between graphs are cloned and - // renamed. - std::unique_ptr ESIMDSplitter = - module_split::getSplitterByKernelType(std::move(MDesc1), - EmitOnlyKernelsAsEntryPoints); - const bool SplitByESIMD = ESIMDSplitter->remainingSplits() > 1; - Modified |= SplitByESIMD; - - if (SplitByESIMD && SplitByScope && - (SplitMode == module_split::SPLIT_PER_KERNEL) && !SplitEsimd) { - // Controversial state reached - SYCL and ESIMD entry points resulting - // from SYCL/ESIMD split (which is done always) are linked back, since - // -split-esimd is not specified, but per-kernel split is requested. - warning("SYCL and ESIMD entry points detected and split mode is " - "per-kernel, so " + - SplitEsimd.ValueStr + " must also be specified"); - } - SmallVector MMs; - - while (ESIMDSplitter->hasMoreSplits()) { - module_split::ModuleDesc MDesc2 = ESIMDSplitter->nextSplit(); - DUMP_ENTRY_POINTS(MDesc2.entries(), MDesc2.Name.c_str(), 3); - Modified |= processSpecConstants(MDesc2); - - // TODO: detach compile-time properties from device globals. - if (DeviceGlobals.getNumOccurrences() > 0) { - Modified |= - runModulePass(MDesc2.getModule()); - } - if (!MDesc2.isSYCL() && LowerEsimd) { - assert(MDesc2.isESIMD() && "NYI"); - // ESIMD lowering also detects large-GRF kernels, so it must happen - // before large-GRF split. - Modified |= lowerEsimdConstructs(MDesc2); - } - MMs.emplace_back(std::move(MDesc2)); - } - if (!SplitEsimd && (MMs.size() > 1)) { - // SYCL/ESIMD splitting is not requested, link back into single module. - assert(MMs.size() == 2); - assert((MMs[0].isESIMD() && MMs[1].isSYCL()) || - (MMs[1].isESIMD() && MMs[0].isSYCL())); - int ESIMDInd = MMs[0].isESIMD() ? 0 : 1; - int SYCLInd = MMs[0].isESIMD() ? 1 : 0; - // ... but before that, make sure no link conflicts will occur. - MMs[ESIMDInd].renameDuplicatesOf(MMs[SYCLInd].getModule(), ".esimd"); - module_split::ModuleDesc M2 = - link(std::move(MMs[0]), std::move(MMs[1])); - M2.restoreLinkageOfDirectInvokeSimdTargets(); - string_vector Names; - M2.saveEntryPointNames(Names); - M2.cleanup(); // may remove some entry points, need to save/rebuild - M2.rebuildEntryPoints(Names); - MMs.clear(); - MMs.emplace_back(std::move(M2)); - DUMP_ENTRY_POINTS(MMs.back().entries(), MMs.back().Name.c_str(), 3); - Modified = true; - } + MDesc.fixupLinkageOfDirectInvokeSimdTargets(); + + // Do SYCL/ESIMD splitting. It happens always, as ESIMD and SYCL must + // undergo different set of LLVMIR passes. After this they are linked back + // together to form single module with disjoint SYCL and ESIMD call graphs + // unless -split-esimd option is specified. The graphs become disjoint + // when linked back because functions shared between graphs are cloned and + // renamed. + std::unique_ptr ESIMDSplitter = + module_split::getSplitterByKernelType(std::move(MDesc), + EmitOnlyKernelsAsEntryPoints); + const bool SplitByESIMD = ESIMDSplitter->remainingSplits() > 1; + Modified |= SplitByESIMD; + + if (SplitByESIMD && SplitByScope && + (SplitMode == module_split::SPLIT_PER_KERNEL) && !SplitEsimd) { + // Controversial state reached - SYCL and ESIMD entry points resulting + // from SYCL/ESIMD split (which is done always) are linked back, since + // -split-esimd is not specified, but per-kernel split is requested. + warning("SYCL and ESIMD entry points detected and split mode is " + "per-kernel, so " + + SplitEsimd.ValueStr + " must also be specified"); + } + SmallVector MMs; - bool SplitOccurred = SplitByScope || SplitByLargeGRF || SplitByESIMD || - SplitByOptionalFeatures; + while (ESIMDSplitter->hasMoreSplits()) { + module_split::ModuleDesc MDesc2 = ESIMDSplitter->nextSplit(); + DUMP_ENTRY_POINTS(MDesc2.entries(), MDesc2.Name.c_str(), 3); + Modified |= processSpecConstants(MDesc2); - if (IROutputOnly) { - if (SplitOccurred) { - error("some modules had to be split, '-" + IROutputOnly.ArgStr + - "' can't be used"); - } - saveModuleIR(MMs.front().getModule(), OutputFilename); - return Table; + // TODO: detach compile-time properties from device globals. + if (DeviceGlobals.getNumOccurrences() > 0) { + Modified |= + runModulePass(MDesc2.getModule()); } - // Empty IR file name directs saveModule to generate one and save IR to - // it: - std::string OutIRFileName = ""; - - if (!Modified && (OutputFilename.getNumOccurrences() == 0)) { - assert(!SplitOccurred); - OutIRFileName = InputFilename; // ... non-empty means "skip IR writing" - errs() << "sycl-post-link NOTE: no modifications to the input LLVM IR " - "have been made\n"; + if (!MDesc2.isSYCL() && LowerEsimd) { + assert(MDesc2.isESIMD() && "NYI"); + // ESIMD lowering also detects large-GRF kernels, so it must happen + // before large-GRF split. + Modified |= lowerEsimdConstructs(MDesc2); } - for (module_split::ModuleDesc &IrMD : MMs) { - IrPropSymFilenameTriple T = saveModule(IrMD, ID, OutIRFileName); - addTableRow(*Table, T); + MMs.emplace_back(std::move(MDesc2)); + } + if (!SplitEsimd && (MMs.size() > 1)) { + // SYCL/ESIMD splitting is not requested, link back into single module. + assert(MMs.size() == 2); + assert((MMs[0].isESIMD() && MMs[1].isSYCL()) || + (MMs[1].isESIMD() && MMs[0].isSYCL())); + int ESIMDInd = MMs[0].isESIMD() ? 0 : 1; + int SYCLInd = MMs[0].isESIMD() ? 1 : 0; + // ... but before that, make sure no link conflicts will occur. + MMs[ESIMDInd].renameDuplicatesOf(MMs[SYCLInd].getModule(), ".esimd"); + module_split::ModuleDesc M2 = link(std::move(MMs[0]), std::move(MMs[1])); + M2.restoreLinkageOfDirectInvokeSimdTargets(); + string_vector Names; + M2.saveEntryPointNames(Names); + M2.cleanup(); // may remove some entry points, need to save/rebuild + M2.rebuildEntryPoints(Names); + MMs.clear(); + MMs.emplace_back(std::move(M2)); + DUMP_ENTRY_POINTS(MMs.back().entries(), MMs.back().Name.c_str(), 3); + Modified = true; + } + + bool SplitOccurred = + SplitByScope || SplitByESIMD || SplitByOptionalFeatures; + + if (IROutputOnly) { + if (SplitOccurred) { + error("some modules had to be split, '-" + IROutputOnly.ArgStr + + "' can't be used"); } + saveModuleIR(MMs.front().getModule(), OutputFilename); + return Table; + } + // Empty IR file name directs saveModule to generate one and save IR to + // it: + std::string OutIRFileName = ""; + + if (!Modified && (OutputFilename.getNumOccurrences() == 0)) { + assert(!SplitOccurred); + OutIRFileName = InputFilename; // ... non-empty means "skip IR writing" + errs() << "sycl-post-link NOTE: no modifications to the input LLVM IR " + "have been made\n"; } + for (module_split::ModuleDesc &IrMD : MMs) { + IrPropSymFilenameTriple T = saveModule(IrMD, ID, OutIRFileName); + addTableRow(*Table, T); + } + ++ID; } return Table;