[SYCL][ABI Break] Add support for per-kernel auto GRF mode specification, and reimplement feature using kernel properties

llvm/lib/Passes/PassBuilder.cpp

@@ -84,7 +84,6 @@
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerInvokeSimd.h"
-#include "llvm/SYCLLowerIR/LowerKernelProps.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/LowerWGScope.h"
 #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"

llvm/lib/Passes/PassRegistry.def

@@ -129,7 +129,6 @@ MODULE_PASS("memprof-module", ModuleMemProfilerPass())
 MODULE_PASS("poison-checking", PoisonCheckingPass())
 MODULE_PASS("pseudo-probe-update", PseudoProbeUpdatePass())
 MODULE_PASS("LowerESIMD", SYCLLowerESIMDPass())
-MODULE_PASS("lower-kernel-props", SYCLLowerKernelPropsPass())
 MODULE_PASS("ESIMDLowerVecArg", ESIMDLowerVecArgPass())
 MODULE_PASS("esimd-opt-call-conv", ESIMDOptimizeVecArgCallConvPass())
 MODULE_PASS("esimd-verifier", ESIMDVerifierPass())

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -59,7 +59,6 @@ add_llvm_component_library(LLVMSYCLLowerIR
   ESIMD/LowerESIMDSlmReservation.cpp
   HostPipes.cpp
   LowerInvokeSimd.cpp
-  LowerKernelProps.cpp
   LowerWGLocalMemory.cpp
   LowerWGScope.cpp
   MutatePrintfAddrspace.cpp

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -10,6 +10,7 @@
 
 #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
 #include "llvm/SYCLLowerIR/DeviceGlobals.h"
+#include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
 #include "llvm/SYCLLowerIR/HostPipes.h"
 
 #include "llvm/ADT/APInt.h"
@@ -27,6 +28,7 @@ namespace {
 
 constexpr StringRef SYCL_HOST_ACCESS_ATTR = "sycl-host-access";
 constexpr StringRef SYCL_PIPELINED_ATTR = "sycl-pipelined";
+constexpr StringRef SYCL_REGISTER_ALLOC_MODE_ATTR = "sycl-register-alloc-mode";
 
 constexpr StringRef SPIRV_DECOR_MD_KIND = "spirv.Decorations";
 constexpr StringRef SPIRV_PARAM_DECOR_MD_KIND = "spirv.ParameterDecorations";
@@ -175,14 +177,15 @@ MDNode *attributeToDecorateMetadata(LLVMContext &Ctx, const Attribute &Attr) {
 /// Tries to generate a SPIR-V execution mode metadata node from an attribute.
 /// If the attribute is unknown \c None will be returned.
 ///
-/// @param M     [in] the LLVM module.
+/// @param F     [in] the LLVM function.
 /// @param Attr  [in] the LLVM attribute to generate metadata for.
 ///
 /// @returns a pair with the name of the resulting metadata and a pointer to
 ///          the metadata node with its values if the attribute has a
 ///          corresponding SPIR-V execution mode. Otherwise \c None is returned.
 std::optional<std::pair<std::string, MDNode *>>
-attributeToExecModeMetadata(Module &M, const Attribute &Attr) {
+attributeToExecModeMetadata(Function &F, const Attribute &Attr) {
+  Module &M = *F.getParent();
   LLVMContext &Ctx = M.getContext();
   const DataLayout &DLayout = M.getDataLayout();
 
@@ -267,6 +270,15 @@ attributeToExecModeMetadata(Module &M, const Attribute &Attr) {
     return std::pair<std::string, MDNode *>("ip_interface",
                                             getIpInterface("csr", Ctx, Attr));
 
+  if (AttrKindStr == SYCL_REGISTER_ALLOC_MODE_ATTR &&
+      !llvm::esimd::isESIMD(F)) {
+    uint32_t RegAllocModeVal = getAttributeAsInteger<uint32_t>(Attr);
+    Metadata *AttrMDArgs[] = {ConstantAsMetadata::get(Constant::getIntegerValue(
+        Type::getInt32Ty(Ctx), APInt(32, RegAllocModeVal)))};
+    return std::pair<std::string, MDNode *>("RegisterAllocMode",
+                                            MDNode::get(Ctx, AttrMDArgs));
+  }
+
   return std::nullopt;
 }
 
@@ -420,7 +432,7 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M,
       } else if (MDNode *SPIRVMetadata =
                      attributeToDecorateMetadata(Ctx, Attribute))
         MDOps.push_back(SPIRVMetadata);
-      else if (auto NamedMetadata = attributeToExecModeMetadata(M, Attribute))
+      else if (auto NamedMetadata = attributeToExecModeMetadata(F, Attribute))
         NamedMDOps.push_back(*NamedMetadata);
     }
 

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -59,7 +59,6 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::exp<.+>",
     "^sycl::_V1::bit_cast<.+>",
     "^sycl::_V1::operator.+<.+>",
-    "^sycl::_V1::ext::intel::experimental::set_kernel_properties",
     "^sycl::_V1::ext::oneapi::sub_group::.+",
     "^sycl::_V1::ext::oneapi::experimental::spec_constant<.+>::.+",
     "^sycl::_V1::ext::oneapi::experimental::this_sub_group",

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

@@ -1969,8 +1969,6 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F,
         ToErase.push_back(CI);
         continue;
       }
-      assert(!Name.startswith("__sycl_set_kernel_properties") &&
-             "__sycl_set_kernel_properties must have been lowered");
 
       if (Name.empty() ||
           (!Name.startswith(ESIMD_INTRIN_PREF1) && !isDevicelibFunction(Name)))

llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/register-alloc-mode.ll

@@ -0,0 +1,35 @@
+; Check we create RegisterAllocMode metadata if there is a non-ESIMD kernel with that property
+; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s --check-prefix CHECK-IR
+
+; Function Attrs: convergent norecurse
+define weak_odr dso_local spir_kernel void @sycl_regallocmode() #1 {
+; CHECK-IR-NOT: !RegisterAllocMode
+; CHECK-IR: sycl_regallocmode() #[[#Attr1:]]{{.*}}!RegisterAllocMode ![[#MDVal:]] {
+; CHECK-IR-NOT: !RegisterAllocMode
+; CHECK-IR: ![[#MDVal]] = !{i32 2}
+entry:
+  ret void
+}
+
+; Function Attrs: convergent norecurse
+define weak_odr dso_local spir_kernel void @sycl_noregallocmode() #0 {
+entry:
+  ret void
+}
+
+; Function Attrs: convergent norecurse
+define weak_odr dso_local spir_kernel void @esimd_regallocmode() #1 !sycl_explicit_simd !1 {
+entry:
+  ret void
+}
+
+; Function Attrs: convergent norecurse
+define weak_odr dso_local spir_kernel void @esimd_noregallocmode() #0 {
+entry:
+  ret void
+}
+
+attributes #0 = { convergent norecurse }
+attributes #1 = { convergent norecurse "sycl-register-alloc-mode"="2" }
+
+!1 = !{}

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

@@ -1,34 +1,41 @@
-; This test checks handling of the
-;   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
-;   "large GRF" and callgraphs for entry points which are not
-; - Compiler adds 'isLargeGRF' property to the ESIMD device binary
-;   images requesting "large GRF" 
+; This test checks handling of RegisterAllocMode in SYCL post link
 
 ; 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_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR  --implicit-check-not='RegisterAllocMode'
+; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR
 ; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
+; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-SYCL-LargeGRF-IR
+; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-SYCL-LargeGRF-PROP
+; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-LargeGRF-SYM
+; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-SYCL-PROP
 ; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM
-; RUN: FileCheck %s -input-file=%t_esimd_1.sym --check-prefixes CHECK-ESIMD-SYM
+; RUN: FileCheck %s -input-file=%t_esimd_1.prop --check-prefixes CHECK-ESIMD-PROP
 ; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM
 
 ; CHECK: [Code|Properties|Symbols]
 ; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
+; CHECK: {{.*}}_esimd_0.ll|{{.*}}_esimd_0.prop|{{.*}}_esimd_0.sym
 ; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
 ; CHECK: {{.*}}_esimd_1.ll|{{.*}}_esimd_1.prop|{{.*}}_esimd_1.sym
 
 ; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
-; CHECK-ESIMD-LargeGRF-PROP: isLargeGRF=1|1
+; CHECK-ESIMD-LargeGRF-PROP: sycl-register-alloc-mode=1|2
+
+; CHECK-SYCL-LargeGRF-PROP: sycl-register-alloc-mode=1|2
+
+; CHECK-SYCL-PROP-NOT: sycl-register-alloc-mode
 
 ; CHECK-SYCL-SYM: __SYCL_kernel
 ; CHECK-SYCL-SYM-EMPTY:
 
+; CHECK-SYCL-LargeGRF-SYM: __SYCL_kernel_large_grf
+; CHECK-SYCL-LargeGRF-SYM-EMPTY:
+
 ; CHECK-ESIMD-SYM: __ESIMD_kernel
 ; CHECK-ESIMD-SYM-EMPTY:
 
+; CHECK-ESIMD-PROP-NOT: sycl-register-alloc-mode
+
 ; CHECK-ESIMD-LargeGRF-SYM: __ESIMD_large_grf_kernel
 ; CHECK-ESIMD-LargeGRF-SYM-EMPTY:
 
@@ -42,29 +49,26 @@ entry:
   ret void
 }
 
-define weak_odr dso_local spir_kernel void @__ESIMD_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
+define weak_odr dso_local spir_kernel void @__SYCL_kernel_large_grf() #1 {
+; CHECK-SYCL-LargeGRF-IR: define {{.*}} spir_kernel void @__SYCL_kernel_large_grf() #[[#Attr:]]
 entry:
   ret void
 }
 
-define dso_local spir_func void @_Z17large_grf_markerv() {
+define weak_odr dso_local spir_kernel void @__ESIMD_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
+
 entry:
-  call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
-; -- Check that ESIMD lowering removed the marker call above:
-; 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_large_grf_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
-; CHECK-ESIMD-LargeGRF-IR: @__ESIMD_large_grf_kernel() {{.*}}
+define weak_odr dso_local spir_kernel void @__ESIMD_large_grf_kernel() #1 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
+; CHECK-ESIMD-LargeGRF-IR: @__ESIMD_large_grf_kernel()
 entry:
-  call spir_func void @_Z17large_grf_markerv()
   ret void
 }
 
 attributes #0 = { "sycl-module-id"="a.cpp" }
+attributes #1 = { "sycl-module-id"="a.cpp" "sycl-register-alloc-mode"="2" }
 
 !0 = !{}
 !1 = !{i32 1}