[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 <nick.sarnie@intel.com>

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 {

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<llvm::ConstantInt>(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:

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}

llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll → llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll

@@ -1,38 +1,38 @@
 ; 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:
 
 ; 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
 }
 

llvm/test/tools/sycl-post-link/sycl-double-grf.ll → 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
 }
 

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<std::string> &Names,
 }
 
 std::unique_ptr<ModuleSplitterBase>
-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");

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<std::string> &Dest) const;
   void rebuildFromNames(const std::vector<std::string> &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<ModuleSplitterBase>
-getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints);
+getLargeGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints);
 
 #ifndef NDEBUG
 void dumpEntryPoints(const EntryPointSet &C, const char *msg = "", int Tab = 0);

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<StringRef> 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<Module> 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<SYCLLowerKernelPropsPass>(*M);
@@ -774,15 +774,15 @@ processInputModule(std::unique_ptr<Module> M) {
     module_split::ModuleDesc MDesc = ScopedSplitter->nextSplit();
     DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1);
 
-    std::unique_ptr<module_split::ModuleSplitterBase> DoubleGRFSplitter =
-        module_split::getDoubleGRFSplitter(std::move(MDesc),
-                                           EmitOnlyKernelsAsEntryPoints);
-    const bool SplitByDoubleGRF = DoubleGRFSplitter->totalSplits() > 1;
-    Modified |= SplitByDoubleGRF;
+    std::unique_ptr<module_split::ModuleSplitterBase> 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<Module> 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<Module> 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) {