From ab2a42c64c10249c40de9916a5936cb9d82592a4 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Mon, 7 Nov 2022 23:30:25 -0500 Subject: [PATCH] [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 --- .../llvm/SYCLLowerIR/LowerKernelProps.h | 2 +- llvm/lib/SYCLLowerIR/LowerKernelProps.cpp | 6 ++-- llvm/test/SYCLLowerIR/lower_kernel_props.ll | 22 ++++++------ ...-double-grf.ll => sycl-esimd-large-grf.ll} | 34 +++++++++---------- .../{sycl-double-grf.ll => sycl-large-grf.ll} | 32 ++++++++--------- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 12 +++---- llvm/tools/sycl-post-link/ModuleSplitter.h | 14 ++++---- llvm/tools/sycl-post-link/sycl-post-link.cpp | 32 ++++++++--------- .../intel/experimental/kernel_properties.hpp | 22 ++++++------ .../program_manager/program_manager.cpp | 6 ++-- 10 files changed, 93 insertions(+), 89 deletions(-) rename llvm/test/tools/sycl-post-link/{sycl-esimd-double-grf.ll => sycl-esimd-large-grf.ll} (56%) rename llvm/test/tools/sycl-post-link/{sycl-double-grf.ll => sycl-large-grf.ll} (54%) diff --git a/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h b/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h index 81e12f9c93d9e..72d405d8e61a3 100644 --- a/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h +++ b/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h @@ -14,7 +14,7 @@ namespace sycl { namespace kernel_props { -constexpr char ATTR_DOUBLE_GRF[] = "double-grf"; +constexpr char ATTR_LARGE_GRF[] = "large-grf"; } } // namespace sycl namespace llvm { diff --git a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp index 2eb227b29e34c..8915b7819e05d 100644 --- a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp +++ b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp @@ -29,7 +29,7 @@ constexpr char SET_KERNEL_PROPS_FUNC_NAME[] = // Kernel property identifiers. Should match ones in // sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp -enum property_ids { use_double_grf = 0 }; +enum property_ids { use_large_grf = 0 }; void processSetKernelPropertiesCall(CallInst &CI) { auto F = CI.getFunction(); @@ -43,11 +43,11 @@ void processSetKernelPropertiesCall(CallInst &CI) { uint64_t PropID = cast(ArgV)->getZExtValue(); switch (PropID) { - case property_ids::use_double_grf: + case property_ids::use_large_grf: // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. llvm::sycl::utils::traverseCallgraphUp(F, [](Function *GraphNode) { - GraphNode->addFnAttr(::sycl::kernel_props::ATTR_DOUBLE_GRF); + GraphNode->addFnAttr(::sycl::kernel_props::ATTR_LARGE_GRF); }); break; default: diff --git a/llvm/test/SYCLLowerIR/lower_kernel_props.ll b/llvm/test/SYCLLowerIR/lower_kernel_props.ll index aa36cd8c94c06..4681437db1013 100644 --- a/llvm/test/SYCLLowerIR/lower_kernel_props.ll +++ b/llvm/test/SYCLLowerIR/lower_kernel_props.ll @@ -3,18 +3,18 @@ ; intrinsic by LowerKernelProps pass - it should: ; - determine kernels calling this intrinsic (walk up the call graph) ; - remove the intrinsic call -; - mark the kernel with corresponding attribute (only "double-grf" for now) +; - mark the kernel with corresponding attribute (only "large-grf" for now) ; RUN: opt -passes=lower-kernel-props -S %s -o - | FileCheck %s -; ModuleID = 'double_grf.bc' +; ModuleID = 'large_grf.bc' source_filename = "llvm-link" 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" target triple = "spir64-unknown-unknown" -define dso_local spir_func void @_Z17double_grf_markerv() { -; CHECK: define dso_local spir_func void @_Z17double_grf_markerv() -; -- '0' constant argument means "double GRF" property: +define dso_local spir_func void @_Z17large_grf_markerv() { +; CHECK: define dso_local spir_func void @_Z17large_grf_markerv() +; -- '0' constant argument means "large GRF" property: call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0) ; -- Check that LowerKernelProps removed the marker call above: ; CHECK-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi @@ -25,20 +25,20 @@ define dso_local spir_func void @_Z17double_grf_markerv() { declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef) ; -- This kernel calls the marker function indirectly -define weak_odr dso_local spir_kernel void @__double_grf_kernel1() !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { -; CHECK: {{.*}} spir_kernel void @__double_grf_kernel1() #0 - call spir_func void @_Z17double_grf_markerv() +define weak_odr dso_local spir_kernel void @__large_grf_kernel1() !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { +; CHECK: {{.*}} spir_kernel void @__large_grf_kernel1() #0 + call spir_func void @_Z17large_grf_markerv() ret void } ; -- This kernel calls the marker function directly -define weak_odr dso_local spir_kernel void @__double_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { -; CHECK: {{.*}} spir_kernel void @__double_grf_kernel2() #0 +define weak_odr dso_local spir_kernel void @__large_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { +; CHECK: {{.*}} spir_kernel void @__large_grf_kernel2() #0 call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0) ret void } -attributes #0 = { "double-grf" } +attributes #0 = { "large-grf" } !0 = !{} !1 = !{i32 1} diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll b/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll similarity index 56% rename from llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll rename to llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll index 5d6104a40505f..ccc61ca20195b 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll @@ -1,27 +1,27 @@ ; This test checks handling of the -; set_kernel_properties(kernel_properties::use_double_grf); +; set_kernel_properties(kernel_properties::use_large_grf); ; by the post-link-tool: ; - ESIMD/SYCL splitting happens as usual ; - ESIMD module is further split into callgraphs for entry points requesting -; "double GRF" and callgraphs for entry points which are not -; - Compiler adds 'isDoubleGRF' property to the ESIMD device binary -; images requesting "double GRF" +; "large GRF" and callgraphs for entry points which are not +; - Compiler adds 'isLargeGRF' property to the ESIMD device binary +; images requesting "large GRF" ; 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_x2grf_0.ll --check-prefixes CHECK-ESIMD-2xGRF-IR -; RUN: FileCheck %s -input-file=%t_esimd_x2grf_0.prop --check-prefixes CHECK-ESIMD-2xGRF-PROP +; 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_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_x2grf_0.sym --check-prefixes CHECK-ESIMD-2xGRF-SYM +; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM ; CHECK: [Code|Properties|Symbols] -; CHECK: {{.*}}esimd_x2grf_0.ll|{{.*}}esimd_x2grf_0.prop|{{.*}}esimd_x2grf_0.sym +; 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-2xGRF-PROP: isEsimdImage=1|1 -; CHECK-ESIMD-2xGRF-PROP: isDoubleGRF=1|1 +; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1 +; CHECK-ESIMD-LargeGRF-PROP: isLargeGRF=1|1 ; CHECK-SYCL-SYM: __SYCL_kernel ; CHECK-SYCL-SYM-EMPTY: @@ -29,10 +29,10 @@ ; CHECK-ESIMD-SYM: __ESIMD_kernel ; CHECK-ESIMD-SYM-EMPTY: -; CHECK-ESIMD-2xGRF-SYM: __ESIMD_double_grf_kernel -; CHECK-ESIMD-2xGRF-SYM-EMPTY: +; CHECK-ESIMD-LargeGRF-SYM: __ESIMD_large_grf_kernel +; CHECK-ESIMD-LargeGRF-SYM-EMPTY: -; ModuleID = 'double_grf.bc' +; ModuleID = 'large_grf.bc' source_filename = "llvm-link" 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" target triple = "spir64-unknown-unknown" @@ -47,19 +47,19 @@ entry: ret void } -define dso_local spir_func void @_Z17double_grf_markerv() { +define dso_local spir_func void @_Z17large_grf_markerv() { entry: call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0) ; -- Check that ESIMD lowering removed the marker call above: -; CHECK-ESIMD-2xGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi +; CHECK-ESIMD-LargeGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi ret void } declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef) -define weak_odr dso_local spir_kernel void @__ESIMD_double_grf_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { +define weak_odr dso_local spir_kernel void @__ESIMD_large_grf_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { entry: - call spir_func void @_Z17double_grf_markerv() + call spir_func void @_Z17large_grf_markerv() ret void } diff --git a/llvm/test/tools/sycl-post-link/sycl-double-grf.ll b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll similarity index 54% rename from llvm/test/tools/sycl-post-link/sycl-double-grf.ll rename to llvm/test/tools/sycl-post-link/sycl-large-grf.ll index dcc5695ce48e9..a4ae724faef15 100644 --- a/llvm/test/tools/sycl-post-link/sycl-double-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-large-grf.ll @@ -1,32 +1,32 @@ ; This test checks handling of the -; set_kernel_properties(kernel_properties::use_double_grf); +; set_kernel_properties(kernel_properties::use_large_grf); ; by the post-link-tool: ; - ESIMD/SYCL splitting happens as usual ; - ESIMD module is further split into callgraphs for entry points requesting -; "double GRF" and callgraphs for entry points which are not -; - Compiler adds 'isDoubleGRF' property to the device binary -; images requesting "double GRF" +; "large GRF" and callgraphs for entry points which are not +; - Compiler adds 'isLargeGRF' property to the device binary +; images requesting "large GRF" ; 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_x2grf_0.ll --check-prefixes CHECK-2xGRF-IR -; RUN: FileCheck %s -input-file=%t_x2grf_0.prop --check-prefixes CHECK-2xGRF-PROP +; 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_0.sym --check-prefixes CHECK-SYCL-SYM -; RUN: FileCheck %s -input-file=%t_x2grf_0.sym --check-prefixes CHECK-2xGRF-SYM +; RUN: FileCheck %s -input-file=%t_large_grf_0.sym --check-prefixes CHECK-LARGE-GRF-SYM ; CHECK: [Code|Properties|Symbols] -; CHECK: {{.*}}_x2grf_0.ll|{{.*}}_x2grf_0.prop|{{.*}}_x2grf_0.sym +; CHECK: {{.*}}_large_grf_0.ll|{{.*}}_large_grf_0.prop|{{.*}}_large_grf_0.sym ; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym -; CHECK-2xGRF-PROP: isDoubleGRF=1|1 +; CHECK-LARGE-GRF-PROP: isLargeGRF=1|1 ; CHECK-SYCL-SYM: __SYCL_kernel ; CHECK-SYCL-SYM-EMPTY: -; CHECK-2xGRF-SYM: __double_grf_kernel -; CHECK-2xGRF-SYM-EMPTY: +; CHECK-LARGE-GRF-SYM: __large_grf_kernel +; CHECK-LARGE-GRF-SYM-EMPTY: -; ModuleID = 'double_grf.bc' +; ModuleID = 'large_grf.bc' source_filename = "llvm-link" 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" target triple = "spir64-unknown-unknown" @@ -36,19 +36,19 @@ entry: ret void } -define dso_local spir_func void @_Z17double_grf_markerv() { +define dso_local spir_func void @_Z17large_grf_markerv() { entry: call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0) ; -- Check that LowerKernelProps lowering removed the marker call above: -; CHECK-2xGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi +; CHECK-LARGE-GRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi ret void } declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef) -define weak_odr dso_local spir_kernel void @__double_grf_kernel() #0 { +define weak_odr dso_local spir_kernel void @__large_grf_kernel() #0 { entry: - call spir_func void @_Z17double_grf_markerv() + call spir_func void @_Z17large_grf_markerv() ret void } diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index a4ed14841a196..ada6c5007ed7d 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -711,8 +711,8 @@ void ModuleDesc::dump() const { llvm::errs() << "split_module::ModuleDesc[" << Name << "] {\n"; llvm::errs() << " ESIMD:" << toString(EntryPoints.Props.HasESIMD) << ", SpecConstMet:" << (Props.SpecConstsMet ? "YES" : "NO") - << ", DoubleGRF:" - << (EntryPoints.Props.UsesDoubleGRF ? "YES" : "NO") << "\n"; + << ", LargeGRF:" + << (EntryPoints.Props.UsesLargeGRF ? "YES" : "NO") << "\n"; dumpEntryPoints(entries(), EntryPoints.GroupId.str().c_str(), 1); llvm::errs() << "}\n"; } @@ -744,12 +744,12 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, } std::unique_ptr -getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) { +getLargeGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) { EntryPointGroupVec Groups = groupEntryPointsByAttribute( - MD, sycl::kernel_props::ATTR_DOUBLE_GRF, EmitOnlyKernelsAsEntryPoints, + MD, sycl::kernel_props::ATTR_LARGE_GRF, EmitOnlyKernelsAsEntryPoints, [](EntryPointGroup &G) { - if (G.GroupId == sycl::kernel_props::ATTR_DOUBLE_GRF) { - G.Props.UsesDoubleGRF = true; + if (G.GroupId == sycl::kernel_props::ATTR_LARGE_GRF) { + G.Props.UsesLargeGRF = true; } }); assert(!Groups.empty() && "At least one group is expected"); diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index f362c2c1973da..7088909c8400a 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -55,8 +55,8 @@ struct EntryPointGroup { struct Properties { // Whether all EPs are ESIMD, SYCL or there are both kinds. SyclEsimdSplitStatus HasESIMD = SyclEsimdSplitStatus::SYCL_AND_ESIMD; - // Whether any of the EPs use double GRF mode. - bool UsesDoubleGRF = false; + // Whether any of the EPs use large GRF mode. + bool UsesLargeGRF = false; // Scope represented by EPs in a group EntryPointsGroupScope Scope = Scope_Global; @@ -65,7 +65,7 @@ struct EntryPointGroup { Res.HasESIMD = HasESIMD == Other.HasESIMD ? HasESIMD : SyclEsimdSplitStatus::SYCL_AND_ESIMD; - Res.UsesDoubleGRF = UsesDoubleGRF || Other.UsesDoubleGRF; + Res.UsesLargeGRF = UsesLargeGRF || Other.UsesLargeGRF; // Scope remains global return Res; } @@ -90,8 +90,8 @@ struct EntryPointGroup { bool isSycl() const { return Props.HasESIMD == SyclEsimdSplitStatus::SYCL_ONLY; } - // Tells if some entry points use double GRF mode. - bool isDoubleGRF() const { return Props.UsesDoubleGRF; } + // Tells if some entry points use large GRF mode. + bool isLargeGRF() const { return Props.UsesLargeGRF; } void saveNames(std::vector &Dest) const; void rebuildFromNames(const std::vector &Names, const Module &M); @@ -146,7 +146,7 @@ class ModuleDesc { bool isESIMD() const { return EntryPoints.isEsimd(); } bool isSYCL() const { return EntryPoints.isSycl(); } - bool isDoubleGRF() const { return EntryPoints.isDoubleGRF(); } + bool isLargeGRF() const { return EntryPoints.isLargeGRF(); } const EntryPointSet &entries() const { return EntryPoints.Functions; } const EntryPointGroup &getEntryPointGroup() const { return EntryPoints; } @@ -251,7 +251,7 @@ getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode, bool EmitOnlyKernelsAsEntryPoints); std::unique_ptr -getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); +getLargeGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); #ifndef NDEBUG void dumpEntryPoints(const EntryPointSet &C, const char *msg = "", int Tab = 0); diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 46cbdcceaae56..30ad2da61330c 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -442,8 +442,8 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, if (MD.isESIMD()) { PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isEsimdImage", true}); } - if (MD.isDoubleGRF()) - PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isDoubleGRF", true}); + if (MD.isLargeGRF()) + PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true}); { std::vector FuncNames = getKernelNamesUsingAssert(M); for (const StringRef &FName : FuncNames) @@ -560,8 +560,8 @@ bool lowerEsimdConstructs(module_split::ModuleDesc &MD) { // Compute the filename suffix for the module StringRef getModuleSuffix(const module_split::ModuleDesc &MD) { - if (MD.isDoubleGRF()) { - return MD.isESIMD() ? "_esimd_x2grf" : "_x2grf"; + if (MD.isLargeGRF()) { + return MD.isESIMD() ? "_esimd_large_grf" : "_large_grf"; } return MD.isESIMD() ? "_esimd" : ""; } @@ -735,7 +735,7 @@ processInputModule(std::unique_ptr M) { } Modified |= InvokeSimdMet; - // Lower kernel properties setting APIs before "double GRF" splitting, as: + // Lower kernel properties setting APIs before "large GRF" splitting, as: // - the latter uses the result of the former // - saves processing time Modified |= runModulePass(*M); @@ -774,15 +774,15 @@ processInputModule(std::unique_ptr M) { module_split::ModuleDesc MDesc = ScopedSplitter->nextSplit(); DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1); - std::unique_ptr DoubleGRFSplitter = - module_split::getDoubleGRFSplitter(std::move(MDesc), - EmitOnlyKernelsAsEntryPoints); - const bool SplitByDoubleGRF = DoubleGRFSplitter->totalSplits() > 1; - Modified |= SplitByDoubleGRF; + std::unique_ptr LargeGRFSplitter = + module_split::getLargeGRFSplitter(std::move(MDesc), + EmitOnlyKernelsAsEntryPoints); + const bool SplitByLargeGRF = LargeGRFSplitter->totalSplits() > 1; + Modified |= SplitByLargeGRF; - // Now split further by "esimd-double-grf" attribute. - while (DoubleGRFSplitter->hasMoreSplits()) { - module_split::ModuleDesc MDesc1 = DoubleGRFSplitter->nextSplit(); + // 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(); @@ -821,8 +821,8 @@ processInputModule(std::unique_ptr M) { } if (!MDesc2.isSYCL() && LowerEsimd) { assert(MDesc2.isESIMD() && "NYI"); - // ESIMD lowering also detects double-GRF kernels, so it must happen - // before double-GRF split. + // ESIMD lowering also detects large-GRF kernels, so it must happen + // before large-GRF split. Modified |= lowerEsimdConstructs(MDesc2); } MMs.emplace_back(std::move(MDesc2)); @@ -848,7 +848,7 @@ processInputModule(std::unique_ptr M) { DUMP_ENTRY_POINTS(MMs.back().entries(), MMs.back().Name.c_str(), 3); Modified = true; } - bool SplitOccurred = SplitByScope || SplitByDoubleGRF || SplitByESIMD; + bool SplitOccurred = SplitByScope || SplitByLargeGRF || SplitByESIMD; if (IROutputOnly) { if (SplitOccurred) { diff --git a/sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp b/sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp index 786a396e921a6..1161ec321d4f4 100644 --- a/sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp @@ -34,10 +34,10 @@ template struct proxy { }; } // namespace detail -/// A boolean property which requests the compiler to double the amount of -/// general-purpose registers available to a thread at the expense of reducing -/// the amount of available hardware threads. -struct use_double_grf_tag { +/// A boolean property which requests the compiler to use large register +/// allocation mode at the expense of reducing the amount of available hardware +/// threads. +struct use_large_grf_tag { template friend struct detail::proxy; private: @@ -45,7 +45,9 @@ struct use_double_grf_tag { static constexpr int value = 0; }; -inline constexpr use_double_grf_tag use_double_grf = {}; +__SYCL_DEPRECATED("use_double_grf is deprecated, use use_large_grf instead") +inline constexpr use_large_grf_tag use_double_grf = {}; +inline constexpr use_large_grf_tag use_large_grf = {}; } // namespace kernel_properties @@ -61,14 +63,14 @@ void set_kernel_properties(KernelProps... props) { using Props = __MP11_NS::mp_list; __MP11_NS::mp_for_each([&](auto Prop) { using PropT = decltype(Prop); - constexpr bool IsDoubleGRF = - std::is_same_v; - if constexpr (IsDoubleGRF) { + constexpr bool IsLargeGRF = + std::is_same_v; + if constexpr (IsLargeGRF) { __sycl_set_kernel_properties( kernel_properties::detail::proxy< - kernel_properties::use_double_grf_tag>::value); + kernel_properties::use_large_grf_tag>::value); } else { - static_assert(IsDoubleGRF && + static_assert(IsLargeGRF && "set_kernel_properties: invalid kernel property"); } }); diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 193c187ccf345..7bea37764cd4f 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -404,7 +404,9 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, CompileOpts += std::string(TemporaryStr); } bool isEsimdImage = getUint32PropAsBool(Img, "isEsimdImage"); - bool isDoubleGRF = getUint32PropAsBool(Img, "isDoubleGRF"); + // TODO: Remove isDoubleGRF check in next ABI break + bool isLargeGRF = getUint32PropAsBool(Img, "isLargeGRF") || + getUint32PropAsBool(Img, "isDoubleGRF"); // The -vc-codegen option is always preserved for ESIMD kernels, regardless // of the contents SYCL_PROGRAM_COMPILE_OPTIONS environment variable. if (isEsimdImage) { @@ -416,7 +418,7 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, if (detail::SYCLConfig::get() == 0) CompileOpts += " -disable-finalizer-msg"; } - if (isDoubleGRF) { + if (isLargeGRF) { if (!CompileOpts.empty()) CompileOpts += " "; // TODO: Always use -ze-opt-large-register-file once IGC VC bug ignoring it