Fix test, add feature macro and move extension to experimental

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/doc/extensions/proposed/sycl_ext_oneapi_named_sub_group_sizes.asciidoc renamed to 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/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

@@ -3,22 +3,40 @@
 // 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> ndr{6, 2};
+  sycl::queue Q;
+  sycl::nd_range<1> NDRange{6, 2};
 
-  // CHECK: spir_kernel void @{{.*}}Kernel1()
+  // CHECK: spir_kernel void @{{.*}}SGSizePrimaryKernelFunctor()
   // 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) {});
+  Q.parallel_for(NDRange, SGSizePrimaryKernelFunctor{});
 
-  // CHECK: spir_kernel void @{{.*}}Kernel2()
+  // CHECK: spir_kernel void @{{.*}}SGSizeAutoKernelFunctor()
   // 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) {});
+  Q.parallel_for(NDRange, SGSizeAutoKernelFunctor{});
 }
 
 // CHECK: ![[SGSizeAttr]] = !{i32 -1}