[SYCL] Add experimental implementation of sycl_ext_intel_grf_size (#9882

)

This change implements `sycl_ext_intel_grf_size`, and in particular:

1) Moves the `sycl_ext_intel_grf_size` spec document from the `proposed`
folder to the `experimental` folder, and updates the implementation
status in the document to match.
2) Adds two new kernel properties
`sycl::ext::intel::experimental::grf_size` and
`sycl::ext::intel::experimental::grf_size_automatic`, as per the spec.
`grf_size` adds the `sycl-grf-size` metadata with a value of the
template parameter (`128` or `256`). `grf_size_automatic` adds the
`sycl-grf-size` metadata with a value of `0`.
3) Marks the `sycl::detail::register_alloc_mode` property as deprecated,
and it still works as before.
4) Updates `CompileTimePropertiesPass.cpp` to map the `sycl-grf-size`
metadata added by the front-end to the `RegisterAllocMode` metadata
which `llvm-spirv` looks for. This `RegisterAllocMode` metadata is how
AOT works.
5) Updates `sycl-post-link` to split by the `sycl-grf-size` metadata,
have a `sycl-grf-size` binary property, and do a error check to make
sure the deprecated `sycl::detail::register_alloc_mode` and
`grf_size`/`grf_size_automatic` are not set at the same time on the same
kernel.
6) Updates `program_manager` to deal with the new image property and
pass the right flags
7) Updates exists tests to also test the new properties.

---------

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -29,6 +29,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 SYCL_GRF_SIZE_ATTR = "sycl-grf-size";
 
 constexpr StringRef SPIRV_DECOR_MD_KIND = "spirv.Decorations";
 constexpr StringRef SPIRV_PARAM_DECOR_MD_KIND = "spirv.ParameterDecorations";
@@ -270,11 +271,24 @@ attributeToExecModeMetadata(const Attribute &Attr, Function &F) {
     return std::pair<std::string, MDNode *>("ip_interface",
                                             getIpInterface("csr", Ctx, Attr));
 
-  if (AttrKindStr == SYCL_REGISTER_ALLOC_MODE_ATTR &&
+  if ((AttrKindStr == SYCL_REGISTER_ALLOC_MODE_ATTR ||
+       AttrKindStr == SYCL_GRF_SIZE_ATTR) &&
       !llvm::esimd::isESIMD(F)) {
-    uint32_t RegAllocModeVal = getAttributeAsInteger<uint32_t>(Attr);
-    Metadata *AttrMDArgs[] = {ConstantAsMetadata::get(Constant::getIntegerValue(
-        Type::getInt32Ty(Ctx), APInt(32, RegAllocModeVal)))};
+    // TODO: Remove SYCL_REGISTER_ALLOC_MODE_ATTR support in next ABI break.
+    uint32_t PropVal = getAttributeAsInteger<uint32_t>(Attr);
+    if (AttrKindStr == SYCL_GRF_SIZE_ATTR) {
+      assert((PropVal == 0 || PropVal == 128 || PropVal == 256) &&
+             "Unsupported GRF Size");
+      // Map sycl-grf-size values to RegisterAllocMode values used in SPIR-V.
+      static constexpr int SMALL_GRF_REGALLOCMODE_VAL = 1;
+      static constexpr int LARGE_GRF_REGALLOCMODE_VAL = 2;
+      if (PropVal == 128)
+        PropVal = SMALL_GRF_REGALLOCMODE_VAL;
+      else if (PropVal == 256)
+        PropVal = LARGE_GRF_REGALLOCMODE_VAL;
+    }
+    Metadata *AttrMDArgs[] = {ConstantAsMetadata::get(
+        Constant::getIntegerValue(Type::getInt32Ty(Ctx), APInt(32, PropVal)))};
     return std::pair<std::string, MDNode *>("RegisterAllocMode",
                                             MDNode::get(Ctx, AttrMDArgs));
   }

llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/grf-size.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_grf_size() #1 {
+; CHECK-IR-NOT: !RegisterAllocMode
+; CHECK-IR: sycl_grf_size() #[[#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_no_grf_size() #0 {
+entry:
+  ret void
+}
+
+; Function Attrs: convergent norecurse
+define weak_odr dso_local spir_kernel void @esimd_grf_size() #1 !sycl_explicit_simd !1 {
+entry:
+  ret void
+}
+
+; Function Attrs: convergent norecurse
+define weak_odr dso_local spir_kernel void @esimd_no_grf_size() #0 {
+entry:
+  ret void
+}
+
+attributes #0 = { convergent norecurse }
+attributes #1 = { convergent norecurse "sycl-grf-size"="256" }
+
+!1 = !{}

llvm/test/tools/sycl-post-link/grf-size-conflict.ll

@@ -0,0 +1,16 @@
+; This test confirms an error with sycl-register-alloc-mode and sycl-grf-size on the same kernel.
+
+; RUN: not sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s 2>&1 | FileCheck %s
+
+; CHECK: Unsupported use of both register_alloc_mode and grf_size
+
+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 weak_odr dso_local spir_kernel void @__SYCL_kernel() #0 {
+entry:
+  ret void
+}
+
+attributes #0 = { "sycl-module-id"="a.cpp" "sycl-grf-size"="256" "sycl-register-alloc-mode"="0"}

llvm/test/tools/sycl-post-link/sycl-grf-size.ll

@@ -0,0 +1,84 @@
+; This test checks handling of sycl-grf-size 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='__ESIMD_kernel()'
+; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
+; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM
+; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-SYCL-LargeGRF-IR --implicit-check-not='__SYCL_kernel()'
+; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-SYCL-LargeGRF-PROP
+; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-LargeGRF-SYM
+; RUN: FileCheck %s -input-file=%t_3.ll --check-prefixes CHECK-SYCL-IR --implicit-check-not='__SYCL_kernel_large_grf()'
+; RUN: FileCheck %s -input-file=%t_3.prop --check-prefixes CHECK-SYCL-PROP
+; RUN: FileCheck %s -input-file=%t_3.sym --check-prefixes CHECK-SYCL-SYM
+; RUN: FileCheck %s -input-file=%t_esimd_2.ll --check-prefixes CHECK-ESIMD-IR --implicit-check-not='__ESIMD_large_grf_kernel()'
+; RUN: FileCheck %s -input-file=%t_esimd_2.prop --check-prefixes CHECK-ESIMD-PROP
+
+; CHECK: [Code|Properties|Symbols]
+; CHECK: {{.*}}_esimd_0.ll|{{.*}}_esimd_0.prop|{{.*}}_esimd_0.sym
+; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
+; CHECK: {{.*}}_esimd_2.ll|{{.*}}_esimd_2.prop|{{.*}}_esimd_2.sym
+
+; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
+; CHECK-ESIMD-LargeGRF-PROP: sycl-grf-size=1|256
+
+; CHECK-SYCL-LargeGRF-PROP: sycl-grf-size=1|256
+
+; CHECK-SYCL-LargeGRF-IR: define {{.*}} spir_kernel void @__SYCL_kernel_large_grf() #[[SYCLAttr:]]
+; CHECK-SYCL-LargeGRF-IR: attributes #[[SYCLAttr]]
+
+; CHECK-SYCL-PROP-NOT: sycl-grf-size
+
+; CHECK-SYCL-SYM: __SYCL_kernel
+; CHECK-SYCL-SYM-EMPTY:
+
+; CHECK-SYCL-IR: __SYCL_kernel() #[[SYCLAttr:]]
+; CHECK-SYCL-IR: attributes #[[SYCLAttr]]
+
+; CHECK-SYCL-LargeGRF-SYM: __SYCL_kernel_large_grf
+; CHECK-SYCL-LargeGRF-SYM-EMPTY:
+
+; CHECK-ESIMD-SYM: __ESIMD_kernel
+; CHECK-ESIMD-SYM-EMPTY:
+
+; CHECK-ESIMD-IR: __ESIMD_kernel() #[[ESIMDAttr:]]
+; CHECK-ESIMD-IR: attributes #[[ESIMDAttr]]
+
+; CHECK-ESIMD-PROP-NOT: sycl-grf-size
+
+; CHECK-ESIMD-LargeGRF-SYM: __ESIMD_large_grf_kernel
+; CHECK-ESIMD-LargeGRF-SYM-EMPTY:
+
+; CHECK-ESIMD-LargeGRF-IR: @__ESIMD_large_grf_kernel() #[[ESIMDLargeAttr:]]
+; CHECK-ESIMD-LargeGRF-IR: attributes #[[ESIMDLargeAttr]]
+
+; ModuleID = 'large_grf.bc'
+source_filename = "grf"
+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 weak_odr dso_local spir_kernel void @__SYCL_kernel() #0 {
+entry:
+  ret void
+}
+
+define weak_odr dso_local spir_kernel void @__SYCL_kernel_large_grf() #1 {
+entry:
+  ret void
+}
+
+define weak_odr dso_local spir_kernel void @__ESIMD_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
+entry:
+  ret void
+}
+
+define weak_odr dso_local spir_kernel void @__ESIMD_large_grf_kernel() #1 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
+entry:
+  ret void
+}
+
+attributes #0 = { "sycl-module-id"="a.cpp" }
+attributes #1 = { "sycl-module-id"="a.cpp" "sycl-grf-size"="256" }
+
+!0 = !{}
+!1 = !{i32 1}

llvm/tools/sycl-post-link/ModuleSplitter.cpp

@@ -854,6 +854,7 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
     // Note: Add more rules at the end of the list to avoid chaning orders of
     // output files in existing tests.
     Categorizer.registerSimpleStringAttributeRule("sycl-register-alloc-mode");
+    Categorizer.registerSimpleStringAttributeRule("sycl-grf-size");
     Categorizer.registerListOfIntegersInMetadataSortedRule("sycl_used_aspects");
     Categorizer.registerListOfIntegersInMetadataRule("reqd_work_group_size");
     Categorizer.registerListOfIntegersInMetadataRule(

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -454,12 +454,12 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
   if (MD.isESIMD()) {
     PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isEsimdImage", true});
   }
-
+  bool HasRegAllocMode = false;
   {
     StringRef RegAllocModeAttr = "sycl-register-alloc-mode";
     uint32_t RegAllocModeVal;
 
-    bool HasRegAllocMode = llvm::any_of(MD.entries(), [&](const Function *F) {
+    HasRegAllocMode = llvm::any_of(MD.entries(), [&](const Function *F) {
       if (!F->hasFnAttribute(RegAllocModeAttr))
         return false;
       const auto &Attr = F->getFnAttribute(RegAllocModeAttr);
@@ -472,6 +472,25 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
     }
   }
 
+  {
+    StringRef GRFSizeAttr = "sycl-grf-size";
+    uint32_t GRFSizeVal;
+
+    bool HasGRFSize = llvm::any_of(MD.entries(), [&](const Function *F) {
+      if (!F->hasFnAttribute(GRFSizeAttr))
+        return false;
+      const auto &Attr = F->getFnAttribute(GRFSizeAttr);
+      GRFSizeVal = getAttributeAsInteger<uint32_t>(Attr);
+      return true;
+    });
+    if (HasGRFSize) {
+      if (HasRegAllocMode)
+        error("Unsupported use of both register_alloc_mode and "
+              "grf_size");
+      PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({GRFSizeAttr, GRFSizeVal});
+    }
+  }
+
   // FIXME: Remove 'if' below when possible
   // GPU backend has a problem with accepting optimization level options in form
   // described by Level Zero specification (-ze-opt-level=1) when 'invoke_simd'

sycl/doc/extensions/proposed/sycl_ext_intel_grf_size.asciidoc → sycl/doc/extensions/experimental/sycl_ext_intel_grf_size.asciidoc

@@ -49,12 +49,11 @@ This extension also depends on the following other SYCL extensions:
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
-
+This is an experimental extension specification, intended to provide early access
+to features and gather community feedback. Interfaces defined in this specification
+are implemented in DPC++, but they are not finalized and may change incompatibly in
+future versions of DPC++ without prior notice. **Shipping software products should not
+rely on APIs defined in this specification.**
 
 == Backend support status
 

sycl/include/sycl/detail/kernel_properties.hpp

@@ -29,7 +29,10 @@ struct register_alloc_mode_key {
 };
 
 template <register_alloc_mode_enum Mode>
-inline constexpr register_alloc_mode_key::value_t<Mode> register_alloc_mode;
+inline constexpr register_alloc_mode_key::value_t<Mode> register_alloc_mode
+    __SYCL_DEPRECATED("register_alloc_mode is deprecated, "
+                      "use sycl::ext::intel::experimental::grf_size or "
+                      "sycl::ext::intel::experimental::grf_size_automatic");
 } // namespace detail
 
 namespace ext::oneapi::experimental {

sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp

@@ -0,0 +1,92 @@
+//==- grf_size_properties.hpp - GRF size kernel properties for Intel GPUs -==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===--------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/ext/oneapi/properties/property.hpp>
+#include <sycl/ext/oneapi/properties/property_value.hpp>
+
+#define SYCL_EXT_INTEL_GRF_SIZE 1
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext::intel::experimental {
+struct grf_size_key {
+  template <unsigned int Size>
+  using value_t = oneapi::experimental::property_value<
+      grf_size_key, std::integral_constant<unsigned int, Size>>;
+};
+
+struct grf_size_automatic_key {
+  using value_t = oneapi::experimental::property_value<grf_size_automatic_key>;
+};
+
+template <unsigned int Size>
+inline constexpr grf_size_key::value_t<Size> grf_size;
+
+inline constexpr grf_size_automatic_key::value_t grf_size_automatic;
+
+} // namespace ext::intel::experimental
+namespace ext::oneapi::experimental {
+template <>
+struct is_property_key<sycl::ext::intel::experimental::grf_size_key>
+    : std::true_type {};
+
+template <>
+struct is_property_key<sycl::ext::intel::experimental::grf_size_automatic_key>
+    : std::true_type {};
+
+namespace detail {
+template <>
+struct PropertyToKind<sycl::ext::intel::experimental::grf_size_key> {
+  static constexpr PropKind Kind = PropKind::GRFSize;
+};
+
+template <>
+struct IsCompileTimeProperty<sycl::ext::intel::experimental::grf_size_key>
+    : std::true_type {};
+
+template <>
+struct PropertyToKind<sycl::ext::intel::experimental::grf_size_automatic_key> {
+  static constexpr PropKind Kind = PropKind::GRFSizeAutomatic;
+};
+
+template <>
+struct IsCompileTimeProperty<
+    sycl::ext::intel::experimental::grf_size_automatic_key> : std::true_type {};
+
+template <unsigned int Size>
+struct PropertyMetaInfo<
+    sycl::ext::intel::experimental::grf_size_key::value_t<Size>> {
+  static_assert(Size == 128 || Size == 256, "Unsupported GRF size");
+  static constexpr const char *name = "sycl-grf-size";
+  static constexpr unsigned int value = Size;
+};
+template <>
+struct PropertyMetaInfo<
+    sycl::ext::intel::experimental::grf_size_automatic_key::value_t> {
+  static constexpr const char *name = "sycl-grf-size";
+  static constexpr unsigned int value = 0;
+};
+
+template <typename Properties>
+struct ConflictingProperties<sycl::ext::intel::experimental::grf_size_key,
+                             Properties>
+    : ContainsProperty<sycl::ext::intel::experimental::grf_size_automatic_key,
+                       Properties> {};
+
+template <typename Properties>
+struct ConflictingProperties<
+    sycl::ext::intel::experimental::grf_size_automatic_key, Properties>
+    : ContainsProperty<sycl::ext::intel::experimental::grf_size_key,
+                       Properties> {};
+
+} // namespace detail
+} // namespace ext::oneapi::experimental
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/ext/oneapi/properties/properties.hpp

@@ -131,6 +131,8 @@ template <typename PropertiesT> class properties {
                 "Properties in property list are not sorted.");
   static_assert(detail::SortedAllUnique<PropertiesT>::value,
                 "Duplicate properties in property list.");
+  static_assert(detail::NoConflictingProperties<PropertiesT>::value,
+                "Conflicting properties in property list.");
 
 public:
   template <typename... PropertyValueTs>

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -195,8 +195,10 @@ enum PropKind : uint32_t {
   UsesValid = 29,
   UseRootSync = 30,
   RegisterAllocMode = 31,
+  GRFSize = 32,
+  GRFSizeAutomatic = 33,
   // PropKindSize must always be the last value.
-  PropKindSize = 32,
+  PropKindSize = 34,
 };
 
 // This trait must be specialized for all properties and must have a unique