[SYCL] Correct spec about constness of get(properties_tag) member and add respective warning (#17947)

HPS-1 · web-flow · commit 507679af571e · 2025-04-14T14:29:21.000Z
Currently some of our docs suggest using `get(properties_tag) {...}` member in kernel functors to specify kernel properties while some suggest using `get(properties_tag) const {...}`. But actually `get(properties_tag) {...}` might not work under certain conditions, so we should update the incorrect docs that they suggest readers to use `get(properties_tag) const {...}` rather than `get(properties_tag) {...}`. Also as @Pennycook suggested, added a diagnostic warning to indicate the user when: -they have declared a `get(properties_tag)` member in the kernel functor AND -they declared `get(properties_tag){...}` rather than `get(properties_tag) const {...}` Updated the tests containing the non-const version of `get(properties_tag)` as well. --------- Signed-off-by: Hu, Peisen <peisen.hu@intel.com>
diff --git a/sycl/doc/design/CompileTimeProperties.md b/sycl/doc/design/CompileTimeProperties.md
@@ -298,7 +298,7 @@ void foo(handler &cgh) {
 ```
 
 The second way an application can specify kernel properties is by adding a
-member function named `get(sycl::ext::oneapi::properties_tag)` to a named
+const member function named `get(sycl::ext::oneapi::properties_tag)` to a named
 kernel function object:
 
 ```
@@ -309,7 +309,7 @@ class MyKernel {
  public:
   void operator()() {/* ... */}
 
-  auto get(properties_tag) {
+  auto get(properties_tag) const {
     return properties{sub_group_size<32>, device_has<aspect::fp16>};
   }
 };
diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
@@ -295,7 +295,7 @@ that it depends upon for correctness.
 
 To enable this use-case, this extension adds a mechanism for implementations to
 extract a property list from a kernel functor, if a kernel functor declares
-a member function named `get` accepting a `sycl::ext::oneapi::experimental::properties_tag`
+a const member function named `get` accepting a `sycl::ext::oneapi::experimental::properties_tag`
 tag type and returning an instance of `sycl::ext::oneapi::experimental::properties`.
 
 ```c++
@@ -338,7 +338,7 @@ struct KernelFunctor {
     a[i] = b[i] + c[i];
   }
 
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::work_group_size<8, 8>,
                                                        sycl::ext::oneapi::experimental::sub_group_size<8>};
   }
diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc
@@ -319,7 +319,7 @@ struct KernelFunctor {
     *a = *b + *c;
   }
 
-  auto get(properties_tag) {
+  auto get(properties_tag) const {
     return properties{streaming_interface_accept_downstream_stall};
   }
 
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
@@ -1596,6 +1596,17 @@ class __SYCL_EXPORT handler {
             KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
       }
 #endif
+      // Note: the static_assert below need to be run on both the host and the
+      // device ends to avoid test issues, so don't put it into the #ifdef
+      // __SYCL_DEVICE_ONLY__ directive above print out diagnostic message if
+      // the kernel functor has a get(properties_tag) member, but it's not const
+      static_assert(
+          (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
+              const KernelType &>::value) ||
+              !(ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
+                  KernelType>::value),
+          "get(sycl::ext::oneapi::experimental::properties_tag) member in "
+          "kernel functor class must be declared as a const member function");
       auto L = [&](auto &&...args) {
         if constexpr (WrapAsVal == WrapAs::single_task) {
           h->kernel_single_task<KernelName, KernelType, MergedProps...>(
diff --git a/sycl/include/syclcompat/launch_policy.hpp b/sycl/include/syclcompat/launch_policy.hpp
@@ -219,7 +219,7 @@ struct KernelFunctor {
       : _kernel_properties{kernel_props}, _local_acc{local_acc},
         _argument_tuple(std::make_tuple(args...)) {}
 
-  auto get(sycl_exp::properties_tag) { return _kernel_properties; }
+  auto get(sycl_exp::properties_tag) const { return _kernel_properties; }
 
   __syclcompat_inline__ void
   operator()(syclcompat::detail::range_to_item_t<Range>) const {
diff --git a/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp b/sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp
@@ -52,7 +52,7 @@ template <size_t I> struct KernelFunctorWithMaxWGSizeProp {
   void operator()(nd_item<1>) const {}
   void operator()(item<1>) const {}
 
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::max_linear_work_group_size<I>};
   }
diff --git a/sycl/test-e2e/Basic/max_work_group_size_props.cpp b/sycl/test-e2e/Basic/max_work_group_size_props.cpp
@@ -43,7 +43,7 @@ template <size_t... Is> struct KernelFunctorWithMaxWGSizeProp {
   void operator()(nd_item<sizeof...(Is)>) const {}
   void operator()(item<sizeof...(Is)>) const {}
 
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::max_work_group_size<Is...>};
   }
diff --git a/sycl/test-e2e/Basic/sub_group_size_prop.cpp b/sycl/test-e2e/Basic/sub_group_size_prop.cpp
@@ -24,7 +24,7 @@ template <size_t SGSize> struct KernelFunctorWithSGSizeProp {
       Acc[0] = SG.get_local_linear_range();
   }
 
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::sub_group_size<SGSize>};
   }
diff --git a/sycl/test-e2e/Basic/work_group_size_prop.cpp b/sycl/test-e2e/Basic/work_group_size_prop.cpp
@@ -39,7 +39,7 @@ template <size_t... Is> struct KernelFunctorWithWGSizeProp {
   void operator()(nd_item<sizeof...(Is)>) const {}
   void operator()(item<sizeof...(Is)>) const {}
 
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::work_group_size<Is...>};
   }
diff --git a/sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp b/sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp
@@ -19,7 +19,7 @@ template <size_t SGSize> struct KernelFunctorWithSGSizeProp {
       Acc[0] = SG.get_local_linear_range();
   }
 
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::sub_group_size<SGSize>};
   }
diff --git a/sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp b/sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp
@@ -34,7 +34,7 @@ template <size_t... Is> struct KernelFunctorWithWGSizeProp {
   void operator()(nd_item<sizeof...(Is)>) const {}
   void operator()(item<sizeof...(Is)>) const {}
 
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::work_group_size<Is...>};
   }
diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp
@@ -44,7 +44,7 @@ void testQueriesAndProperties() {
               q, wgRange, wgRange.size() * sizeof(int));
   struct TestKernel0 {
     void operator()() const {}
-    auto get(sycl::ext::oneapi::experimental::properties_tag) {
+    auto get(sycl::ext::oneapi::experimental::properties_tag) const {
       return sycl::ext::oneapi::experimental::properties{
           sycl::ext::oneapi::experimental::use_root_sync};
     }
@@ -133,7 +133,7 @@ template <typename T> struct TestKernel2 {
           root.get_local_linear_range() == root.get_local_range().size();
     }
   }
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::use_root_sync};
   }
diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp
@@ -9,7 +9,9 @@ constexpr auto Props = sycl::ext::oneapi::experimental::properties{
 };
 struct TestKernelLaunchBounds {
   void operator()() const {}
-  auto get(sycl::ext::oneapi::experimental::properties_tag) { return Props; }
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return Props;
+  }
 };
 
 int main() {
diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp
@@ -13,17 +13,23 @@ constexpr auto Props3 = sycl::ext::oneapi::experimental::properties{
 
 struct TestKernel_Props1 {
   void operator()() const {}
-  auto get(sycl::ext::oneapi::experimental::properties_tag) { return Props1; }
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return Props1;
+  }
 };
 
 struct TestKernel_Props2 {
   void operator()() const {}
-  auto get(sycl::ext::oneapi::experimental::properties_tag) { return Props2; }
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return Props2;
+  }
 };
 
 struct TestKernel_Props3 {
   void operator()() const {}
-  auto get(sycl::ext::oneapi::experimental::properties_tag) { return Props3; }
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return Props3;
+  }
 };
 
 int main() {
diff --git a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp
@@ -76,38 +76,38 @@ template <typename T> struct K_funcIndirectlyUsingFP16 {
   T *Props;
   K_funcIndirectlyUsingFP16(T Props_param) { Props = &Props_param; };
   void operator()() const { int a = funcIndirectlyUsingFP16(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcIndirectlyUsingFP16_Warn16 {
   T *Props;
   K_funcIndirectlyUsingFP16_Warn16(T Props_param) { Props = &Props_param; };
   // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}}
   void operator()() const { int a = funcIndirectlyUsingFP16(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingFP16AndFP64 {
   T *Props;
   K_funcUsingFP16AndFP64(T Props_param) { Props = &Props_param; };
   void operator()() const { int a = funcUsingFP16AndFP64(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingFP16AndFP64_Warn16 {
   T *Props;
   K_funcUsingFP16AndFP64_Warn16(T Props_param) { Props = &Props_param; };
   // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}}
   void operator()() const { int a = funcUsingFP16AndFP64(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingFP16AndFP64_Warn64 {
   T *Props;
   K_funcUsingFP16AndFP64_Warn64(T Props_param) { Props = &Props_param; };
   // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}}
   void operator()() const { int a = funcUsingFP16AndFP64(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingFP16AndFP64_Warn1664 {
@@ -116,7 +116,7 @@ template <typename T> struct K_funcUsingFP16AndFP64_Warn1664 {
   // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}}
   // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}}
   void operator()() const { int a = funcUsingFP16AndFP64(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingFP16AndFP64_False {
@@ -127,52 +127,52 @@ template <typename T> struct K_funcUsingFP16AndFP64_False {
       int a = funcUsingFP16AndFP64(1, 2);
     }
   }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingCPUHasFP64 {
   T *Props;
   K_funcUsingCPUHasFP64(T Props_param) { Props = &Props_param; };
   void operator()() const { int a = funcUsingCPUHasFP64(1); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcIndirectlyUsingCPU {
   T *Props;
   K_funcIndirectlyUsingCPU(T Props_param) { Props = &Props_param; };
   void operator()() const { int a = funcIndirectlyUsingCPU(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcIndirectlyUsingCPU_WarnCPU {
   T *Props;
   K_funcIndirectlyUsingCPU_WarnCPU(T Props_param) { Props = &Props_param; };
   // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}}
   void operator()() const { int a = funcIndirectlyUsingCPU(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingCPUAndFP64 {
   T *Props;
   K_funcUsingCPUAndFP64(T Props_param) { Props = &Props_param; };
   void operator()() const { int a = funcUsingCPUAndFP64(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingCPUAndFP64_WarnCPU {
   T *Props;
   K_funcUsingCPUAndFP64_WarnCPU(T Props_param) { Props = &Props_param; };
   // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}}
   void operator()() const { int a = funcUsingCPUAndFP64(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingCPUAndFP64_Warn64 {
   T *Props;
   K_funcUsingCPUAndFP64_Warn64(T Props_param) { Props = &Props_param; };
   // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}}
   void operator()() const { int a = funcUsingCPUAndFP64(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingCPUAndFP64_Warn64CPU {
@@ -181,7 +181,7 @@ template <typename T> struct K_funcUsingCPUAndFP64_Warn64CPU {
   // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}}
   // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}}
   void operator()() const { int a = funcUsingCPUAndFP64(1, 2); }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 template <typename T> struct K_funcUsingCPUAndFP64_False {
@@ -192,7 +192,7 @@ template <typename T> struct K_funcUsingCPUAndFP64_False {
       int a = funcUsingCPUAndFP64(1, 2);
     }
   }
-  auto get(properties_tag) { return *Props; }
+  auto get(properties_tag) const { return *Props; }
 };
 
 int main() {
diff --git a/sycl/test/extensions/properties/properties_kernel_negative.cpp b/sycl/test/extensions/properties/properties_kernel_negative.cpp
@@ -4,23 +4,23 @@
 
 template <size_t... Is> struct KernelFunctorWithWGSize {
   void operator()() const {}
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::work_group_size<Is...>};
   }
 };
 
 template <size_t... Is> struct KernelFunctorWithWGSizeHint {
   void operator()() const {}
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::work_group_size_hint<Is...>};
   }
 };
 
 template <uint32_t I> struct KernelFunctorWithSGSize {
   void operator()() const {}
-  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
     return sycl::ext::oneapi::experimental::properties{
         sycl::ext::oneapi::experimental::sub_group_size<I>};
   }
@@ -373,11 +373,49 @@ void check_max_linear_work_group_size() {
       []() {});
 }
 
+struct TestKernelNonConstGetter {
+  TestKernelNonConstGetter() {}
+  void operator()() const { return; }
+  auto get(sycl::ext::oneapi::experimental::properties_tag) {
+    return sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::use_root_sync};
+  }
+};
+
+struct TestKernelConstGetter {
+  TestKernelConstGetter() {}
+  void operator()() const { return; }
+  auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+    return sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::use_root_sync};
+  }
+};
+
+struct TestKernelNoGetter {
+  TestKernelNoGetter() {}
+  void operator()() const { return; }
+};
+
+void check_non_const_getter_warning() {
+  sycl::queue Q;
+
+  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: get(sycl::ext::oneapi::experimental::properties_tag) member in kernel functor class must be declared as a const member function}}
+  Q.single_task(TestKernelNonConstGetter());
+
+  // No error expected for kernel functor with a const get(properties_tag)
+  // method
+  Q.single_task(TestKernelConstGetter());
+
+  // No error expected for kernel functor with no get(properties_tag) method
+  Q.single_task(TestKernelNoGetter());
+}
+
 int main() {
   check_work_group_size();
   check_work_group_size_hint();
   check_sub_group_size();
   check_max_work_group_size();
   check_max_linear_work_group_size();
+  check_non_const_getter_warning();
   return 0;
 }
diff --git a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp
diff --git a/sycl/test/virtual-functions/diagnostics-positive.cpp b/sycl/test/virtual-functions/diagnostics-positive.cpp
diff --git a/sycl/test/virtual-functions/properties-positive.cpp b/sycl/test/virtual-functions/properties-positive.cpp
diff --git a/unified-runtime/test/conformance/device_code/fixed_sg_size.cpp b/unified-runtime/test/conformance/device_code/fixed_sg_size.cpp
diff --git a/unified-runtime/test/conformance/device_code/fixed_wg_size.cpp b/unified-runtime/test/conformance/device_code/fixed_wg_size.cpp
diff --git a/unified-runtime/test/conformance/device_code/max_wg_size.cpp b/unified-runtime/test/conformance/device_code/max_wg_size.cpp
diff --git a/unified-runtime/test/conformance/device_code/subgroup.cpp b/unified-runtime/test/conformance/device_code/subgroup.cpp

Original file line number	Diff line number	Diff line change
`@@ -295,7 +295,7 @@ that it depends upon for correctness.`
`295`	`295`
`296`	`296`	`To enable this use-case, this extension adds a mechanism for implementations to`
`297`	`297`	`extract a property list from a kernel functor, if a kernel functor declares`
`298`		-a member function named `get` accepting a `sycl::ext::oneapi::experimental::properties_tag`
	`298`	+a const member function named `get` accepting a `sycl::ext::oneapi::experimental::properties_tag`
`299`	`299`	tag type and returning an instance of `sycl::ext::oneapi::experimental::properties`.
`300`	`300`
`301`	`301`	```c++
`@@ -338,7 +338,7 @@ struct KernelFunctor {`
`338`	`338`	`a[i] = b[i] + c[i];`
`339`	`339`	`}`
`340`	`340`
`341`		`- auto get(sycl::ext::oneapi::experimental::properties_tag) {`
	`341`	`+ auto get(sycl::ext::oneapi::experimental::properties_tag) const {`
`342`	`342`	`return sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::work_group_size<8, 8>,`
`343`	`343`	`sycl::ext::oneapi::experimental::sub_group_size<8>};`
`344`	`344`	`}`
Original file line number	Diff line number	Diff line change
`@@ -319,7 +319,7 @@ struct KernelFunctor {`
`319`	`319`	`a = b + *c;`
`320`	`320`	`}`
`321`	`321`
`322`		`- auto get(properties_tag) {`
	`322`	`+ auto get(properties_tag) const {`
`323`	`323`	`return properties{streaming_interface_accept_downstream_stall};`
`324`	`324`	`}`
`325`	`325`
Original file line number	Diff line number	Diff line change
`@@ -52,7 +52,7 @@ template <size_t I> struct KernelFunctorWithMaxWGSizeProp {`
`52`	`52`	`void operator()(nd_item<1>) const {}`
`53`	`53`	`void operator()(item<1>) const {}`
`54`	`54`
`55`		`- auto get(sycl::ext::oneapi::experimental::properties_tag) {`
	`55`	`+ auto get(sycl::ext::oneapi::experimental::properties_tag) const {`
`56`	`56`	`return sycl::ext::oneapi::experimental::properties{`
`57`	`57`	`sycl::ext::oneapi::experimental::max_linear_work_group_size<I>};`
`58`	`58`	`}`
Original file line number	Diff line number	Diff line change
`@@ -43,7 +43,7 @@ template <size_t... Is> struct KernelFunctorWithMaxWGSizeProp {`
`43`	`43`	`void operator()(nd_item<sizeof...(Is)>) const {}`
`44`	`44`	`void operator()(item<sizeof...(Is)>) const {}`
`45`	`45`
`46`		`- auto get(sycl::ext::oneapi::experimental::properties_tag) {`
	`46`	`+ auto get(sycl::ext::oneapi::experimental::properties_tag) const {`
`47`	`47`	`return sycl::ext::oneapi::experimental::properties{`
`48`	`48`	`sycl::ext::oneapi::experimental::max_work_group_size<Is...>};`
`49`	`49`	`}`
Original file line number	Diff line number	Diff line change
`@@ -24,7 +24,7 @@ template <size_t SGSize> struct KernelFunctorWithSGSizeProp {`
`24`	`24`	`Acc[0] = SG.get_local_linear_range();`
`25`	`25`	`}`
`26`	`26`
`27`		`- auto get(sycl::ext::oneapi::experimental::properties_tag) {`
	`27`	`+ auto get(sycl::ext::oneapi::experimental::properties_tag) const {`
`28`	`28`	`return sycl::ext::oneapi::experimental::properties{`
`29`	`29`	`sycl::ext::oneapi::experimental::sub_group_size<SGSize>};`
`30`	`30`	`}`
Original file line number	Diff line number	Diff line change
`@@ -39,7 +39,7 @@ template <size_t... Is> struct KernelFunctorWithWGSizeProp {`
`39`	`39`	`void operator()(nd_item<sizeof...(Is)>) const {}`
`40`	`40`	`void operator()(item<sizeof...(Is)>) const {}`
`41`	`41`
`42`		`- auto get(sycl::ext::oneapi::experimental::properties_tag) {`
	`42`	`+ auto get(sycl::ext::oneapi::experimental::properties_tag) const {`
`43`	`43`	`return sycl::ext::oneapi::experimental::properties{`
`44`	`44`	`sycl::ext::oneapi::experimental::work_group_size<Is...>};`
`45`	`45`	`}`
Original file line number	Diff line number	Diff line change
`@@ -19,7 +19,7 @@ template <size_t SGSize> struct KernelFunctorWithSGSizeProp {`
`19`	`19`	`Acc[0] = SG.get_local_linear_range();`
`20`	`20`	`}`
`21`	`21`
`22`		`- auto get(sycl::ext::oneapi::experimental::properties_tag) {`
	`22`	`+ auto get(sycl::ext::oneapi::experimental::properties_tag) const {`
`23`	`23`	`return sycl::ext::oneapi::experimental::properties{`
`24`	`24`	`sycl::ext::oneapi::experimental::sub_group_size<SGSize>};`
`25`	`25`	`}`