[SYCL][Docs] Add sycl_ext_named_sub_group_sizes kernel properties

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

@@ -51,11 +51,12 @@ 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 {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
 
 
 == Overview

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

@@ -127,6 +127,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")
 #include <sycl/ext/oneapi/experimental/group_helpers_sorters.hpp>
 #include <sycl/ext/oneapi/experimental/group_load_store.hpp>
 #include <sycl/ext/oneapi/experimental/group_sort.hpp>
+#include <sycl/ext/oneapi/experimental/named_sub_group_sizes.hpp>
 #include <sycl/ext/oneapi/experimental/prefetch.hpp>
 #include <sycl/ext/oneapi/experimental/profiling_tag.hpp>
 #include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>

sycl/source/feature_test.hpp.in

@@ -124,6 +124,7 @@ inline namespace _V1 {
 #define SYCL_KHR_DEFAULT_CONTEXT 1
 #define SYCL_EXT_INTEL_EVENT_MODE 1
 #define SYCL_EXT_ONEAPI_TANGLE 1
+#define SYCL_EXT_ONEAPI_NAMED_SUB_GROUP_SIZES 1
 
 // Unfinished KHR extensions. These extensions are only available if the
 // __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS macro is defined.

sycl/test/extensions/properties/properties_kernel_named_sub_group_size.cpp

@@ -0,0 +1,42 @@
+// 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>
+
+struct SGSizePrimaryKernelFunctor {
+  SGSizePrimaryKernelFunctor() {}
+
+  void operator()(sycl::nd_item<1>) const {}
+
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::sub_group_size_primary};
+  }
+};
+
+struct SGSizeAutoKernelFunctor {
+  SGSizeAutoKernelFunctor() {}
+
+  void operator()(sycl::nd_item<1>) const {}
+
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::sub_group_size_automatic};
+  }
+};
+
+int main() {
+  sycl::queue Q;
+  sycl::nd_range<1> NDRange{6, 2};
+
+  // CHECK: spir_kernel void @{{.*}}SGSizePrimaryKernelFunctor()
+  // CHECK-SAME: !intel_reqd_sub_group_size ![[SGSizeAttr:[0-9]+]]
+  Q.parallel_for(NDRange, SGSizePrimaryKernelFunctor{});
+
+  // CHECK: spir_kernel void @{{.*}}SGSizeAutoKernelFunctor()
+  // CHECK-NOT: intel_reqd_sub_group_size
+  // CHECK-SAME: {
+  Q.parallel_for(NDRange, SGSizeAutoKernelFunctor{});
+}
+
+// CHECK: ![[SGSizeAttr]] = !{i32 -1}