Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Apr 14, 2025
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
23ad991
[SYCL] Correct spec about constness of properties_tag getter and add …
HPS-1 Apr 10, 2025
4ddecb4
[SYCL] Add test case for diagnostic message
HPS-1 Apr 11, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions sycl/doc/design/CompileTimeProperties.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:

```
Expand All @@ -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>};
}
};
Expand Down
4 changes: 2 additions & 2 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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++
Expand Down Expand Up @@ -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 {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@Pennycook , I think there was some other PR where it was decided that kernel properties can't be runtime. Should we change this to static constexpr instead?

return sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::work_group_size<8, 8>,
sycl::ext::oneapi::experimental::sub_group_size<8>};
}
Expand Down
2 changes: 1 addition & 1 deletion ...oc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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};
}

Expand Down
11 changes: 11 additions & 0 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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...>(
Expand Down
2 changes: 1 addition & 1 deletion sycl/include/syclcompat/launch_policy.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 {
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/max_linear_work_group_size_props.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>};
}
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/max_work_group_size_props.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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...>};
}
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/sub_group_size_prop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>};
}
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Basic/work_group_size_prop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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...>};
}
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Graph/Inputs/sub_group_prop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>};
}
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Graph/Inputs/work_group_size_prop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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...>};
}
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/GroupAlgorithm/root_group.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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};
}
Expand Down Expand Up @@ -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};
}
Expand Down
4 changes: 3 additions & 1 deletion sycl/test/check_device_code/extensions/properties/properties_kernel_launch_bounds.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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() {
Expand Down
12 changes: 9 additions & 3 deletions sycl/test/check_device_code/extensions/properties/properties_kernel_max_work_group_size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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() {
Expand Down
30 changes: 15 additions & 15 deletions sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 {
Expand All @@ -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 {
Expand All @@ -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 {
Expand All @@ -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 {
Expand All @@ -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() {
Expand Down
44 changes: 41 additions & 3 deletions sycl/test/extensions/properties/properties_kernel_negative.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>};
}
Expand Down Expand Up @@ -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;
}
Loading
Loading