@@ -261,42 +261,6 @@ EntryPointGroupVec groupEntryPointsByScope(ModuleDesc &MD,
261261 return EntryPointGroups;
262262}
263263
264- template <class EntryPoinGroupFunc >
265- EntryPointGroupVec
266- groupEntryPointsByAttribute (ModuleDesc &MD, StringRef AttrName,
267- bool EmitOnlyKernelsAsEntryPoints,
268- EntryPoinGroupFunc F) {
269- EntryPointGroupVec EntryPointGroups{};
270- std::map<StringRef, EntryPointSet> EntryPointMap;
271- Module &M = MD.getModule ();
272-
273- // Only process module entry points:
274- for (auto &F : M.functions ()) {
275- if (!isEntryPoint (F, EmitOnlyKernelsAsEntryPoints) ||
276- !MD.isEntryPointCandidate (F)) {
277- continue ;
278- }
279- if (F.hasFnAttribute (AttrName)) {
280- EntryPointMap[AttrName].insert (&F);
281- } else {
282- EntryPointMap[" " insert (&F);
283- }
284- }
285- if (!EntryPointMap.empty ()) {
286- EntryPointGroups.reserve (EntryPointMap.size ());
287- for (auto &EPG : EntryPointMap) {
288- EntryPointGroups.emplace_back (EPG.first , std::move (EPG.second ),
289- MD.getEntryPointGroup ().Props );
290- F (EntryPointGroups.back ());
291- }
292- } else {
293- // No entry points met, record this.
294- EntryPointGroups.emplace_back (GLOBAL_SCOPE_NAME, EntryPointSet{});
295- F (EntryPointGroups.back ());
296- }
297- return EntryPointGroups;
298- }
299-
300264// Represents a call graph between functions in a module. Nodes are functions,
301265// edges are "calls" relation.
302266class CallGraph {
@@ -741,33 +705,16 @@ void EntryPointGroup::rebuildFromNames(const std::vector<std::string> &Names,
741705 });
742706}
743707
744- std::unique_ptr<ModuleSplitterBase>
745- getLargeGRFSplitter (ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) {
746- EntryPointGroupVec Groups = groupEntryPointsByAttribute (
747- MD, sycl::kernel_props::ATTR_LARGE_GRF, EmitOnlyKernelsAsEntryPoints,
748- [](EntryPointGroup &G) {
749- if (G.GroupId == sycl::kernel_props::ATTR_LARGE_GRF) {
750- G.Props .UsesLargeGRF = true ;
751- }
752- });
753- assert (!Groups.empty () && " At least one group is expected"
754- assert (Groups.size () <= 2 && " At most 2 groups are expected"
755-
756- if (Groups.size () > 1 )
757- return std::make_unique<ModuleSplitter>(std::move (MD), std::move (Groups));
758- else
759- return std::make_unique<ModuleCopier>(std::move (MD), std::move (Groups));
760- }
761-
762708namespace {
763709// Data structure, which represent a combination of all possible optional
764710// features used in a function.
765711//
766712// It has extra methods to be useable as a key in llvm::DenseMap.
767713struct UsedOptionalFeatures {
768714 SmallVector<int , 4 > Aspects;
769- // TODO: extend this further with reqd-sub-group-size, reqd-work-group-size,
770- // large-grf and other properties
715+ bool UsesLargeGRF = false ;
716+ // TODO: extend this further with reqd-sub-group-size, reqd-work-group-size
717+ // and other properties
771718
772719 UsedOptionalFeatures () = default ;
773720
@@ -785,19 +732,27 @@ struct UsedOptionalFeatures {
785732 llvm::sort (Aspects);
786733 }
787734
735+ if (F->hasFnAttribute (sycl::kernel_props::ATTR_LARGE_GRF))
736+ UsesLargeGRF = true ;
737+
788738 llvm::hash_code AspectsHash =
789739 llvm::hash_combine_range (Aspects.begin (), Aspects.end ());
790- Hash = static_cast <unsigned >(llvm::hash_combine (AspectsHash));
740+ llvm::hash_code LargeGRFHash = llvm::hash_value (UsesLargeGRF);
741+ Hash = static_cast <unsigned >(llvm::hash_combine (AspectsHash, LargeGRFHash));
791742 }
792743
793- std::string getName (StringRef BaseName) const {
744+ std::string generateModuleName (StringRef BaseName) const {
794745 if (Aspects.empty ())
795746 return BaseName.str () + " -no-aspects"
796747
797748 std::string Ret = BaseName.str () + " -aspects"
798749 for (int A : Aspects) {
799750 Ret += " -" std::to_string (A);
800751 }
752+
753+ if (UsesLargeGRF)
754+ Ret += " -large-grf"
755+
801756 return Ret;
802757 }
803758
@@ -833,7 +788,7 @@ struct UsedOptionalFeatures {
833788 return false ;
834789 }
835790
836- return IsEmpty == Other.IsEmpty ;
791+ return IsEmpty == Other.IsEmpty && UsesLargeGRF == Other. UsesLargeGRF ;
837792 }
838793
839794 unsigned hash () const { return static_cast <unsigned >(Hash); }
@@ -885,9 +840,18 @@ getSplitterByOptionalFeatures(ModuleDesc &&MD,
885840 Groups.emplace_back (GLOBAL_SCOPE_NAME, EntryPointSet{});
886841 } else {
887842 Groups.reserve (PropertiesToFunctionsMap.size ());
888- for (auto &EPG : PropertiesToFunctionsMap) {
889- Groups.emplace_back (EPG.first .getName (MD.getEntryPointGroup ().GroupId ),
890- std::move (EPG.second ), MD.getEntryPointGroup ().Props );
843+ for (auto &It : PropertiesToFunctionsMap) {
844+ const UsedOptionalFeatures &Features = It.first ;
845+ EntryPointSet &EntryPoints = It.second ;
846+
847+ // Start with properties of a source module
848+ EntryPointGroup::Properties MDProps = MD.getEntryPointGroup ().Props ;
849+ // Propagate LargeGRF flag to entry points group
850+ if (Features.UsesLargeGRF )
851+ MDProps.UsesLargeGRF = true ;
852+ Groups.emplace_back (
853+ Features.generateModuleName (MD.getEntryPointGroup ().GroupId ),
854+ std::move (EntryPoints), MDProps);
891855 }
892856 }
893857
0 commit comments