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 1 commit
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?

Copy link
Contributor

Choose a reason for hiding this comment

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

Adding @gmlueck and @rolandschulz per @Pennycook 's suggestion.

Also note that we need to modify

llvm/sycl/include/sycl/handler.hpp

Lines 339 to 340 in daca6a3

template <typename TransformedArgType, int Dims, typename KernelType>
class RoundedRangeKernel {
to propagate the properties, making this constexpr static would make that much easier.

Copy link
Contributor

Choose a reason for hiding this comment

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

I agree that kernel properties will always be compile-time. For a case where a runtime property is required, I think we would use a launch property instead. Therefore, it seems like it's OK to make this static constexpr.

Copy link
Contributor

Choose a reason for hiding this comment

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

I found the discussion @aelovikov-intel was referencing over in #17633, where we ended up adding both static constexpr and auto versions to deal with compile-time and run-time properties.

Committing to only compile-time properties would solve the problem here and simplify that extension, so I'm also happy to make this static constexpr.

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
8 changes: 8 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 @@ -1594,6 +1594,14 @@ class __SYCL_EXPORT handler {
const KernelType &>::value) {
h->processProperties<detail::isKernelESIMD<KernelName>()>(
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) {
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
6 changes: 3 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
8 changes: 4 additions & 4 deletions sycl/test/extensions/properties/properties_kernel_negative_device.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@
template <size_t... Is> 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<Is...>};
}
Expand All @@ -14,7 +14,7 @@ template <size_t... Is> struct KernelFunctorWithWGSizeWithAttr {
template <size_t... Is> 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<Is...>};
}
Expand All @@ -23,7 +23,7 @@ template <size_t... Is> struct KernelFunctorWithWGSizeHintWithAttr {
template <uint32_t I> 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<I>};
}
Expand All @@ -32,7 +32,7 @@ template <uint32_t I> struct KernelFunctorWithSGSizeWithAttr {
template <sycl::aspect Aspect> 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<Aspect>};
}
Expand Down
Loading