From 23ad9912b4b1fc8a1b36a0694750acd9d8b3056a Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Wed, 9 Apr 2025 21:32:25 -0700 Subject: [PATCH 1/2] [SYCL] Correct spec about constness of properties_tag getter and add respective warning Signed-off-by: Hu, Peisen --- sycl/doc/design/CompileTimeProperties.md | 4 +-- ...sycl_ext_oneapi_kernel_properties.asciidoc | 4 +-- ..._fpga_kernel_interface_properties.asciidoc | 2 +- sycl/include/sycl/handler.hpp | 8 +++++ sycl/include/syclcompat/launch_policy.hpp | 2 +- .../max_linear_work_group_size_props.cpp | 2 +- .../Basic/max_work_group_size_props.cpp | 2 +- sycl/test-e2e/Basic/sub_group_size_prop.cpp | 2 +- sycl/test-e2e/Basic/work_group_size_prop.cpp | 2 +- sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp | 2 +- .../Graph/Inputs/work_group_size_prop.cpp | 2 +- sycl/test-e2e/GroupAlgorithm/root_group.cpp | 4 +-- .../properties_kernel_launch_bounds.cpp | 4 ++- .../properties_kernel_max_work_group_size.cpp | 12 ++++++-- .../properties_kernel_device_has_warning.cpp | 30 +++++++++---------- .../properties/properties_kernel_negative.cpp | 6 ++-- .../properties_kernel_negative_device.cpp | 8 ++--- .../diagnostics-positive.cpp | 8 ++--- .../virtual-functions/properties-positive.cpp | 10 +++---- .../conformance/device_code/fixed_sg_size.cpp | 2 +- .../conformance/device_code/fixed_wg_size.cpp | 2 +- .../conformance/device_code/max_wg_size.cpp | 2 +- .../test/conformance/device_code/subgroup.cpp | 2 +- 23 files changed, 69 insertions(+), 53 deletions(-) diff --git a/sycl/doc/design/CompileTimeProperties.md b/sycl/doc/design/CompileTimeProperties.md index 02c50cd3b6f84..9c65356329ac9 100644 --- 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}; } }; diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc index 2ca9ac1b55d2a..341a1473d33b5 100644 --- 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 index e8c69f2b9ed66..371e87d5525ed 100644 --- 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 index 3466cf5362262..b2f35037ddcf4 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1594,6 +1594,14 @@ class __SYCL_EXPORT handler { const KernelType &>::value) { h->processProperties()>( KernelFunc.get(ext::oneapi::experimental::properties_tag{})); + } else { + // 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< + KernelType>::value), + "get(sycl::ext::oneapi::experimental::properties_tag) member in " + "kernel functor class must be declared as a const member function"); } #endif auto L = [&](auto &&...args) { diff --git a/sycl/include/syclcompat/launch_policy.hpp b/sycl/include/syclcompat/launch_policy.hpp index e5042620a8e0d..eca71e9ba4393 100644 --- 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) 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 index afe3ebd0d2557..d335b0c34b5a2 100644 --- 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 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}; } diff --git a/sycl/test-e2e/Basic/max_work_group_size_props.cpp b/sycl/test-e2e/Basic/max_work_group_size_props.cpp index 6694cb1d35d3f..a4d23db90863e 100644 --- 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 struct KernelFunctorWithMaxWGSizeProp { void operator()(nd_item) const {} void operator()(item) 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}; } diff --git a/sycl/test-e2e/Basic/sub_group_size_prop.cpp b/sycl/test-e2e/Basic/sub_group_size_prop.cpp index ae8281903a92b..dbff1c0c18ef5 100644 --- 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 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}; } diff --git a/sycl/test-e2e/Basic/work_group_size_prop.cpp b/sycl/test-e2e/Basic/work_group_size_prop.cpp index 9cf04c4d2ea66..08fa3e8aca972 100644 --- 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 struct KernelFunctorWithWGSizeProp { void operator()(nd_item) const {} void operator()(item) 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}; } diff --git a/sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp b/sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp index adaf6e1977ea4..a954e073cd87e 100644 --- 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 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}; } diff --git a/sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp b/sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp index 7fd3d8eef1856..f29e1a6455abd 100644 --- 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 struct KernelFunctorWithWGSizeProp { void operator()(nd_item) const {} void operator()(item) 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}; } diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp index aec138d2299b7..c0614c69d7991 100644 --- 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 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 index 75fbf5a18ab73..d242c8425dcb3 100644 --- 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 index e81bd5d20c452..8b3787be51f4c 100644 --- 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 index 55ac13958bce9..a3c095ae03975 100644 --- a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp +++ b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp @@ -76,7 +76,7 @@ template 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 struct K_funcIndirectlyUsingFP16_Warn16 { @@ -84,14 +84,14 @@ template struct K_funcIndirectlyUsingFP16_Warn16 { 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 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 struct K_funcUsingFP16AndFP64_Warn16 { @@ -99,7 +99,7 @@ template struct K_funcUsingFP16AndFP64_Warn16 { 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 struct K_funcUsingFP16AndFP64_Warn64 { @@ -107,7 +107,7 @@ template struct K_funcUsingFP16AndFP64_Warn64 { 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 struct K_funcUsingFP16AndFP64_Warn1664 { @@ -116,7 +116,7 @@ template 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 struct K_funcUsingFP16AndFP64_False { @@ -127,21 +127,21 @@ template struct K_funcUsingFP16AndFP64_False { int a = funcUsingFP16AndFP64(1, 2); } } - auto get(properties_tag) { return *Props; } + auto get(properties_tag) const { return *Props; } }; template 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 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 struct K_funcIndirectlyUsingCPU_WarnCPU { @@ -149,14 +149,14 @@ template struct K_funcIndirectlyUsingCPU_WarnCPU { 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 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 struct K_funcUsingCPUAndFP64_WarnCPU { @@ -164,7 +164,7 @@ template struct K_funcUsingCPUAndFP64_WarnCPU { 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 struct K_funcUsingCPUAndFP64_Warn64 { @@ -172,7 +172,7 @@ template struct K_funcUsingCPUAndFP64_Warn64 { 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 struct K_funcUsingCPUAndFP64_Warn64CPU { @@ -181,7 +181,7 @@ template 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 struct K_funcUsingCPUAndFP64_False { @@ -192,7 +192,7 @@ template 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 index 12c93cb312eed..bd3c95f48a366 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative.cpp @@ -4,7 +4,7 @@ template 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}; } @@ -12,7 +12,7 @@ template struct KernelFunctorWithWGSize { template 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}; } @@ -20,7 +20,7 @@ template struct KernelFunctorWithWGSizeHint { template 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}; } diff --git a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp index 230626e97e248..23bd5b31846b0 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp @@ -5,7 +5,7 @@ template struct KernelFunctorWithWGSizeWithAttr { // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::reqd_work_group_size(32)]] () 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}; } @@ -14,7 +14,7 @@ template struct KernelFunctorWithWGSizeWithAttr { template struct KernelFunctorWithWGSizeHintWithAttr { // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::work_group_size_hint(32)]] () 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}; } @@ -23,7 +23,7 @@ template struct KernelFunctorWithWGSizeHintWithAttr { template struct KernelFunctorWithSGSizeWithAttr { // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::reqd_sub_group_size(32)]] () 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}; } @@ -32,7 +32,7 @@ template struct KernelFunctorWithSGSizeWithAttr { template struct KernelFunctorWithDeviceHasWithAttr { // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::device_has(sycl::aspect::cpu)]] () 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::device_has}; } diff --git a/sycl/test/virtual-functions/diagnostics-positive.cpp b/sycl/test/virtual-functions/diagnostics-positive.cpp index 812161457784c..555dd9aeb48bb 100644 --- a/sycl/test/virtual-functions/diagnostics-positive.cpp +++ b/sycl/test/virtual-functions/diagnostics-positive.cpp @@ -66,7 +66,7 @@ struct TestKernel_props_empty { Ptr->foo(); } - auto get(oneapi::properties_tag) { return props_empty; } + auto get(oneapi::properties_tag) const { return props_empty; } }; struct TestKernel_props_void { @@ -78,7 +78,7 @@ struct TestKernel_props_void { Ptr->bar(); } - auto get(oneapi::properties_tag) { return props_void; } + auto get(oneapi::properties_tag) const { return props_void; } }; struct TestKernel_props_int { @@ -89,7 +89,7 @@ struct TestKernel_props_int { auto *Ptr = reinterpret_cast(Storage); foo(Ptr); } - auto get(oneapi::properties_tag) { return props_int; } + auto get(oneapi::properties_tag) const { return props_int; } }; struct TestKernel_props_base { @@ -99,7 +99,7 @@ struct TestKernel_props_base { auto *Ptr = reinterpret_cast(Storage); bar(Ptr); } - auto get(oneapi::properties_tag) { return props_base; } + auto get(oneapi::properties_tag) const { return props_base; } }; int main() { diff --git a/sycl/test/virtual-functions/properties-positive.cpp b/sycl/test/virtual-functions/properties-positive.cpp index 507a284e4de21..a660658118040 100644 --- a/sycl/test/virtual-functions/properties-positive.cpp +++ b/sycl/test/virtual-functions/properties-positive.cpp @@ -49,27 +49,27 @@ oneapi::properties props_multiple{oneapi::assume_indirect_calls_to}; struct TestKernel_props_empty { void operator()() const {} - auto get(oneapi::properties_tag) { return props_empty; } + auto get(oneapi::properties_tag) const { return props_empty; } }; struct TestKernel_props_void { void operator()() const {} - auto get(oneapi::properties_tag) { return props_void; } + auto get(oneapi::properties_tag) const { return props_void; } }; struct TestKernel_props_int { void operator()() const {} - auto get(oneapi::properties_tag) { return props_int; } + auto get(oneapi::properties_tag) const { return props_int; } }; struct TestKernel_props_base { void operator()() const {} - auto get(oneapi::properties_tag) { return props_base; } + auto get(oneapi::properties_tag) const { return props_base; } }; struct TestKernel_props_multiple { void operator()() const {} - auto get(oneapi::properties_tag) { return props_multiple; } + auto get(oneapi::properties_tag) const { return props_multiple; } }; int main() { diff --git a/unified-runtime/test/conformance/device_code/fixed_sg_size.cpp b/unified-runtime/test/conformance/device_code/fixed_sg_size.cpp index 3db4f54291c2d..9a65fd6fe4ebe 100644 --- a/unified-runtime/test/conformance/device_code/fixed_sg_size.cpp +++ b/unified-runtime/test/conformance/device_code/fixed_sg_size.cpp @@ -10,7 +10,7 @@ struct KernelFunctor { void operator()(sycl::nd_item<3>) const {} void operator()(sycl::item<3>) 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<8>}; } diff --git a/unified-runtime/test/conformance/device_code/fixed_wg_size.cpp b/unified-runtime/test/conformance/device_code/fixed_wg_size.cpp index 90a62d7cc9cd6..23f17eaa8a891 100644 --- a/unified-runtime/test/conformance/device_code/fixed_wg_size.cpp +++ b/unified-runtime/test/conformance/device_code/fixed_wg_size.cpp @@ -10,7 +10,7 @@ struct KernelFunctor { void operator()(sycl::nd_item<3>) const {} void operator()(sycl::item<3>) 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<8, 4, 2>}; } diff --git a/unified-runtime/test/conformance/device_code/max_wg_size.cpp b/unified-runtime/test/conformance/device_code/max_wg_size.cpp index 2e69d82778f52..fae5adb2223b1 100644 --- a/unified-runtime/test/conformance/device_code/max_wg_size.cpp +++ b/unified-runtime/test/conformance/device_code/max_wg_size.cpp @@ -10,7 +10,7 @@ struct KernelFunctor { void operator()(sycl::nd_item<3>) const {} void operator()(sycl::item<3>) 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<8, 4, 2>, sycl::ext::oneapi::experimental::max_linear_work_group_size<64>}; diff --git a/unified-runtime/test/conformance/device_code/subgroup.cpp b/unified-runtime/test/conformance/device_code/subgroup.cpp index ff668b03ff98d..3fed93641b9f0 100644 --- a/unified-runtime/test/conformance/device_code/subgroup.cpp +++ b/unified-runtime/test/conformance/device_code/subgroup.cpp @@ -12,7 +12,7 @@ struct KernelFunctor { KernelFunctor(sycl::accessor Acc) : Acc(Acc) {} - 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<8>}; } From 4ddecb467b2790cc44872183ba48b9b7ea3d2d5f Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Fri, 11 Apr 2025 13:53:45 -0700 Subject: [PATCH 2/2] [SYCL] Add test case for diagnostic message Signed-off-by: Hu, Peisen --- sycl/include/sycl/handler.hpp | 19 ++++++---- .../properties/properties_kernel_negative.cpp | 38 +++++++++++++++++++ 2 files changed, 49 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index b2f35037ddcf4..df3f7c43806d6 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1594,16 +1594,19 @@ class __SYCL_EXPORT handler { const KernelType &>::value) { h->processProperties()>( KernelFunc.get(ext::oneapi::experimental::properties_tag{})); - } else { - // 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< - KernelType>::value), - "get(sycl::ext::oneapi::experimental::properties_tag) member in " - "kernel functor class must be declared as a const member function"); } #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( diff --git a/sycl/test/extensions/properties/properties_kernel_negative.cpp b/sycl/test/extensions/properties/properties_kernel_negative.cpp index bd3c95f48a366..03f57711e0842 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative.cpp @@ -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; }