[SYCL] Add sycl_ext_named_sub_group_sizes kernel properties

sycl/include/sycl/ext/oneapi/experimental/named_sub_group_sizes.hpp

@@ -0,0 +1,45 @@
+//== named_sub_group_sizes.hpp --- SYCL extension for named sub-group sizes ==//
+//
+// 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/kernel_properties/properties.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+struct named_sub_group_size {
+  static constexpr uint32_t primary = -1;
+  static constexpr uint32_t automatic = -2;
+};
+
+inline constexpr sub_group_size_key::value_t<named_sub_group_size::primary>
+    sub_group_size_primary;
+
+inline constexpr sub_group_size_key::value_t<named_sub_group_size::automatic>
+    sub_group_size_automatic;
+
+namespace detail {
+template <>
+struct PropertyMetaInfo<
+    sub_group_size_key::value_t<named_sub_group_size::automatic>> {
+  // sub_group_size_automatic means that the kernel can be compiled with
+  // any sub-group size. That is, if the kernel has the sub_group_size_automatic
+  // property, then no sycl-sub-group-size IR attribute needs to be attached.
+  // Specializing PropertyMetaInfo for sub_group_size_automatic and setting
+  // name to an empty string will result in no sycl-sub-group-size IR being
+  // attached.
+  static constexpr const char *name = "";
+  static constexpr const char *value = 0;
+};
+} // namespace detail
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/sycl.hpp

@@ -88,6 +88,7 @@
 #include <sycl/ext/oneapi/experimental/composite_device.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
 #include <sycl/ext/oneapi/experimental/fixed_size_group.hpp>
+#include <sycl/ext/oneapi/experimental/named_sub_group_sizes.hpp>
 #include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
 #include <sycl/ext/oneapi/experimental/prefetch.hpp>
 #include <sycl/ext/oneapi/experimental/root_group.hpp>

sycl/test/extensions/properties/properties_kernel_named_sub_group_size.cpp

@@ -0,0 +1,24 @@
+// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+// expected-no-diagnostics
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::queue q;
+  sycl::nd_range<1> ndr{6, 2};
+
+  // CHECK: spir_kernel void @{{.*}}Kernel1()
+  // CHECK-SAME: !intel_reqd_sub_group_size ![[SGSizeAttr:[0-9]+]]
+  sycl::ext::oneapi::experimental::properties P1{
+      sycl::ext::oneapi::experimental::sub_group_size_primary};
+  q.parallel_for<class Kernel1>(ndr, P1, [=](auto id) {});
+
+  // CHECK: spir_kernel void @{{.*}}Kernel2()
+  // CHECK-NOT: intel_reqd_sub_group_size
+  // CHECK-SAME: {
+  sycl::ext::oneapi::experimental::properties P2{
+      sycl::ext::oneapi::experimental::sub_group_size_automatic};
+  q.parallel_for<class Kernel2>(ndr, P2, [=](auto id) {});
+}
+
+// CHECK: ![[SGSizeAttr]] = !{i32 -1}