Skip to content

Commit 102df2a

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents b867fe5 + 1a0d7f1 commit 102df2a

File tree

32 files changed

+245
-316
lines changed

32 files changed

+245
-316
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2087,6 +2087,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler {
20872087

20882088
const ConstantArrayType *CAT =
20892089
SemaRef.getASTContext().getAsConstantArrayType(ArrayTy);
2090+
assert(CAT && "Should only be called on constant-size array.");
20902091
QualType ModifiedArray = SemaRef.getASTContext().getConstantArrayType(
20912092
ModifiedArrayElement, CAT->getSize(),
20922093
const_cast<Expr *>(CAT->getSizeExpr()), CAT->getSizeModifier(),

‎clang/test/Frontend/sycl-propagate-aspect-warning.cpp

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -5,50 +5,50 @@
55

66
#include "sycl.hpp"
77

8-
// expected-note@+1 4 {{propagated from call to function 'Struct1::Struct1()'}}
8+
// expected-note-re@+1 4 {{propagated from call to function '{{.*}}Struct1::Struct1({{.*}})'}}
99
struct [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16)]] Struct1 {
1010
int a = 0;
1111
};
1212

13-
// expected-note@+1 4 {{propagated from call to function 'func5(int)'}}
13+
// expected-note-re@+1 4 {{propagated from call to function '{{.*}}func5(int)'}}
1414
[[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] int func5(int a) {
1515
return 0;
1616
}
1717

18-
// expected-note@+1 4 {{propagated from call to function 'func3(int, int, int)'}}
18+
// expected-note-re@+1 4 {{propagated from call to function '{{.*}}func3(int, int, int)'}}
1919
int func3(int a, int b, int c) {
2020
Struct1 s;
2121
s.a = 1;
2222
return s.a;
2323
}
2424

25-
// expected-note@+1 4 {{propagated from call to function 'func4(int, int)'}}
25+
// expected-note-re@+1 4 {{propagated from call to function '{{.*}}func4(int, int)'}}
2626
int func4(int a, int b) {
2727
double x = 3.0;
2828
return func3(a, b, (int)x);
2929
}
3030

31-
// expected-note@+1 2 {{propagated from call to function 'func2(int, int)'}}
31+
// expected-note-re@+1 2 {{propagated from call to function '{{.*}}func2(int, int)'}}
3232
int func2(int a, int b) { return func3(a, b, 1); }
3333

34-
// expected-warning@+1 {{function 'func1(int)' uses aspect 'fp16' not listed in its 'sycl::device_has' attribute}}
34+
// expected-warning-re@+1 {{function '{{.*}}func1(int)' uses aspect 'fp16' not listed in its 'sycl::device_has' attribute}}
3535
[[sycl::device_has(sycl::aspect::fp64)]] int func1(int a) { return func2(a, 1); }
3636

37-
// expected-note@+1 4 {{propagated from call to function 'func6(int, int, int)'}}
37+
// expected-note-re@+1 4 {{propagated from call to function '{{.*}}func6(int, int, int)'}}
3838
int func6(int a, int b, int c) {
3939
return func5(a);
4040
}
4141

42-
// expected-note@+1 4 {{propagated from call to function 'func7(int, int)'}}
42+
// expected-note-re@+1 4 {{propagated from call to function '{{.*}}func7(int, int)'}}
4343
int func7(int a, int b) {
4444
double x = 3.0;
4545
return func6(a, b, (int)x);
4646
}
4747

48-
// expected-note@+1 2 {{propagated from call to function 'func8(int, int)'}}
48+
// expected-note-re@+1 2 {{propagated from call to function '{{.*}}func8(int, int)'}}
4949
int func8(int a, int b) { return func6(a, b, 1); }
5050

51-
// expected-warning@+1 {{function 'func9(int)' uses aspect 'cpu' not listed in its 'sycl::device_has' attribute}}
51+
// expected-warning-re@+1 {{function '{{.*}}func9(int)' uses aspect 'cpu' not listed in its 'sycl::device_has' attribute}}
5252
[[sycl::device_has(sycl::aspect::fp64)]] int func9(int a) { return func8(a, 1); }
5353

5454
int main() {

‎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);

0 commit comments

Comments
 (0)