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

This change reimplements the GRF mode specification using native SYCL
kernel properties, does an ABI break, and adds the `automatic` option.

We add a new property named `sycl::detail::register_alloc_mode`, and it
takes in an enum `sycl::detail::register_alloc_mode_enum` which
currently has two values: `automatic` and `large`.

This can be applied to kernels as below:

```
    properties prop{register_alloc_mode<register_alloc_mode_enum::automatic>};
...
 cgh.parallel_for<class Foo>(
          Size, prop, [=](id<1> i) { PA[i] += 2; });
```

We do some `register_alloc_mode` specific work in
`CompileTimePropertiesPass`, we add function metadata named
`RegisterAllocMode` based on the function attribute
`sycl-register-alloc-mode`. This is because llvm-spirv looks for the
metadata. This metadata is how AOT works.

We also do some `register_alloc_mode` specific work in `sycl-post-link`:
1) Split based on the value of the `sycl-register-alloc-mode` function
attribute added in the front end
2) Add a binary property named `sycl-register-alloc-mode` used in the
SYCL runtime.


Note that ESIMD does not work at all yet for `automatic` yet, not for
JIT nor AOT. This is because the VC backend does not support auto GRF. I
have made a feature request for the VC team to add this. `large`
continues to work for JIT only.

---------

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h

@@ -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/PassBuilder.cpp

@@ -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/Passes/PassRegistry.def

@@ -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/CMakeLists.txt

@@ -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/CompileTimePropertiesPass.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/ESIMDVerifier.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/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

@@ -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 = !{}