Skip to content

Commit ab2a42c

Browse files
sarnexsarnex
authored
[SYCL] Rename DoubleGRF to LargeGRF (#7284)
This change renames double GRF to large GRF both for users and internally in the compiler. We're doing this because we got direct feedback from customer facing engineers that we should use the large GRF terminology, and it also makes the naming consistent with other compiler work we are doing. For the user, ``` set_kernel_properties(kernel_properties::use_double_grf); ``` still works, it will just throw a deprecated warning and will be removed in a future release. The new way is ``` set_kernel_properties(kernel_properties::use_large_grf); ``` There should be no ABI break because we still check the previous image property name in the program manager, so applications built with an old compiler work using the runtime from a new compiler. I confirmed this with manual testing. I will update the system test here to test the new flag as well: https://github.com/intel/llvm-test-suite/blob/intel/SYCL/DeviceCodeSplit/double-grf.cpp Signed-off-by: Sarnie, Nick <[email protected]>
1 parent 823f2b2 commit ab2a42c

File tree

10 files changed

+93
-89
lines changed

10 files changed

+93
-89
lines changed

‎llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@
1414

1515
namespace sycl {
1616
namespace kernel_props {
17-
constexpr char ATTR_DOUBLE_GRF[] = "double-grf";
17+
constexpr char ATTR_LARGE_GRF[] = "large-grf";
1818
}
1919
} // namespace sycl
2020
namespace llvm {

‎llvm/lib/SYCLLowerIR/LowerKernelProps.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ constexpr char SET_KERNEL_PROPS_FUNC_NAME[] =
2929

3030
// Kernel property identifiers. Should match ones in
3131
// sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp
32-
enum property_ids { use_double_grf = 0 };
32+
enum property_ids { use_large_grf = 0 };
3333

3434
void processSetKernelPropertiesCall(CallInst &CI) {
3535
auto F = CI.getFunction();
@@ -43,11 +43,11 @@ void processSetKernelPropertiesCall(CallInst &CI) {
4343
uint64_t PropID = cast<llvm::ConstantInt>(ArgV)->getZExtValue();
4444

4545
switch (PropID) {
46-
case property_ids::use_double_grf:
46+
case property_ids::use_large_grf:
4747
// TODO: Keep track of traversed functions to avoid repeating traversals
4848
// over same function.
4949
llvm::sycl::utils::traverseCallgraphUp(F, [](Function *GraphNode) {
50-
GraphNode->addFnAttr(::sycl::kernel_props::ATTR_DOUBLE_GRF);
50+
GraphNode->addFnAttr(::sycl::kernel_props::ATTR_LARGE_GRF);
5151
});
5252
break;
5353
default:

‎llvm/test/SYCLLowerIR/lower_kernel_props.ll

Lines changed: 11 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -3,18 +3,18 @@
33
; intrinsic by LowerKernelProps pass - it should:
44
; - determine kernels calling this intrinsic (walk up the call graph)
55
; - remove the intrinsic call
6-
; - mark the kernel with corresponding attribute (only "double-grf" for now)
6+
; - mark the kernel with corresponding attribute (only "large-grf" for now)
77

88
; RUN: opt -passes=lower-kernel-props -S %s -o - | FileCheck %s
99

10-
; ModuleID = 'double_grf.bc'
10+
; ModuleID = 'large_grf.bc'
1111
source_filename = "llvm-link"
1212
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"
1313
target triple = "spir64-unknown-unknown"
1414

15-
define dso_local spir_func void @_Z17double_grf_markerv() {
16-
; CHECK: define dso_local spir_func void @_Z17double_grf_markerv()
17-
; -- '0' constant argument means "double GRF" property:
15+
define dso_local spir_func void @_Z17large_grf_markerv() {
16+
; CHECK: define dso_local spir_func void @_Z17large_grf_markerv()
17+
; -- '0' constant argument means "large GRF" property:
1818
call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
1919
; -- Check that LowerKernelProps removed the marker call above:
2020
; CHECK-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi
@@ -25,20 +25,20 @@ define dso_local spir_func void @_Z17double_grf_markerv() {
2525
declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef)
2626

2727
; -- This kernel calls the marker function indirectly
28-
define weak_odr dso_local spir_kernel void @__double_grf_kernel1() !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
29-
; CHECK: {{.*}} spir_kernel void @__double_grf_kernel1() #0
30-
call spir_func void @_Z17double_grf_markerv()
28+
define weak_odr dso_local spir_kernel void @__large_grf_kernel1() !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
29+
; CHECK: {{.*}} spir_kernel void @__large_grf_kernel1() #0
30+
call spir_func void @_Z17large_grf_markerv()
3131
ret void
3232
}
3333

3434
; -- This kernel calls the marker function directly
35-
define weak_odr dso_local spir_kernel void @__double_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
36-
; CHECK: {{.*}} spir_kernel void @__double_grf_kernel2() #0
35+
define weak_odr dso_local spir_kernel void @__large_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
36+
; CHECK: {{.*}} spir_kernel void @__large_grf_kernel2() #0
3737
call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
3838
ret void
3939
}
4040

41-
attributes #0 = { "double-grf" }
41+
attributes #0 = { "large-grf" }
4242

4343
!0 = !{}
4444
!1 = !{i32 1}

‎llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll renamed to ‎llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll

Lines changed: 17 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1,38 +1,38 @@
11
; This test checks handling of the
2-
; set_kernel_properties(kernel_properties::use_double_grf);
2+
; set_kernel_properties(kernel_properties::use_large_grf);
33
; by the post-link-tool:
44
; - ESIMD/SYCL splitting happens as usual
55
; - ESIMD module is further split into callgraphs for entry points requesting
6-
; "double GRF" and callgraphs for entry points which are not
7-
; - Compiler adds 'isDoubleGRF' property to the ESIMD device binary
8-
; images requesting "double GRF"
6+
; "large GRF" and callgraphs for entry points which are not
7+
; - Compiler adds 'isLargeGRF' property to the ESIMD device binary
8+
; images requesting "large GRF"
99

1010
; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S %s -o %t.table
1111
; RUN: FileCheck %s -input-file=%t.table
12-
; RUN: FileCheck %s -input-file=%t_esimd_x2grf_0.ll --check-prefixes CHECK-ESIMD-2xGRF-IR
13-
; RUN: FileCheck %s -input-file=%t_esimd_x2grf_0.prop --check-prefixes CHECK-ESIMD-2xGRF-PROP
12+
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR
13+
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
1414
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM
1515
; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-SYM
16-
; RUN: FileCheck %s -input-file=%t_esimd_x2grf_0.sym --check-prefixes CHECK-ESIMD-2xGRF-SYM
16+
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM
1717

1818
; CHECK: [Code|Properties|Symbols]
19-
; CHECK: {{.*}}esimd_x2grf_0.ll|{{.*}}esimd_x2grf_0.prop|{{.*}}esimd_x2grf_0.sym
19+
; CHECK: {{.*}}esimd_large_grf_0.ll|{{.*}}esimd_large_grf_0.prop|{{.*}}esimd_large_grf_0.sym
2020
; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
2121
; CHECK: {{.*}}esimd_0.ll|{{.*}}esimd_0.prop|{{.*}}esimd_0.sym
2222

23-
; CHECK-ESIMD-2xGRF-PROP: isEsimdImage=1|1
24-
; CHECK-ESIMD-2xGRF-PROP: isDoubleGRF=1|1
23+
; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
24+
; CHECK-ESIMD-LargeGRF-PROP: isLargeGRF=1|1
2525

2626
; CHECK-SYCL-SYM: __SYCL_kernel
2727
; CHECK-SYCL-SYM-EMPTY:
2828

2929
; CHECK-ESIMD-SYM: __ESIMD_kernel
3030
; CHECK-ESIMD-SYM-EMPTY:
3131

32-
; CHECK-ESIMD-2xGRF-SYM: __ESIMD_double_grf_kernel
33-
; CHECK-ESIMD-2xGRF-SYM-EMPTY:
32+
; CHECK-ESIMD-LargeGRF-SYM: __ESIMD_large_grf_kernel
33+
; CHECK-ESIMD-LargeGRF-SYM-EMPTY:
3434

35-
; ModuleID = 'double_grf.bc'
35+
; ModuleID = 'large_grf.bc'
3636
source_filename = "llvm-link"
3737
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"
3838
target triple = "spir64-unknown-unknown"
@@ -47,19 +47,19 @@ entry:
4747
ret void
4848
}
4949

50-
define dso_local spir_func void @_Z17double_grf_markerv() {
50+
define dso_local spir_func void @_Z17large_grf_markerv() {
5151
entry:
5252
call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
5353
; -- Check that ESIMD lowering removed the marker call above:
54-
; CHECK-ESIMD-2xGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi
54+
; CHECK-ESIMD-LargeGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi
5555
ret void
5656
}
5757

5858
declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef)
5959

60-
define weak_odr dso_local spir_kernel void @__ESIMD_double_grf_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
60+
define weak_odr dso_local spir_kernel void @__ESIMD_large_grf_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
6161
entry:
62-
call spir_func void @_Z17double_grf_markerv()
62+
call spir_func void @_Z17large_grf_markerv()
6363
ret void
6464
}
6565

‎llvm/test/tools/sycl-post-link/sycl-double-grf.ll renamed to ‎llvm/test/tools/sycl-post-link/sycl-large-grf.ll

Lines changed: 16 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -1,32 +1,32 @@
11
; This test checks handling of the
2-
; set_kernel_properties(kernel_properties::use_double_grf);
2+
; set_kernel_properties(kernel_properties::use_large_grf);
33
; by the post-link-tool:
44
; - ESIMD/SYCL splitting happens as usual
55
; - ESIMD module is further split into callgraphs for entry points requesting
6-
; "double GRF" and callgraphs for entry points which are not
7-
; - Compiler adds 'isDoubleGRF' property to the device binary
8-
; images requesting "double GRF"
6+
; "large GRF" and callgraphs for entry points which are not
7+
; - Compiler adds 'isLargeGRF' property to the device binary
8+
; images requesting "large GRF"
99

1010
; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S %s -o %t.table
1111
; RUN: FileCheck %s -input-file=%t.table
12-
; RUN: FileCheck %s -input-file=%t_x2grf_0.ll --check-prefixes CHECK-2xGRF-IR
13-
; RUN: FileCheck %s -input-file=%t_x2grf_0.prop --check-prefixes CHECK-2xGRF-PROP
12+
; RUN: FileCheck %s -input-file=%t_large_grf_0.ll --check-prefixes CHECK-LARGE-GRF-IR
13+
; RUN: FileCheck %s -input-file=%t_large_grf_0.prop --check-prefixes CHECK-LARGE-GRF-PROP
1414
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM
15-
; RUN: FileCheck %s -input-file=%t_x2grf_0.sym --check-prefixes CHECK-2xGRF-SYM
15+
; RUN: FileCheck %s -input-file=%t_large_grf_0.sym --check-prefixes CHECK-LARGE-GRF-SYM
1616

1717
; CHECK: [Code|Properties|Symbols]
18-
; CHECK: {{.*}}_x2grf_0.ll|{{.*}}_x2grf_0.prop|{{.*}}_x2grf_0.sym
18+
; CHECK: {{.*}}_large_grf_0.ll|{{.*}}_large_grf_0.prop|{{.*}}_large_grf_0.sym
1919
; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
2020

21-
; CHECK-2xGRF-PROP: isDoubleGRF=1|1
21+
; CHECK-LARGE-GRF-PROP: isLargeGRF=1|1
2222

2323
; CHECK-SYCL-SYM: __SYCL_kernel
2424
; CHECK-SYCL-SYM-EMPTY:
2525

26-
; CHECK-2xGRF-SYM: __double_grf_kernel
27-
; CHECK-2xGRF-SYM-EMPTY:
26+
; CHECK-LARGE-GRF-SYM: __large_grf_kernel
27+
; CHECK-LARGE-GRF-SYM-EMPTY:
2828

29-
; ModuleID = 'double_grf.bc'
29+
; ModuleID = 'large_grf.bc'
3030
source_filename = "llvm-link"
3131
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"
3232
target triple = "spir64-unknown-unknown"
@@ -36,19 +36,19 @@ entry:
3636
ret void
3737
}
3838

39-
define dso_local spir_func void @_Z17double_grf_markerv() {
39+
define dso_local spir_func void @_Z17large_grf_markerv() {
4040
entry:
4141
call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
4242
; -- Check that LowerKernelProps lowering removed the marker call above:
43-
; CHECK-2xGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi
43+
; CHECK-LARGE-GRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi
4444
ret void
4545
}
4646

4747
declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef)
4848

49-
define weak_odr dso_local spir_kernel void @__double_grf_kernel() #0 {
49+
define weak_odr dso_local spir_kernel void @__large_grf_kernel() #0 {
5050
entry:
51-
call spir_func void @_Z17double_grf_markerv()
51+
call spir_func void @_Z17large_grf_markerv()
5252
ret void
5353
}
5454

‎llvm/tools/sycl-post-link/ModuleSplitter.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -711,8 +711,8 @@ void ModuleDesc::dump() const {
711711
llvm::errs() << "split_module::ModuleDesc[" << Name << "] {\n";
712712
llvm::errs() << " ESIMD:" << toString(EntryPoints.Props.HasESIMD)
713713
<< ", SpecConstMet:" << (Props.SpecConstsMet ? "YES" : "NO")
714-
<< ", DoubleGRF:"
715-
<< (EntryPoints.Props.UsesDoubleGRF ? "YES" : "NO") << "\n";
714+
<< ", LargeGRF:"
715+
<< (EntryPoints.Props.UsesLargeGRF ? "YES" : "NO") << "\n";
716716
dumpEntryPoints(entries(), EntryPoints.GroupId.str().c_str(), 1);
717717
llvm::errs() << "}\n";
718718
}
@@ -744,12 +744,12 @@ void EntryPointGroup::rebuildFromNames(const std::vector<std::string> &Names,
744744
}
745745

746746
std::unique_ptr<ModuleSplitterBase>
747-
getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) {
747+
getLargeGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) {
748748
EntryPointGroupVec Groups = groupEntryPointsByAttribute(
749-
MD, sycl::kernel_props::ATTR_DOUBLE_GRF, EmitOnlyKernelsAsEntryPoints,
749+
MD, sycl::kernel_props::ATTR_LARGE_GRF, EmitOnlyKernelsAsEntryPoints,
750750
[](EntryPointGroup &G) {
751-
if (G.GroupId == sycl::kernel_props::ATTR_DOUBLE_GRF) {
752-
G.Props.UsesDoubleGRF = true;
751+
if (G.GroupId == sycl::kernel_props::ATTR_LARGE_GRF) {
752+
G.Props.UsesLargeGRF = true;
753753
}
754754
});
755755
assert(!Groups.empty() && "At least one group is expected");

‎llvm/tools/sycl-post-link/ModuleSplitter.h

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -55,8 +55,8 @@ struct EntryPointGroup {
5555
struct Properties {
5656
// Whether all EPs are ESIMD, SYCL or there are both kinds.
5757
SyclEsimdSplitStatus HasESIMD = SyclEsimdSplitStatus::SYCL_AND_ESIMD;
58-
// Whether any of the EPs use double GRF mode.
59-
bool UsesDoubleGRF = false;
58+
// Whether any of the EPs use large GRF mode.
59+
bool UsesLargeGRF = false;
6060
// Scope represented by EPs in a group
6161
EntryPointsGroupScope Scope = Scope_Global;
6262

@@ -65,7 +65,7 @@ struct EntryPointGroup {
6565
Res.HasESIMD = HasESIMD == Other.HasESIMD
6666
? HasESIMD
6767
: SyclEsimdSplitStatus::SYCL_AND_ESIMD;
68-
Res.UsesDoubleGRF = UsesDoubleGRF || Other.UsesDoubleGRF;
68+
Res.UsesLargeGRF = UsesLargeGRF || Other.UsesLargeGRF;
6969
// Scope remains global
7070
return Res;
7171
}
@@ -90,8 +90,8 @@ struct EntryPointGroup {
9090
bool isSycl() const {
9191
return Props.HasESIMD == SyclEsimdSplitStatus::SYCL_ONLY;
9292
}
93-
// Tells if some entry points use double GRF mode.
94-
bool isDoubleGRF() const { return Props.UsesDoubleGRF; }
93+
// Tells if some entry points use large GRF mode.
94+
bool isLargeGRF() const { return Props.UsesLargeGRF; }
9595

9696
void saveNames(std::vector<std::string> &Dest) const;
9797
void rebuildFromNames(const std::vector<std::string> &Names, const Module &M);
@@ -146,7 +146,7 @@ class ModuleDesc {
146146

147147
bool isESIMD() const { return EntryPoints.isEsimd(); }
148148
bool isSYCL() const { return EntryPoints.isSycl(); }
149-
bool isDoubleGRF() const { return EntryPoints.isDoubleGRF(); }
149+
bool isLargeGRF() const { return EntryPoints.isLargeGRF(); }
150150

151151
const EntryPointSet &entries() const { return EntryPoints.Functions; }
152152
const EntryPointGroup &getEntryPointGroup() const { return EntryPoints; }
@@ -251,7 +251,7 @@ getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode,
251251
bool EmitOnlyKernelsAsEntryPoints);
252252

253253
std::unique_ptr<ModuleSplitterBase>
254-
getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints);
254+
getLargeGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints);
255255

256256
#ifndef NDEBUG
257257
void dumpEntryPoints(const EntryPointSet &C, const char *msg = "", int Tab = 0);

‎llvm/tools/sycl-post-link/sycl-post-link.cpp

Lines changed: 16 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -442,8 +442,8 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
442442
if (MD.isESIMD()) {
443443
PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isEsimdImage", true});
444444
}
445-
if (MD.isDoubleGRF())
446-
PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isDoubleGRF", true});
445+
if (MD.isLargeGRF())
446+
PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true});
447447
{
448448
std::vector<StringRef> FuncNames = getKernelNamesUsingAssert(M);
449449
for (const StringRef &FName : FuncNames)
@@ -560,8 +560,8 @@ bool lowerEsimdConstructs(module_split::ModuleDesc &MD) {
560560

561561
// Compute the filename suffix for the module
562562
StringRef getModuleSuffix(const module_split::ModuleDesc &MD) {
563-
if (MD.isDoubleGRF()) {
564-
return MD.isESIMD() ? "_esimd_x2grf" : "_x2grf";
563+
if (MD.isLargeGRF()) {
564+
return MD.isESIMD() ? "_esimd_large_grf" : "_large_grf";
565565
}
566566
return MD.isESIMD() ? "_esimd" : "";
567567
}
@@ -735,7 +735,7 @@ processInputModule(std::unique_ptr<Module> M) {
735735
}
736736
Modified |= InvokeSimdMet;
737737

738-
// Lower kernel properties setting APIs before "double GRF" splitting, as:
738+
// Lower kernel properties setting APIs before "large GRF" splitting, as:
739739
// - the latter uses the result of the former
740740
// - saves processing time
741741
Modified |= runModulePass<SYCLLowerKernelPropsPass>(*M);
@@ -774,15 +774,15 @@ processInputModule(std::unique_ptr<Module> M) {
774774
module_split::ModuleDesc MDesc = ScopedSplitter->nextSplit();
775775
DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1);
776776

777-
std::unique_ptr<module_split::ModuleSplitterBase> DoubleGRFSplitter =
778-
module_split::getDoubleGRFSplitter(std::move(MDesc),
779-
EmitOnlyKernelsAsEntryPoints);
780-
const bool SplitByDoubleGRF = DoubleGRFSplitter->totalSplits() > 1;
781-
Modified |= SplitByDoubleGRF;
777+
std::unique_ptr<module_split::ModuleSplitterBase> LargeGRFSplitter =
778+
module_split::getLargeGRFSplitter(std::move(MDesc),
779+
EmitOnlyKernelsAsEntryPoints);
780+
const bool SplitByLargeGRF = LargeGRFSplitter->totalSplits() > 1;
781+
Modified |= SplitByLargeGRF;
782782

783-
// Now split further by "esimd-double-grf" attribute.
784-
while (DoubleGRFSplitter->hasMoreSplits()) {
785-
module_split::ModuleDesc MDesc1 = DoubleGRFSplitter->nextSplit();
783+
// Now split further by "large-grf" attribute.
784+
while (LargeGRFSplitter->hasMoreSplits()) {
785+
module_split::ModuleDesc MDesc1 = LargeGRFSplitter->nextSplit();
786786
DUMP_ENTRY_POINTS(MDesc1.entries(), MDesc1.Name.c_str(), 2);
787787
MDesc1.fixupLinkageOfDirectInvokeSimdTargets();
788788

@@ -821,8 +821,8 @@ processInputModule(std::unique_ptr<Module> M) {
821821
}
822822
if (!MDesc2.isSYCL() && LowerEsimd) {
823823
assert(MDesc2.isESIMD() && "NYI");
824-
// ESIMD lowering also detects double-GRF kernels, so it must happen
825-
// before double-GRF split.
824+
// ESIMD lowering also detects large-GRF kernels, so it must happen
825+
// before large-GRF split.
826826
Modified |= lowerEsimdConstructs(MDesc2);
827827
}
828828
MMs.emplace_back(std::move(MDesc2));
@@ -848,7 +848,7 @@ processInputModule(std::unique_ptr<Module> M) {
848848
DUMP_ENTRY_POINTS(MMs.back().entries(), MMs.back().Name.c_str(), 3);
849849
Modified = true;
850850
}
851-
bool SplitOccurred = SplitByScope || SplitByDoubleGRF || SplitByESIMD;
851+
bool SplitOccurred = SplitByScope || SplitByLargeGRF || SplitByESIMD;
852852

853853
if (IROutputOnly) {
854854
if (SplitOccurred) {

0 commit comments

Comments
 (0)