From 5125ae12e9194d6de810e4dfa9dea2eb91b20017 Mon Sep 17 00:00:00 2001 From: y Date: Fri, 10 Apr 2026 09:53:48 -0700 Subject: [PATCH 01/11] [SYCL] Add lightweight free-function kernel property headers Introduce lightweight kernel-property headers for free-function kernel annotation and launch-tuning use cases: - function_properties.hpp provides a standalone path for execution-kind annotations such as nd_range_kernel and single_task_kernel. - function_launch_properties.hpp provides a lightweight path for launch properties such as work_group_size, work_group_size_hint, and sub_group_size. - kernel_properties/properties.hpp remains the umbrella path for full property-list semantics and cross-property conflict checking. This reduces compile-time cost for kernel-language and JIT-style usage while preserving the full umbrella interface for users who need complete property-list behavior. Measured header compile times: - function_properties.hpp: ~26 ms - function_launch_properties path: ~58 ms - kernel_properties/properties.hpp umbrella: ~68 ms --- ..._ext_oneapi_free_function_kernels.asciidoc | 37 ++ ...sycl_ext_oneapi_kernel_properties.asciidoc | 34 ++ .../oneapi/experimental/virtual_functions.hpp | 23 ++ .../function_launch_properties.hpp | 329 ++++++++++++++++++ .../kernel_properties/function_properties.hpp | 98 ++++++ .../oneapi/kernel_properties/properties.hpp | 252 +------------- .../invalid_compile_time_properties.cpp | 6 +- .../properties/properties_kernel.cpp | 33 +- .../properties/properties_kernel_negative.cpp | 100 +++--- .../include_deps/sycl_detail_core.hpp.cpp | 6 +- .../sycl_khr_includes_handler.hpp.cpp | 2 + .../sycl_khr_includes_kernel_bundle.hpp.cpp | 2 + .../sycl_khr_includes_queue.hpp.cpp | 6 +- .../sycl_khr_includes_reduction.hpp.cpp | 2 + .../sycl_khr_includes_stream.hpp.cpp | 2 + .../sycl_khr_includes_usm.hpp.cpp | 6 +- 16 files changed, 630 insertions(+), 308 deletions(-) create mode 100644 sycl/include/sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp create mode 100644 sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc index 5ef1c663ee431..3359b3e0be1b4 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc @@ -112,6 +112,36 @@ supports. feature-test macro always has this value. |=== +=== Headers + +In the {dpcpp} implementation, applications that use this extension have a +choice of headers depending on which properties they need: + +* `#include ` is + sufficient for `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, + `nd_range_kernel`, and `single_task_kernel`. + +* `#include ` + additionally provides the launch-configuration properties that may be applied + directly to a free function kernel declaration, such as + `work_group_size`, `work_group_size_hint`, `sub_group_size`, + `max_work_group_size`, and `max_linear_work_group_size`. + +* `#include ` provides the + full kernel properties interface, including property lists such as + `properties{...}` and embedded kernel properties via `properties_tag`. + +Applications may include the umbrella header in all cases. The lighter-weight +headers are intended for code that only needs function annotations on free +function kernels and where reducing compile time is a priority. + +The lighter-weight headers retain validation of each individual property, such +as arity and non-zero size requirements, but they do not provide the +cross-property validation that is performed when properties are combined via +the full property-list machinery. For example, checks for conflicting +combinations such as `work_group_size` together with `max_work_group_size` are +available only through the umbrella header. + === Defining a free function kernel A free function kernel is a normal C++ function definition, where the function @@ -803,6 +833,13 @@ in link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[ sycl_ext_oneapi_kernel_properties] by applying the properties to the function declaration as illustrated below. +In the {dpcpp} implementation, code that uses only function annotations may +include either `function_properties.hpp` or +`function_launch_properties.hpp` instead of the full +`kernel_properties/properties.hpp` header. Applications that need property-list +semantics, including cross-property validation, should include the umbrella +header. + ``` SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size<32>)) 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 341a1473d33b5..a36020e31bba4 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -115,6 +115,40 @@ supports. feature-test macro always has this value. |=== +=== Headers + +In the {dpcpp} implementation, the full API described in this extension is +provided by: + +* `#include ` + +This umbrella header should be used when code needs property lists such as +`properties{...}`, embedded kernel properties via `properties_tag`, or any of +the non-launch kernel properties described in this extension. + +The implementation also provides lighter-weight headers for code that only +needs free-function annotations: + +* `#include ` for + `nd_range_kernel`, `single_task_kernel`, and + `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`. + +* `#include ` + for the launch-configuration properties that can be applied directly to free + function declarations, including `work_group_size`, + `work_group_size_hint`, `sub_group_size`, `max_work_group_size`, and + `max_linear_work_group_size`. + +These lighter-weight headers are primarily intended for code paths where +reduced compile time is a high priority and the application only needs +function-level annotations. + +They retain validation of each individual property, but they do not provide +the cross-property validation that depends on forming a full property list. +For example, checks for conflicting combinations such as `work_group_size` +together with `max_work_group_size` or `max_linear_work_group_size` are +performed only when using the umbrella header. + === Kernel Properties The kernel properties below correspond to kernel attributes defined in diff --git a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp index 9b13f6e3ed123..4ac20d852ec81 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp @@ -1,5 +1,6 @@ #pragma once +#include #include #include #include @@ -49,6 +50,17 @@ struct PropertyMetaInfo> { #endif }; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "indirectly-callable"; + static constexpr const char *value = +#ifdef __SYCL_DEVICE_ONLY__ + __builtin_sycl_unique_stable_name(Set); +#else + ""; +#endif +}; + #ifdef __SYCL_DEVICE_ONLY__ // Helper to concatenate several lists of characters into a single string. // Lists are separated from each other with comma within the resulting string. @@ -105,6 +117,17 @@ struct PropertyMetaInfo> { #endif }; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "calls-indirectly"; + static constexpr const char *value = +#ifdef __SYCL_DEVICE_ONLY__ + UniqueStableNameListStr::value; +#else + ""; +#endif +}; + } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp new file mode 100644 index 0000000000000..94fe9124145f0 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp @@ -0,0 +1,329 @@ +//==--- function_launch_properties.hpp - SYCL function launch properties --==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +// This header extends the split kernel-property path with launch-tuning +// annotations. Keep these public property_value definitions local to this split +// header so the umbrella include reuses the same types and preserves +// decltype(...) identity without pulling the full kernel-properties machinery +// into standalone launch-property use. + +#include +#include +#include + +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +template struct ConflictingProperties; + +template struct LaunchAllNonZero { + static constexpr bool value = true; +}; +template struct LaunchAllNonZero { + static constexpr bool value = X > 0 && LaunchAllNonZero::value; +}; + +inline constexpr size_t LaunchDecimalBase = 10; + +template struct LaunchSizeList {}; +template struct LaunchCharList {}; + +template struct LaunchCharsToStr { + static constexpr const char value[] = {Chars..., '\0'}; +}; + +template +struct LaunchSizeListToStrHelper; + +template +struct LaunchSizeListToStrHelper, + LaunchCharList, Chars...> + : LaunchSizeListToStrHelper< + LaunchSizeList, + LaunchCharList, '0' + (Value % LaunchDecimalBase), + Chars...> {}; + +template +struct LaunchSizeListToStrHelper, + LaunchCharList, Chars...> + : LaunchSizeListToStrHelper, + LaunchCharList> { +}; + +template +struct LaunchSizeListToStrHelper, + LaunchCharList> + : LaunchSizeListToStrHelper, + LaunchCharList> {}; + +template +struct LaunchSizeListToStrHelper, + LaunchCharList, Chars...> + : LaunchCharsToStr {}; + +template +struct LaunchSizeListToStrHelper, + LaunchCharList> + : LaunchCharsToStr {}; + +template <> +struct LaunchSizeListToStrHelper, LaunchCharList<>> + : LaunchCharsToStr<> {}; + +template +struct LaunchSizeListToStr + : LaunchSizeListToStrHelper, LaunchCharList<>> {}; + +} // namespace detail + +struct work_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct work_group_size_hint_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct sub_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value>; +}; + +struct max_work_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct max_linear_work_group_size_key + : detail::compile_time_property_key< + detail::PropKind::MaxLinearWorkGroupSize> { + template + using value_t = property_value>; +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::WorkGroupSize, work_group_size_key> { + static_assert( + sizeof...(Dims) + 1 <= 3, + "work_group_size property currently only supports up to three values."); + static_assert(detail::LaunchAllNonZero::value, + "work_group_size property must only contain non-zero values."); + + constexpr size_t operator[](int Dim) const { + return std::array{Dim0, Dims...}[Dim]; + } + +private: + constexpr size_t size() const { return sizeof...(Dims) + 1; } + + template friend struct detail::ConflictingProperties; +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::WorkGroupSizeHint, work_group_size_hint_key> { + static_assert(sizeof...(Dims) + 1 <= 3, + "work_group_size_hint property currently only supports up to " + "three values."); + static_assert( + detail::LaunchAllNonZero::value, + "work_group_size_hint property must only contain non-zero values."); + + constexpr size_t operator[](int Dim) const { + return std::array{Dim0, Dims...}[Dim]; + } +}; + +template +struct property_value> + : detail::property_base< + property_value>, + detail::PropKind::SubGroupSize, sub_group_size_key> { + static_assert(Size != 0, + "sub_group_size property must contain a non-zero value."); + + using value_t = std::integral_constant; + static constexpr uint32_t value = Size; +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::MaxWorkGroupSize, max_work_group_size_key> { + static_assert( + sizeof...(Dims) + 1 <= 3, + "max_work_group_size currently only supports up to three values."); + static_assert(detail::LaunchAllNonZero::value, + "max_work_group_size must only contain non-zero values."); + + constexpr size_t operator[](int Dim) const { + return std::array{Dim0, Dims...}[Dim]; + } + +private: + constexpr size_t size() const { return sizeof...(Dims) + 1; } + + template friend struct detail::ConflictingProperties; +}; + +template +struct property_value> + : detail::property_base< + property_value>, + detail::PropKind::MaxLinearWorkGroupSize, + max_linear_work_group_size_key> { + static_assert(Size != 0, + "max_linear_work_group_size must contain a non-zero value."); + + using value_t = std::integral_constant; + static constexpr size_t value = Size; +}; + +namespace detail { + +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size"; + static constexpr const char *value = + LaunchSizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size"; + static constexpr const char *value = + LaunchSizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size-hint"; + static constexpr const char *value = + LaunchSizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo< + work_group_size_hint_key::value_t> { + static constexpr const char *name = "sycl-work-group-size-hint"; + static constexpr const char *value = + LaunchSizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-sub-group-size"; + static constexpr uint32_t value = Size; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-sub-group-size"; + static constexpr uint32_t value = Size; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-max-work-group-size"; + static constexpr const char *value = + LaunchSizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo< + max_work_group_size_key::value_t> { + static constexpr const char *name = "sycl-max-work-group-size"; + static constexpr const char *value = + LaunchSizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-max-linear-work-group-size"; + static constexpr size_t value = Size; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-max-linear-work-group-size"; + static constexpr size_t value = Size; +}; + +} // namespace detail + +template +inline constexpr work_group_size_key::value_t work_group_size; + +template +inline constexpr work_group_size_hint_key::value_t + work_group_size_hint; + +template +inline constexpr sub_group_size_key::value_t sub_group_size; + +template +inline constexpr max_work_group_size_key::value_t + max_work_group_size; + +template +inline constexpr max_linear_work_group_size_key::value_t + max_linear_work_group_size; + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl \ No newline at end of file diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp new file mode 100644 index 0000000000000..1b8a133e05b9d --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp @@ -0,0 +1,98 @@ +//==--- function_properties.hpp - SYCL standalone function properties -----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +// This header is the lightweight split entry point for free-function kernel +// annotations. Keep the public property_value definitions here, rather than +// only in kernel_properties/properties.hpp, so standalone users and umbrella +// users observe the same decltype(...) while avoiding the heavier property-list +// machinery on this path. + +#include + +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +template +using remove_cvref_t = std::remove_cv_t>; + +template struct FunctionPropertyMetaInfo; + +} // namespace detail + +struct nd_range_kernel_key + : detail::compile_time_property_key { + template + using value_t = + property_value>; +}; + +struct single_task_kernel_key + : detail::compile_time_property_key { + using value_t = property_value; +}; + +template +struct property_value> + : detail::property_base>, + detail::PropKind::NDRangeKernel, + nd_range_kernel_key> { + static_assert(Dims >= 1 && Dims <= 3, + "nd_range_kernel must use dimension 1, 2, or 3."); + + using value_t = int; + static constexpr int dimensions = Dims; +}; + +template <> +struct property_value + : detail::property_base, + detail::PropKind::SingleTaskKernel, + single_task_kernel_key> {}; + +template +inline constexpr nd_range_kernel_key::value_t nd_range_kernel; + +inline constexpr single_task_kernel_key::value_t single_task_kernel; + +namespace detail { + +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-nd-range-kernel"; + static constexpr int value = Dims; +}; + +template <> struct FunctionPropertyMetaInfo { + static constexpr const char *name = "sycl-single-task-kernel"; + static constexpr int value = 0; +}; + +} // namespace detail +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl + +#ifdef __SYCL_DEVICE_ONLY__ +#define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) \ + [[__sycl_detail__::add_ir_attributes_function( \ + sycl::ext::oneapi::experimental::detail::FunctionPropertyMetaInfo< \ + sycl::ext::oneapi::experimental::detail::remove_cvref_t< \ + decltype(PROP)>>::name, \ + sycl::ext::oneapi::experimental::detail::FunctionPropertyMetaInfo< \ + sycl::ext::oneapi::experimental::detail::remove_cvref_t< \ + decltype(PROP)>>::value)]] +#else +#define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) +#endif \ No newline at end of file diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp index defb1589679b9..16cf6aed572b8 100644 --- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp @@ -13,45 +13,16 @@ #include // for uint32_t #include // for aspect #include // for forward_progress_guarantee enum +#include #include #include // for true_type #include // for declval namespace sycl { inline namespace _V1 { namespace ext::oneapi::experimental { -namespace detail { -// Trait for checking that all size_t values are non-zero. -template struct AllNonZero { - static constexpr bool value = true; -}; -template struct AllNonZero { - static constexpr bool value = X > 0 && AllNonZero::value; -}; -} // namespace detail struct properties_tag {}; -struct work_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct work_group_size_hint_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct sub_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value>; -}; - struct device_has_key : detail::compile_time_property_key { template @@ -59,92 +30,6 @@ struct device_has_key std::integral_constant...>; }; -struct nd_range_kernel_key - : detail::compile_time_property_key { - template - using value_t = - property_value>; -}; - -struct single_task_kernel_key - : detail::compile_time_property_key { - using value_t = property_value; -}; - -struct max_work_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct max_linear_work_group_size_key - : detail::compile_time_property_key< - detail::PropKind::MaxLinearWorkGroupSize> { - template - using value_t = property_value>; -}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::WorkGroupSize, work_group_size_key> { - static_assert( - sizeof...(Dims) + 1 <= 3, - "work_group_size property currently only supports up to three values."); - static_assert(detail::AllNonZero::value, - "work_group_size property must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - return std::array{Dim0, Dims...}[Dim]; - } - -private: - constexpr size_t size() const { return sizeof...(Dims) + 1; } - - template friend struct detail::ConflictingProperties; -}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::WorkGroupSizeHint, work_group_size_hint_key> { - static_assert(sizeof...(Dims) + 1 <= 3, - "work_group_size_hint property currently " - "only supports up to three values."); - static_assert( - detail::AllNonZero::value, - "work_group_size_hint property must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - return std::array{Dim0, Dims...}[Dim]; - } -}; - -template -struct property_value> - : detail::property_base< - property_value>, - detail::PropKind::SubGroupSize, sub_group_size_key> { - static_assert(Size != 0, - "sub_group_size_key property must contain a non-zero value."); - - using value_t = std::integral_constant; - static constexpr uint32_t value = Size; -}; - template struct property_value...> @@ -155,84 +40,9 @@ struct property_value value{Aspects...}; }; -template -struct property_value> - : detail::property_base>, - detail::PropKind::NDRangeKernel, - nd_range_kernel_key> { - static_assert( - Dims >= 1 && Dims <= 3, - "nd_range_kernel_key property must use dimension of 1, 2 or 3."); - - using value_t = int; - static constexpr int dimensions = Dims; -}; - -template <> -struct property_value - : detail::property_base, - detail::PropKind::SingleTaskKernel, - single_task_kernel_key> {}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::MaxWorkGroupSize, max_work_group_size_key> { - static_assert(sizeof...(Dims) + 1 <= 3, - "max_work_group_size property currently " - "only supports up to three values."); - static_assert( - detail::AllNonZero::value, - "max_work_group_size property must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - return std::array{Dim0, Dims...}[Dim]; - } - -private: - constexpr size_t size() const { return sizeof...(Dims) + 1; } - - template friend struct detail::ConflictingProperties; -}; - -template <> -struct property_value - : detail::property_base, - detail::PropKind::MaxLinearWorkGroupSize, - max_linear_work_group_size_key> {}; - -template -inline constexpr work_group_size_key::value_t work_group_size; - -template -inline constexpr work_group_size_hint_key::value_t - work_group_size_hint; - -template -inline constexpr sub_group_size_key::value_t sub_group_size; - template inline constexpr device_has_key::value_t device_has; -template -inline constexpr nd_range_kernel_key::value_t nd_range_kernel; - -inline constexpr single_task_kernel_key::value_t single_task_kernel; - -template -inline constexpr max_work_group_size_key::value_t - max_work_group_size; - -template -inline constexpr max_linear_work_group_size_key::value_t - max_linear_work_group_size; - struct work_group_progress_key : detail::compile_time_property_key { template work_item_progress; namespace detail { - -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; template struct HasCompileTimeEffect> : std::true_type {}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-work-group-size"; - static constexpr const char *value = SizeListToStr::value; -}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-work-group-size-hint"; - static constexpr const char *value = SizeListToStr::value; -}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-sub-group-size"; - static constexpr uint32_t value = Size; -}; template struct PropertyMetaInfo> { static constexpr const char *name = "sycl-device-has"; static constexpr const char *value = SizeListToStr(Aspects)...>::value; }; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-nd-range-kernel"; - static constexpr int value = Dims; -}; -template <> struct PropertyMetaInfo { - static constexpr const char *name = "sycl-single-task-kernel"; - static constexpr int value = 0; -}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-max-work-group-size"; - static constexpr const char *value = SizeListToStr::value; -}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-max-linear-work-group-size"; - static constexpr size_t value = Size; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-device-has"; + static constexpr const char *value = + SizeListToStr(Aspects)...>::value; }; template @@ -465,14 +236,3 @@ auto RetrieveGetMethodPropertiesOrEmpty(RestT &&...Rest) { } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl - -#ifdef __SYCL_DEVICE_ONLY__ -#define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) \ - [[__sycl_detail__::add_ir_attributes_function( \ - sycl::ext::oneapi::experimental::detail::PropertyMetaInfo< \ - std::remove_cv_t>>::name, \ - sycl::ext::oneapi::experimental::detail::PropertyMetaInfo< \ - std::remove_cv_t>>::value)]] -#else -#define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) -#endif diff --git a/sycl/test/extensions/annotated_usm/invalid_compile_time_properties.cpp b/sycl/test/extensions/annotated_usm/invalid_compile_time_properties.cpp index 1210ff0b602c9..d0f2b08845293 100644 --- a/sycl/test/extensions/annotated_usm/invalid_compile_time_properties.cpp +++ b/sycl/test/extensions/annotated_usm/invalid_compile_time_properties.cpp @@ -39,7 +39,7 @@ void testInvalidCompileTimeProperty(sycl::queue &q) { // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}device_has_key{{.+}}: Found invalid compile-time property in the property list.}} TEST(malloc_device_annotated, N, dev, Ctx, properties{alignment<4>, device_has}) - // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}work_group_size_key{{.+}}: Found invalid compile-time property in the property list.}} + // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}property_value<{{.+}}work_group_size_key{{.+}}: Found invalid compile-time property in the property list.}} TEST(aligned_alloc_device_annotated, 1, N, q, properties{unaliased, work_group_size<1>}) // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}baz_key{{.+}}: Found invalid compile-time property in the property list.}} @@ -87,7 +87,7 @@ void testInvalidCompileTimeProperty(sycl::queue &q) { // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}boo_key{{.+}}: Found invalid compile-time property in the property list.}} TEST(aligned_alloc_shared_annotated, 1, N, dev, Ctx, properties{unaliased, alignment<4>, boo}) - // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}sub_group_size_key{{.+}}: Found invalid compile-time property in the property list.}} + // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}property_value<{{.+}}sub_group_size_key{{.+}}: Found invalid compile-time property in the property list.}} TEST(malloc_annotated, N, q, alloc::device, properties{unaliased, sub_group_size<2>}) // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}boo_key{{.+}}: Found invalid compile-time property in the property list.}} @@ -96,7 +96,7 @@ void testInvalidCompileTimeProperty(sycl::queue &q) { // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}boo_key{{.+}}: Found invalid compile-time property in the property list.}} TEST(aligned_alloc_annotated, 1, N, dev, Ctx, alloc::device, properties{unaliased, boo}) - // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}work_group_size_hint_key{{.+}}: Found invalid compile-time property in the property list.}} + // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}property_value<{{.+}}work_group_size_hint_key{{.+}}: Found invalid compile-time property in the property list.}} TEST(aligned_alloc_annotated, 1, N, q, alloc::device, properties{unaliased, work_group_size_hint<1>}) // expected-error-re@sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp:* {{static assertion failed due to requirement {{.+}}baz_key{{.+}}: Found invalid compile-time property in the property list.}} diff --git a/sycl/test/extensions/properties/properties_kernel.cpp b/sycl/test/extensions/properties/properties_kernel.cpp index 86be8c94cf0d4..b5a6df22b28fc 100644 --- a/sycl/test/extensions/properties/properties_kernel.cpp +++ b/sycl/test/extensions/properties/properties_kernel.cpp @@ -60,6 +60,36 @@ int main() { is_property_value)>::value); static_assert( is_property_value)>::value); + static_assert( + std::is_same_v< + std::remove_cv_t)>, + property_value>>); + static_assert(std::is_same_v, + property_value>); + static_assert( + std::is_same_v< + std::remove_cv_t)>, + property_value, + std::integral_constant>>); + static_assert( + std::is_same_v)>, + property_value, + std::integral_constant, + std::integral_constant>>); + static_assert( + std::is_same_v)>, + property_value>>); + static_assert( + std::is_same_v)>, + property_value, + std::integral_constant>>); + static_assert( + std::is_same_v)>, + property_value>>); static_assert( std::is_same_v)::key_t>); @@ -108,9 +138,6 @@ int main() { static_assert(max_work_group_size<28, 29, 30>[2] == 30); static_assert(max_linear_work_group_size<28>.value == 28); - static_assert(std::is_same_v)::value_t, - std::integral_constant>); - singleAspectDeviceHasChecks(); singleAspectDeviceHasChecks(); singleAspectDeviceHasChecks(); diff --git a/sycl/test/extensions/properties/properties_kernel_negative.cpp b/sycl/test/extensions/properties/properties_kernel_negative.cpp index e3c7314182208..27dac354dca74 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative.cpp @@ -30,53 +30,53 @@ void check_work_group_size() { // expected-error@+1 {{too few template arguments for variable template 'work_group_size'}} auto WGSize0 = sycl::ext::oneapi::experimental::work_group_size<>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0>' requested here}} auto WGSize1 = sycl::ext::oneapi::experimental::work_group_size<0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 0>' requested here}} auto WGSize2 = sycl::ext::oneapi::experimental::work_group_size<0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 0>' requested here}} auto WGSize3 = sycl::ext::oneapi::experimental::work_group_size<1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 1>' requested here}} auto WGSize4 = sycl::ext::oneapi::experimental::work_group_size<0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<0, 0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 0, 0>' requested here}} auto WGSize5 = sycl::ext::oneapi::experimental::work_group_size<0, 0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<1, 0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 0, 0>' requested here}} auto WGSize6 = sycl::ext::oneapi::experimental::work_group_size<1, 0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<0, 1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 1, 0>' requested here}} auto WGSize7 = sycl::ext::oneapi::experimental::work_group_size<0, 1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<0, 0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 0, 1>' requested here}} auto WGSize8 = sycl::ext::oneapi::experimental::work_group_size<0, 0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<1, 1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 1, 0>' requested here}} auto WGSize9 = sycl::ext::oneapi::experimental::work_group_size<1, 1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<0, 1, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 1, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 1, 1>' requested here}} auto WGSize10 = sycl::ext::oneapi::experimental::work_group_size<0, 1, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<1, 0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 0, 1>' requested here}} auto WGSize11 = sycl::ext::oneapi::experimental::work_group_size<1, 0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property currently only supports up to three values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size<1, 1, 1, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property currently only supports up to three values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 1, 1, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 1, 1, 1>' requested here}} auto WGSize12 = sycl::ext::oneapi::experimental::work_group_size<1, 1, 1, 1>; @@ -164,55 +164,55 @@ void check_work_group_size_hint() { // expected-error@+1 {{too few template arguments for variable template 'work_group_size_hint'}} auto WGSize0 = sycl::ext::oneapi::experimental::work_group_size_hint<>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0>' requested here}} auto WGSize1 = sycl::ext::oneapi::experimental::work_group_size_hint<0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 0>' requested here}} auto WGSize2 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 0>' requested here}} auto WGSize3 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 1>' requested here}} auto WGSize4 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 0, 0>' requested here}} auto WGSize5 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 0, 0>' requested here}} auto WGSize6 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 1, 0>' requested here}} auto WGSize7 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 0, 1>' requested here}} auto WGSize8 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 1, 0>' requested here}} auto WGSize9 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1, 1>' must be initialized by a constant expression}} // expected-note@+2 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 1, 1>' requested here}} auto WGSize10 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 1, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0, 1>' must be initialized by a constant expression}} // expected-note@+2 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 0, 1>' requested here}} auto WGSize11 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property currently only supports up to three values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 1, 1, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property currently only supports up to three values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 1, 1, 1>' must be initialized by a constant expression}} // expected-note@+2 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 1, 1, 1>' requested here}} auto WGSize12 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 1, 1, 1>; @@ -301,8 +301,8 @@ void check_sub_group_size() { // expected-error@+1 {{too few template arguments for variable template 'sub_group_size'}} auto WGSize0 = sycl::ext::oneapi::experimental::sub_group_size<>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: sub_group_size_key property must contain a non-zero value.}} - // expected-error@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{constexpr variable 'sub_group_size<0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: sub_group_size property must contain a non-zero value.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'sub_group_size<0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::sub_group_size<0>' requested here}} auto WGSize1 = sycl::ext::oneapi::experimental::sub_group_size<0>; diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index 2b8ac26da9ddc..1aa2adc6b740b 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -113,10 +113,12 @@ // CHECK-NEXT: device_selector.hpp // CHECK-NEXT: kernel_bundle_enums.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp +// CHECK-NEXT: ext/oneapi/properties/property_value.hpp +// CHECK-NEXT: ext/oneapi/properties/property.hpp // CHECK-NEXT: ext/oneapi/properties/properties.hpp // CHECK-NEXT: ext/oneapi/properties/property_utils.hpp -// CHECK-NEXT: ext/oneapi/properties/property.hpp -// CHECK-NEXT: ext/oneapi/properties/property_value.hpp // CHECK-NEXT: event.hpp // CHECK-NEXT: exception_list.hpp // CHECK-NEXT: ext/oneapi/device_global/device_global.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp index a4c53ca4f9439..7194a1acf4c80 100644 --- a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp @@ -135,7 +135,9 @@ // CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/nd_range_view.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp index b82f38f1cfdd6..edabb605d9ff3 100644 --- a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp @@ -143,7 +143,9 @@ // CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/nd_range_view.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp index 8ddf2a9576ed0..ea7d66f1d650c 100644 --- a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp @@ -117,10 +117,12 @@ // CHECK-NEXT: device_selector.hpp // CHECK-NEXT: kernel_bundle_enums.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp +// CHECK-NEXT: ext/oneapi/properties/property_value.hpp +// CHECK-NEXT: ext/oneapi/properties/property.hpp // CHECK-NEXT: ext/oneapi/properties/properties.hpp // CHECK-NEXT: ext/oneapi/properties/property_utils.hpp -// CHECK-NEXT: ext/oneapi/properties/property.hpp -// CHECK-NEXT: ext/oneapi/properties/property_value.hpp // CHECK-NEXT: event.hpp // CHECK-NEXT: exception_list.hpp // CHECK-NEXT: ext/oneapi/device_global/device_global.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp index 6775a70dbd859..90bd4c21c51f6 100644 --- a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp @@ -172,7 +172,9 @@ // CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/nd_range_view.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp index 8b69fc29fe202..abfcf488bf61e 100644 --- a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp @@ -153,7 +153,9 @@ // CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/nd_range_view.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp index 10d253de26117..a646ff9e4bd3a 100644 --- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp @@ -132,10 +132,12 @@ // CHECK-NEXT: detail/optional.hpp // CHECK-NEXT: detail/range_rounding.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp +// CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp +// CHECK-NEXT: ext/oneapi/properties/property_value.hpp +// CHECK-NEXT: ext/oneapi/properties/property.hpp // CHECK-NEXT: ext/oneapi/properties/properties.hpp // CHECK-NEXT: ext/oneapi/properties/property_utils.hpp -// CHECK-NEXT: ext/oneapi/properties/property.hpp -// CHECK-NEXT: ext/oneapi/properties/property_value.hpp // CHECK-NEXT: event.hpp // CHECK-NEXT: exception_list.hpp // CHECK-NEXT: ext/oneapi/device_global/device_global.hpp From d7405eb5c21f16f8c0dac7e094be319ff296ec98 Mon Sep 17 00:00:00 2001 From: y Date: Tue, 14 Apr 2026 10:21:49 -0700 Subject: [PATCH 02/11] Fold lightweight free-function properties into function_properties.hpp --- ..._ext_oneapi_free_function_kernels.asciidoc | 39 +-- ...sycl_ext_oneapi_kernel_properties.asciidoc | 23 +- .../function_launch_properties.hpp | 329 ------------------ .../kernel_properties/function_properties.hpp | 315 ++++++++++++++++- .../oneapi/kernel_properties/properties.hpp | 2 +- .../properties/function_properties_split.cpp | 37 ++ .../properties/properties_kernel_negative.cpp | 100 +++--- .../include_deps/sycl_detail_core.hpp.cpp | 1 - .../sycl_khr_includes_handler.hpp.cpp | 1 - .../sycl_khr_includes_kernel_bundle.hpp.cpp | 1 - .../sycl_khr_includes_queue.hpp.cpp | 1 - .../sycl_khr_includes_reduction.hpp.cpp | 1 - .../sycl_khr_includes_stream.hpp.cpp | 1 - .../sycl_khr_includes_usm.hpp.cpp | 1 - 14 files changed, 424 insertions(+), 428 deletions(-) delete mode 100644 sycl/include/sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp create mode 100644 sycl/test/extensions/properties/function_properties_split.cpp diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc index 3359b3e0be1b4..d33b59a859229 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc @@ -119,28 +119,23 @@ choice of headers depending on which properties they need: * `#include ` is sufficient for `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, - `nd_range_kernel`, and `single_task_kernel`. - -* `#include ` - additionally provides the launch-configuration properties that may be applied - directly to a free function kernel declaration, such as - `work_group_size`, `work_group_size_hint`, `sub_group_size`, - `max_work_group_size`, and `max_linear_work_group_size`. + `nd_range_kernel`, `single_task_kernel`, and the standalone + compile-time kernel properties that may be applied directly to a free + function kernel declaration, such as `work_group_size`, + `work_group_size_hint`, `sub_group_size`, `max_work_group_size`, and + `max_linear_work_group_size`. * `#include ` provides the full kernel properties interface, including property lists such as `properties{...}` and embedded kernel properties via `properties_tag`. -Applications may include the umbrella header in all cases. The lighter-weight -headers are intended for code that only needs function annotations on free -function kernels and where reducing compile time is a priority. +Applications may include the umbrella header in all cases. The lightweight +header is intended for code that defines free-function kernels and only needs +the free-function annotation path. -The lighter-weight headers retain validation of each individual property, such -as arity and non-zero size requirements, but they do not provide the -cross-property validation that is performed when properties are combined via -the full property-list machinery. For example, checks for conflicting -combinations such as `work_group_size` together with `max_work_group_size` are -available only through the umbrella header. +Code that forms property lists such as `properties{...}`, uses +`properties_tag`, or depends on the umbrella property's property-list conflict +checks should include `kernel_properties/properties.hpp`. === Defining a free function kernel @@ -833,12 +828,12 @@ in link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[ sycl_ext_oneapi_kernel_properties] by applying the properties to the function declaration as illustrated below. -In the {dpcpp} implementation, code that uses only function annotations may -include either `function_properties.hpp` or -`function_launch_properties.hpp` instead of the full -`kernel_properties/properties.hpp` header. Applications that need property-list -semantics, including cross-property validation, should include the umbrella -header. +In the {dpcpp} implementation, code that defines free-function kernels may +include `function_properties.hpp` instead of the full +`kernel_properties/properties.hpp` header when it only needs the +free-function annotation path. Applications that use property-list semantics, +including `properties{...}`, `properties_tag`, or the associated +property-list conflict checks, should include the umbrella header. ``` SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) 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 a36020e31bba4..52fc1ea1d07a6 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -130,24 +130,19 @@ The implementation also provides lighter-weight headers for code that only needs free-function annotations: * `#include ` for - `nd_range_kernel`, `single_task_kernel`, and - `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`. - -* `#include ` - for the launch-configuration properties that can be applied directly to free - function declarations, including `work_group_size`, - `work_group_size_hint`, `sub_group_size`, `max_work_group_size`, and - `max_linear_work_group_size`. + `nd_range_kernel`, `single_task_kernel`, + `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, and the standalone compile-time kernel + properties that may be applied directly to a free function declaration, such + as `work_group_size`, `work_group_size_hint`, `sub_group_size`, + `max_work_group_size`, and `max_linear_work_group_size`. These lighter-weight headers are primarily intended for code paths where reduced compile time is a high priority and the application only needs function-level annotations. - -They retain validation of each individual property, but they do not provide -the cross-property validation that depends on forming a full property list. -For example, checks for conflicting combinations such as `work_group_size` -together with `max_work_group_size` or `max_linear_work_group_size` are -performed only when using the umbrella header. +The `function_properties.hpp` header is intended for the free-function kernel +definition path. Code that forms property lists, uses `properties_tag`, or +depends on the umbrella property's property-list conflict checks should include +`kernel_properties/properties.hpp`. === Kernel Properties diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp deleted file mode 100644 index 94fe9124145f0..0000000000000 --- a/sycl/include/sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp +++ /dev/null @@ -1,329 +0,0 @@ -//==--- function_launch_properties.hpp - SYCL function launch properties --==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#pragma once - -// This header extends the split kernel-property path with launch-tuning -// annotations. Keep these public property_value definitions local to this split -// header so the umbrella include reuses the same types and preserves -// decltype(...) identity without pulling the full kernel-properties machinery -// into standalone launch-property use. - -#include -#include -#include - -#include -#include - -namespace sycl { -inline namespace _V1 { -namespace ext::oneapi::experimental { -namespace detail { - -template struct ConflictingProperties; - -template struct LaunchAllNonZero { - static constexpr bool value = true; -}; -template struct LaunchAllNonZero { - static constexpr bool value = X > 0 && LaunchAllNonZero::value; -}; - -inline constexpr size_t LaunchDecimalBase = 10; - -template struct LaunchSizeList {}; -template struct LaunchCharList {}; - -template struct LaunchCharsToStr { - static constexpr const char value[] = {Chars..., '\0'}; -}; - -template -struct LaunchSizeListToStrHelper; - -template -struct LaunchSizeListToStrHelper, - LaunchCharList, Chars...> - : LaunchSizeListToStrHelper< - LaunchSizeList, - LaunchCharList, '0' + (Value % LaunchDecimalBase), - Chars...> {}; - -template -struct LaunchSizeListToStrHelper, - LaunchCharList, Chars...> - : LaunchSizeListToStrHelper, - LaunchCharList> { -}; - -template -struct LaunchSizeListToStrHelper, - LaunchCharList> - : LaunchSizeListToStrHelper, - LaunchCharList> {}; - -template -struct LaunchSizeListToStrHelper, - LaunchCharList, Chars...> - : LaunchCharsToStr {}; - -template -struct LaunchSizeListToStrHelper, - LaunchCharList> - : LaunchCharsToStr {}; - -template <> -struct LaunchSizeListToStrHelper, LaunchCharList<>> - : LaunchCharsToStr<> {}; - -template -struct LaunchSizeListToStr - : LaunchSizeListToStrHelper, LaunchCharList<>> {}; - -} // namespace detail - -struct work_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct work_group_size_hint_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct sub_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value>; -}; - -struct max_work_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct max_linear_work_group_size_key - : detail::compile_time_property_key< - detail::PropKind::MaxLinearWorkGroupSize> { - template - using value_t = property_value>; -}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::WorkGroupSize, work_group_size_key> { - static_assert( - sizeof...(Dims) + 1 <= 3, - "work_group_size property currently only supports up to three values."); - static_assert(detail::LaunchAllNonZero::value, - "work_group_size property must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - return std::array{Dim0, Dims...}[Dim]; - } - -private: - constexpr size_t size() const { return sizeof...(Dims) + 1; } - - template friend struct detail::ConflictingProperties; -}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::WorkGroupSizeHint, work_group_size_hint_key> { - static_assert(sizeof...(Dims) + 1 <= 3, - "work_group_size_hint property currently only supports up to " - "three values."); - static_assert( - detail::LaunchAllNonZero::value, - "work_group_size_hint property must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - return std::array{Dim0, Dims...}[Dim]; - } -}; - -template -struct property_value> - : detail::property_base< - property_value>, - detail::PropKind::SubGroupSize, sub_group_size_key> { - static_assert(Size != 0, - "sub_group_size property must contain a non-zero value."); - - using value_t = std::integral_constant; - static constexpr uint32_t value = Size; -}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::MaxWorkGroupSize, max_work_group_size_key> { - static_assert( - sizeof...(Dims) + 1 <= 3, - "max_work_group_size currently only supports up to three values."); - static_assert(detail::LaunchAllNonZero::value, - "max_work_group_size must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - return std::array{Dim0, Dims...}[Dim]; - } - -private: - constexpr size_t size() const { return sizeof...(Dims) + 1; } - - template friend struct detail::ConflictingProperties; -}; - -template -struct property_value> - : detail::property_base< - property_value>, - detail::PropKind::MaxLinearWorkGroupSize, - max_linear_work_group_size_key> { - static_assert(Size != 0, - "max_linear_work_group_size must contain a non-zero value."); - - using value_t = std::integral_constant; - static constexpr size_t value = Size; -}; - -namespace detail { - -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-work-group-size"; - static constexpr const char *value = - LaunchSizeListToStr::value; -}; -template -struct FunctionPropertyMetaInfo> { - static constexpr const char *name = "sycl-work-group-size"; - static constexpr const char *value = - LaunchSizeListToStr::value; -}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-work-group-size-hint"; - static constexpr const char *value = - LaunchSizeListToStr::value; -}; -template -struct FunctionPropertyMetaInfo< - work_group_size_hint_key::value_t> { - static constexpr const char *name = "sycl-work-group-size-hint"; - static constexpr const char *value = - LaunchSizeListToStr::value; -}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-sub-group-size"; - static constexpr uint32_t value = Size; -}; -template -struct FunctionPropertyMetaInfo> { - static constexpr const char *name = "sycl-sub-group-size"; - static constexpr uint32_t value = Size; -}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-max-work-group-size"; - static constexpr const char *value = - LaunchSizeListToStr::value; -}; -template -struct FunctionPropertyMetaInfo< - max_work_group_size_key::value_t> { - static constexpr const char *name = "sycl-max-work-group-size"; - static constexpr const char *value = - LaunchSizeListToStr::value; -}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-max-linear-work-group-size"; - static constexpr size_t value = Size; -}; -template -struct FunctionPropertyMetaInfo> { - static constexpr const char *name = "sycl-max-linear-work-group-size"; - static constexpr size_t value = Size; -}; - -} // namespace detail - -template -inline constexpr work_group_size_key::value_t work_group_size; - -template -inline constexpr work_group_size_hint_key::value_t - work_group_size_hint; - -template -inline constexpr sub_group_size_key::value_t sub_group_size; - -template -inline constexpr max_work_group_size_key::value_t - max_work_group_size; - -template -inline constexpr max_linear_work_group_size_key::value_t - max_linear_work_group_size; - -} // namespace ext::oneapi::experimental -} // namespace _V1 -} // namespace sycl \ No newline at end of file diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp index 1b8a133e05b9d..dda6371583db6 100644 --- a/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp @@ -9,11 +9,14 @@ #pragma once // This header is the lightweight split entry point for free-function kernel -// annotations. Keep the public property_value definitions here, rather than -// only in kernel_properties/properties.hpp, so standalone users and umbrella -// users observe the same decltype(...) while avoiding the heavier property-list -// machinery on this path. +// annotations and standalone compile-time kernel properties that do not need +// the umbrella property's property-list machinery. Keep these public +// property_value definitions here so standalone users and umbrella users +// observe the same decltype(...) while avoiding the heavier machinery on this +// path. +#include +#include #include #include @@ -27,6 +30,73 @@ template using remove_cvref_t = std::remove_cv_t>; template struct FunctionPropertyMetaInfo; +template struct ConflictingProperties; + +template struct FunctionPropertyAllNonZero { + static constexpr bool value = true; +}; +template +struct FunctionPropertyAllNonZero { + static constexpr bool value = X > 0 && FunctionPropertyAllNonZero::value; +}; + +inline constexpr size_t FunctionPropertyDecimalBase = 10; + +template struct FunctionPropertySizeList {}; +template struct FunctionPropertyCharList {}; + +template struct FunctionPropertyCharsToStr { + static constexpr const char value[] = {Chars..., '\0'}; +}; + +template +struct FunctionPropertySizeListToStrHelper; + +template +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList, + FunctionPropertyCharList, Chars...> + : FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList, + FunctionPropertyCharList, + '0' + (Value % FunctionPropertyDecimalBase), Chars...> {}; + +template +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList<0, Values...>, + FunctionPropertyCharList, Chars...> + : FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList, + FunctionPropertyCharList> {}; + +template +struct FunctionPropertySizeListToStrHelper, + FunctionPropertyCharList> + : FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList, + FunctionPropertyCharList> {}; + +template +struct FunctionPropertySizeListToStrHelper, + FunctionPropertyCharList, + Chars...> + : FunctionPropertyCharsToStr {}; + +template +struct FunctionPropertySizeListToStrHelper, + FunctionPropertyCharList> + : FunctionPropertyCharsToStr {}; + +template <> +struct FunctionPropertySizeListToStrHelper, + FunctionPropertyCharList<>> + : FunctionPropertyCharsToStr<> {}; + +template +struct FunctionPropertySizeListToStr + : FunctionPropertySizeListToStrHelper, + FunctionPropertyCharList<>> {}; } // namespace detail @@ -66,6 +136,162 @@ inline constexpr nd_range_kernel_key::value_t nd_range_kernel; inline constexpr single_task_kernel_key::value_t single_task_kernel; +struct work_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct work_group_size_hint_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct sub_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value>; +}; + +struct max_work_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct max_linear_work_group_size_key + : detail::compile_time_property_key< + detail::PropKind::MaxLinearWorkGroupSize> { + template + using value_t = property_value>; +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::WorkGroupSize, work_group_size_key> { + static_assert( + sizeof...(Dims) + 1 <= 3, + "work_group_size property currently only supports up to three values."); + static_assert(detail::FunctionPropertyAllNonZero::value, + "work_group_size property must only contain non-zero values."); + + constexpr size_t operator[](int Dim) const { + constexpr size_t Values[] = {Dim0, Dims...}; + return Values[Dim]; + } + +private: + constexpr size_t size() const { return sizeof...(Dims) + 1; } + + template friend struct detail::ConflictingProperties; +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::WorkGroupSizeHint, work_group_size_hint_key> { + static_assert(sizeof...(Dims) + 1 <= 3, + "work_group_size_hint property currently only supports up to " + "three values."); + static_assert( + detail::FunctionPropertyAllNonZero::value, + "work_group_size_hint property must only contain non-zero values."); + + constexpr size_t operator[](int Dim) const { + constexpr size_t Values[] = {Dim0, Dims...}; + return Values[Dim]; + } +}; + +template +struct property_value> + : detail::property_base< + property_value>, + detail::PropKind::SubGroupSize, sub_group_size_key> { + static_assert(Size != 0, + "sub_group_size property must contain a non-zero value."); + + using value_t = std::integral_constant; + static constexpr uint32_t value = Size; +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::MaxWorkGroupSize, max_work_group_size_key> { + static_assert( + sizeof...(Dims) + 1 <= 3, + "max_work_group_size currently only supports up to three values."); + static_assert(detail::FunctionPropertyAllNonZero::value, + "max_work_group_size must only contain non-zero values."); + + constexpr size_t operator[](int Dim) const { + constexpr size_t Values[] = {Dim0, Dims...}; + return Values[Dim]; + } + +private: + constexpr size_t size() const { return sizeof...(Dims) + 1; } + + template friend struct detail::ConflictingProperties; +}; + +template +struct property_value> + : detail::property_base< + property_value>, + detail::PropKind::MaxLinearWorkGroupSize, + max_linear_work_group_size_key> { + static_assert(Size != 0, + "max_linear_work_group_size must contain a non-zero value."); + + using value_t = std::integral_constant; + static constexpr size_t value = Size; +}; + +template +inline constexpr work_group_size_key::value_t work_group_size; + +template +inline constexpr work_group_size_hint_key::value_t + work_group_size_hint; + +template +inline constexpr sub_group_size_key::value_t sub_group_size; + +template +inline constexpr max_work_group_size_key::value_t + max_work_group_size; + +template +inline constexpr max_linear_work_group_size_key::value_t + max_linear_work_group_size; + namespace detail { template @@ -79,6 +305,85 @@ template <> struct FunctionPropertyMetaInfo { static constexpr int value = 0; }; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size-hint"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo< + work_group_size_hint_key::value_t> { + static constexpr const char *name = "sycl-work-group-size-hint"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-sub-group-size"; + static constexpr uint32_t value = Size; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-sub-group-size"; + static constexpr uint32_t value = Size; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-max-work-group-size"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo< + max_work_group_size_key::value_t> { + static constexpr const char *name = "sycl-max-work-group-size"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-max-linear-work-group-size"; + static constexpr size_t value = Size; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-max-linear-work-group-size"; + static constexpr size_t value = Size; +}; + } // namespace detail } // namespace ext::oneapi::experimental } // namespace _V1 @@ -95,4 +400,4 @@ template <> struct FunctionPropertyMetaInfo { decltype(PROP)>>::value)]] #else #define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) -#endif \ No newline at end of file +#endif diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp index 16cf6aed572b8..b26d68a4d3be1 100644 --- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp @@ -13,7 +13,7 @@ #include // for uint32_t #include // for aspect #include // for forward_progress_guarantee enum -#include +#include #include #include // for true_type #include // for declval diff --git a/sycl/test/extensions/properties/function_properties_split.cpp b/sycl/test/extensions/properties/function_properties_split.cpp new file mode 100644 index 0000000000000..c72d27b1c115d --- /dev/null +++ b/sycl/test/extensions/properties/function_properties_split.cpp @@ -0,0 +1,37 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -I %sycl_source_dir/include -I %sycl_include %s +// expected-no-diagnostics + +#include + +#include + +using namespace sycl::ext::oneapi::experimental; + +int main() { + static_assert(is_property_value)>::value); + static_assert(is_property_value::value); + static_assert(is_property_value)>::value); + static_assert(is_property_value)>::value); + static_assert(is_property_value)>::value); + static_assert(is_property_value)>::value); + static_assert(is_property_value)>::value); + + static_assert(work_group_size<2, 3>[0] == 2); + static_assert(work_group_size<2, 3>[1] == 3); + static_assert(work_group_size_hint<4, 5, 6>[2] == 6); + static_assert(sub_group_size<7>.value == 7); + static_assert(max_work_group_size<8, 9>[1] == 9); + static_assert(max_linear_work_group_size<10>.value == 10); + + static_assert(std::is_same_v)::key_t, + work_group_size_key>); + static_assert(std::is_same_v)::key_t, + work_group_size_hint_key>); + static_assert(std::is_same_v)::key_t, + sub_group_size_key>); + static_assert(std::is_same_v)::key_t, + max_work_group_size_key>); + static_assert(std::is_same_v)::key_t, + max_linear_work_group_size_key>); + return 0; +} diff --git a/sycl/test/extensions/properties/properties_kernel_negative.cpp b/sycl/test/extensions/properties/properties_kernel_negative.cpp index 27dac354dca74..3ee88369079fa 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative.cpp @@ -30,53 +30,53 @@ void check_work_group_size() { // expected-error@+1 {{too few template arguments for variable template 'work_group_size'}} auto WGSize0 = sycl::ext::oneapi::experimental::work_group_size<>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0>' requested here}} auto WGSize1 = sycl::ext::oneapi::experimental::work_group_size<0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 0>' requested here}} auto WGSize2 = sycl::ext::oneapi::experimental::work_group_size<0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 0>' requested here}} auto WGSize3 = sycl::ext::oneapi::experimental::work_group_size<1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 1>' requested here}} auto WGSize4 = sycl::ext::oneapi::experimental::work_group_size<0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<0, 0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 0, 0>' requested here}} auto WGSize5 = sycl::ext::oneapi::experimental::work_group_size<0, 0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<1, 0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 0, 0>' requested here}} auto WGSize6 = sycl::ext::oneapi::experimental::work_group_size<1, 0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<0, 1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 1, 0>' requested here}} auto WGSize7 = sycl::ext::oneapi::experimental::work_group_size<0, 1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<0, 0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 0, 1>' requested here}} auto WGSize8 = sycl::ext::oneapi::experimental::work_group_size<0, 0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<1, 1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 1, 0>' requested here}} auto WGSize9 = sycl::ext::oneapi::experimental::work_group_size<1, 1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<0, 1, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<0, 1, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<0, 1, 1>' requested here}} auto WGSize10 = sycl::ext::oneapi::experimental::work_group_size<0, 1, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<1, 0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 0, 1>' requested here}} auto WGSize11 = sycl::ext::oneapi::experimental::work_group_size<1, 0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property currently only supports up to three values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size<1, 1, 1, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size property currently only supports up to three values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size<1, 1, 1, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size<1, 1, 1, 1>' requested here}} auto WGSize12 = sycl::ext::oneapi::experimental::work_group_size<1, 1, 1, 1>; @@ -164,55 +164,55 @@ void check_work_group_size_hint() { // expected-error@+1 {{too few template arguments for variable template 'work_group_size_hint'}} auto WGSize0 = sycl::ext::oneapi::experimental::work_group_size_hint<>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0>' requested here}} auto WGSize1 = sycl::ext::oneapi::experimental::work_group_size_hint<0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 0>' requested here}} auto WGSize2 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 0>' requested here}} auto WGSize3 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 1>' requested here}} auto WGSize4 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 0, 0>' requested here}} auto WGSize5 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 0, 0>' requested here}} auto WGSize6 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 0, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 1, 0>' requested here}} auto WGSize7 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 0, 1>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 0, 1>' requested here}} auto WGSize8 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 1, 0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 1, 0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 1, 0>' requested here}} auto WGSize9 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 1, 0>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<0, 1, 1>' must be initialized by a constant expression}} // expected-note@+2 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<0, 1, 1>' requested here}} auto WGSize10 = sycl::ext::oneapi::experimental::work_group_size_hint<0, 1, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property must only contain non-zero values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 0, 1>' must be initialized by a constant expression}} // expected-note@+2 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 0, 1>' requested here}} auto WGSize11 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 0, 1>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property currently only supports up to three values.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 1, 1, 1>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size_hint property currently only supports up to three values.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'work_group_size_hint<1, 1, 1, 1>' must be initialized by a constant expression}} // expected-note@+2 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::work_group_size_hint<1, 1, 1, 1>' requested here}} auto WGSize12 = sycl::ext::oneapi::experimental::work_group_size_hint<1, 1, 1, 1>; @@ -301,8 +301,8 @@ void check_sub_group_size() { // expected-error@+1 {{too few template arguments for variable template 'sub_group_size'}} auto WGSize0 = sycl::ext::oneapi::experimental::sub_group_size<>; - // expected-error-re@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{static assertion failed due to requirement {{.+}}: sub_group_size property must contain a non-zero value.}} - // expected-error@sycl/ext/oneapi/kernel_properties/function_launch_properties.hpp:* {{constexpr variable 'sub_group_size<0>' must be initialized by a constant expression}} + // expected-error-re@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{static assertion failed due to requirement {{.+}}: sub_group_size property must contain a non-zero value.}} + // expected-error@sycl/ext/oneapi/kernel_properties/function_properties.hpp:* {{constexpr variable 'sub_group_size<0>' must be initialized by a constant expression}} // expected-note@+1 {{in instantiation of variable template specialization 'sycl::ext::oneapi::experimental::sub_group_size<0>' requested here}} auto WGSize1 = sycl::ext::oneapi::experimental::sub_group_size<0>; diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp index 1aa2adc6b740b..72d93643164a2 100644 --- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp +++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp @@ -113,7 +113,6 @@ // CHECK-NEXT: device_selector.hpp // CHECK-NEXT: kernel_bundle_enums.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp -// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/properties/property_value.hpp // CHECK-NEXT: ext/oneapi/properties/property.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp index 7194a1acf4c80..7f19980446c7d 100644 --- a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp @@ -137,7 +137,6 @@ // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp -// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/nd_range_view.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp index edabb605d9ff3..4e7588112ff91 100644 --- a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp @@ -145,7 +145,6 @@ // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp -// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/nd_range_view.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp index ea7d66f1d650c..954e8906be1de 100644 --- a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp @@ -117,7 +117,6 @@ // CHECK-NEXT: device_selector.hpp // CHECK-NEXT: kernel_bundle_enums.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp -// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/properties/property_value.hpp // CHECK-NEXT: ext/oneapi/properties/property.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp index 90bd4c21c51f6..a3db1523e08d5 100644 --- a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp @@ -174,7 +174,6 @@ // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp -// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/nd_range_view.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp index abfcf488bf61e..166c8c70a4299 100644 --- a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp @@ -155,7 +155,6 @@ // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp -// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp // CHECK-NEXT: detail/nd_range_view.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp index a646ff9e4bd3a..1478ac6895dc8 100644 --- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp @@ -132,7 +132,6 @@ // CHECK-NEXT: detail/optional.hpp // CHECK-NEXT: detail/range_rounding.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp -// CHECK-NEXT: ext/oneapi/kernel_properties/function_launch_properties.hpp // CHECK-NEXT: ext/oneapi/kernel_properties/function_properties.hpp // CHECK-NEXT: ext/oneapi/properties/property_value.hpp // CHECK-NEXT: ext/oneapi/properties/property.hpp From c17bca0468d4a915bfeb51ec873bf1794cd350a7 Mon Sep 17 00:00:00 2001 From: y Date: Tue, 14 Apr 2026 12:02:32 -0700 Subject: [PATCH 03/11] fix formatting --- .../kernel_properties/function_properties.hpp | 22 +++++++++---------- .../properties/function_properties_split.cpp | 10 +++++---- 2 files changed, 17 insertions(+), 15 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp index dda6371583db6..bd2d7b8c41172 100644 --- a/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp @@ -35,9 +35,9 @@ template struct ConflictingProperties; template struct FunctionPropertyAllNonZero { static constexpr bool value = true; }; -template -struct FunctionPropertyAllNonZero { - static constexpr bool value = X > 0 && FunctionPropertyAllNonZero::value; +template struct FunctionPropertyAllNonZero { + static constexpr bool value = + X > 0 && FunctionPropertyAllNonZero::value; }; inline constexpr size_t FunctionPropertyDecimalBase = 10; @@ -71,21 +71,21 @@ struct FunctionPropertySizeListToStrHelper< FunctionPropertyCharList> {}; template -struct FunctionPropertySizeListToStrHelper, - FunctionPropertyCharList> +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList<0, Values...>, + FunctionPropertyCharList> : FunctionPropertySizeListToStrHelper< FunctionPropertySizeList, FunctionPropertyCharList> {}; template -struct FunctionPropertySizeListToStrHelper, - FunctionPropertyCharList, - Chars...> - : FunctionPropertyCharsToStr {}; +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList<0>, FunctionPropertyCharList, + Chars...> : FunctionPropertyCharsToStr {}; template -struct FunctionPropertySizeListToStrHelper, - FunctionPropertyCharList> +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList<0>, FunctionPropertyCharList> : FunctionPropertyCharsToStr {}; template <> diff --git a/sycl/test/extensions/properties/function_properties_split.cpp b/sycl/test/extensions/properties/function_properties_split.cpp index c72d27b1c115d..e9891185d0725 100644 --- a/sycl/test/extensions/properties/function_properties_split.cpp +++ b/sycl/test/extensions/properties/function_properties_split.cpp @@ -11,10 +11,12 @@ int main() { static_assert(is_property_value)>::value); static_assert(is_property_value::value); static_assert(is_property_value)>::value); - static_assert(is_property_value)>::value); + static_assert( + is_property_value)>::value); static_assert(is_property_value)>::value); static_assert(is_property_value)>::value); - static_assert(is_property_value)>::value); + static_assert( + is_property_value)>::value); static_assert(work_group_size<2, 3>[0] == 2); static_assert(work_group_size<2, 3>[1] == 3); @@ -27,8 +29,8 @@ int main() { work_group_size_key>); static_assert(std::is_same_v)::key_t, work_group_size_hint_key>); - static_assert(std::is_same_v)::key_t, - sub_group_size_key>); + static_assert( + std::is_same_v)::key_t, sub_group_size_key>); static_assert(std::is_same_v)::key_t, max_work_group_size_key>); static_assert(std::is_same_v)::key_t, From faa95e1849cc93b7bee1695b960fe1b8714aa380 Mon Sep 17 00:00:00 2001 From: y Date: Thu, 23 Apr 2026 09:21:13 -0700 Subject: [PATCH 04/11] Update the user-facing documentation to point to the top-level kernel property headers instead of the nested kernel_properties/* paths. Keep the legacy kernel_properties/* headers as forwarding compatibility shims, add smoke tests for both legacy include paths, and mark the shims for deletion once support for the old paths is dropped. --- ..._ext_oneapi_free_function_kernels.asciidoc | 10 +- ...sycl_ext_oneapi_kernel_properties.asciidoc | 12 +- .../sycl/detail/kernel_launch_helper.hpp | 49 +- sycl/include/sycl/detail/range_rounding.hpp | 2 +- .../ext/intel/esimd/memory_properties.hpp | 2 +- .../experimental/cache_control_properties.hpp | 2 +- .../experimental/grf_size_properties.hpp | 2 +- .../oneapi/device_global/device_global.hpp | 2 +- .../annotated_ptr/annotated_ptr.hpp | 2 +- .../annotated_ptr_properties.hpp | 2 +- .../experimental/cluster_group_prop.hpp | 2 +- .../properties.hpp | 2 +- .../oneapi/experimental/enqueue_functions.hpp | 2 +- .../ext/oneapi/experimental/graph/dynamic.hpp | 2 +- .../experimental/group_helpers_sorters.hpp | 2 +- .../oneapi/experimental/group_load_store.hpp | 2 +- .../sycl/ext/oneapi/experimental/prefetch.hpp | 2 +- .../experimental/syclbin_kernel_bundle.hpp | 2 +- .../experimental/syclbin_properties.hpp | 2 +- .../experimental/use_root_sync_prop.hpp | 2 +- .../oneapi/experimental/virtual_functions.hpp | 2 +- .../oneapi/experimental/work_group_memory.hpp | 2 +- .../free_function_kernel_properties.hpp | 389 ++++++++++++++++ .../sycl/ext/oneapi/kernel_properties.hpp | 184 ++++++++ .../kernel_properties/function_properties.hpp | 398 +--------------- .../oneapi/kernel_properties/properties.hpp | 233 +--------- sycl/include/sycl/ext/oneapi/properties.hpp | 421 +++++++++++++++++ .../sycl/ext/oneapi/properties/properties.hpp | 428 +----------------- .../ext/oneapi/work_group_scratch_memory.hpp | 2 +- sycl/include/sycl/handler.hpp | 4 +- sycl/include/sycl/kernel_bundle.hpp | 2 +- sycl/include/sycl/queue.hpp | 2 +- sycl/include/sycl/sycl.hpp | 4 +- .../free_function_kernels.hpp | 2 +- .../KernelCompiler/auto_pch_compile_error.cpp | 2 +- .../test-e2e/KernelCompiler/auto_pch_time.cpp | 2 +- .../conflicting_auto_pch_opts.cpp | 2 +- .../KernelCompiler/empty_preamble.cpp | 2 +- .../KernelCompiler/multi_threaded_rtc.cpp | 2 +- .../multiple_auto_pch_includes.cpp | 2 +- .../persistent_auto_pch_cache_collision.cpp | 2 +- .../persistent_auto_pch_read_error.cpp | 2 +- .../persistent_auto_pch_stress_deletion.cpp | 2 +- .../preamble_define_before_include.cpp | 2 +- .../KernelCompiler/preamble_if_stack.cpp | 4 +- .../KernelCompiler/auto-pch.cpp | 2 +- .../annotated_usm/invalid_usm_kind.cpp | 2 +- .../properties/final_header_include_both.cpp | 18 + .../properties/function_properties_split.cpp | 2 +- .../kernel_function_properties_compat.cpp | 20 + .../properties/kernel_properties_compat.cpp | 20 + .../properties/properties_ctor_negative.cpp | 12 +- .../properties_kernel_cache_config.cpp | 2 +- .../properties/properties_kernel_negative.cpp | 160 +++---- .../include_deps/sycl_detail_core.hpp.cpp | 6 +- .../sycl_khr_includes_handler.hpp.cpp | 6 +- .../sycl_khr_includes_interop_handle.hpp.cpp | 2 +- .../sycl_khr_includes_kernel_bundle.hpp.cpp | 6 +- .../sycl_khr_includes_queue.hpp.cpp | 6 +- .../sycl_khr_includes_reduction.hpp.cpp | 6 +- .../sycl_khr_includes_stream.hpp.cpp | 6 +- .../sycl_khr_includes_usm.hpp.cpp | 6 +- 62 files changed, 1272 insertions(+), 1212 deletions(-) create mode 100644 sycl/include/sycl/ext/oneapi/free_function_kernel_properties.hpp create mode 100644 sycl/include/sycl/ext/oneapi/kernel_properties.hpp create mode 100644 sycl/include/sycl/ext/oneapi/properties.hpp create mode 100644 sycl/test/extensions/properties/final_header_include_both.cpp create mode 100644 sycl/test/extensions/properties/kernel_function_properties_compat.cpp create mode 100644 sycl/test/extensions/properties/kernel_properties_compat.cpp diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc index d33b59a859229..4c972bec21738 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_kernels.asciidoc @@ -117,7 +117,7 @@ supports. In the {dpcpp} implementation, applications that use this extension have a choice of headers depending on which properties they need: -* `#include ` is +* `#include ` is sufficient for `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, `nd_range_kernel`, `single_task_kernel`, and the standalone compile-time kernel properties that may be applied directly to a free @@ -125,7 +125,7 @@ choice of headers depending on which properties they need: `work_group_size_hint`, `sub_group_size`, `max_work_group_size`, and `max_linear_work_group_size`. -* `#include ` provides the +* `#include ` provides the full kernel properties interface, including property lists such as `properties{...}` and embedded kernel properties via `properties_tag`. @@ -135,7 +135,7 @@ the free-function annotation path. Code that forms property lists such as `properties{...}`, uses `properties_tag`, or depends on the umbrella property's property-list conflict -checks should include `kernel_properties/properties.hpp`. +checks should include `sycl/ext/oneapi/kernel_properties.hpp`. === Defining a free function kernel @@ -829,8 +829,8 @@ sycl_ext_oneapi_kernel_properties] by applying the properties to the function declaration as illustrated below. In the {dpcpp} implementation, code that defines free-function kernels may -include `function_properties.hpp` instead of the full -`kernel_properties/properties.hpp` header when it only needs the +include `sycl/ext/oneapi/free_function_kernel_properties.hpp` instead of the +full `sycl/ext/oneapi/kernel_properties.hpp` header when it only needs the free-function annotation path. Applications that use property-list semantics, including `properties{...}`, `properties_tag`, or the associated property-list conflict checks, should include the umbrella header. 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 52fc1ea1d07a6..4b40180c4180e 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -120,7 +120,7 @@ supports. In the {dpcpp} implementation, the full API described in this extension is provided by: -* `#include ` +* `#include ` This umbrella header should be used when code needs property lists such as `properties{...}`, embedded kernel properties via `properties_tag`, or any of @@ -129,7 +129,7 @@ the non-launch kernel properties described in this extension. The implementation also provides lighter-weight headers for code that only needs free-function annotations: -* `#include ` for +* `#include ` for `nd_range_kernel`, `single_task_kernel`, `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, and the standalone compile-time kernel properties that may be applied directly to a free function declaration, such @@ -139,10 +139,10 @@ needs free-function annotations: These lighter-weight headers are primarily intended for code paths where reduced compile time is a high priority and the application only needs function-level annotations. -The `function_properties.hpp` header is intended for the free-function kernel -definition path. Code that forms property lists, uses `properties_tag`, or -depends on the umbrella property's property-list conflict checks should include -`kernel_properties/properties.hpp`. +The `sycl/ext/oneapi/free_function_kernel_properties.hpp` header is intended +for the free-function kernel definition path. Code that forms property lists, +uses `properties_tag`, or depends on the umbrella property's property-list +conflict checks should include `sycl/ext/oneapi/kernel_properties.hpp`. === Kernel Properties diff --git a/sycl/include/sycl/detail/kernel_launch_helper.hpp b/sycl/include/sycl/detail/kernel_launch_helper.hpp index 909f6f2f49ae8..3aff73e58130d 100644 --- a/sycl/include/sycl/detail/kernel_launch_helper.hpp +++ b/sycl/include/sycl/detail/kernel_launch_helper.hpp @@ -19,7 +19,7 @@ #include #include #include -#include +#include #include #include @@ -361,6 +361,52 @@ using KernelPropertyHolderStructTy = sycl::ext::oneapi::experimental::cuda::cluster_size_key<2>, sycl::ext::oneapi::experimental::cuda::cluster_size_key<3>>; +template +constexpr void validateKernelProperties() { + using namespace sycl::ext::oneapi::experimental; + + if constexpr (PropertiesT::template has_property()) { + constexpr auto WGSize = + PropertiesT::template get_property(); + + if constexpr (PropertiesT::template has_property()) { + constexpr auto MaxWGSize = + PropertiesT::template get_property(); + constexpr auto WGDimensions = decltype(WGSize)::dimensions; + constexpr auto MaxWGDimensions = decltype(MaxWGSize)::dimensions; + + static_assert( + WGDimensions == MaxWGDimensions, + "work_group_size and max_work_group_size dimensionality must match"); + if constexpr (WGDimensions == MaxWGDimensions) { + static_assert( + WGDimensions < 1 || WGSize[0] <= MaxWGSize[0], + "work_group_size must not exceed max_work_group_size"); + static_assert( + WGDimensions < 2 || WGSize[1] <= MaxWGSize[1], + "work_group_size must not exceed max_work_group_size"); + static_assert( + WGDimensions < 3 || WGSize[2] <= MaxWGSize[2], + "work_group_size must not exceed max_work_group_size"); + } + } + + if constexpr (PropertiesT::template has_property< + max_linear_work_group_size_key>()) { + constexpr auto Dimensions = decltype(WGSize)::dimensions; + constexpr auto LinearSize = + WGSize[0] * (Dimensions > 1 ? WGSize[1] : 1) * + (Dimensions > 2 ? WGSize[2] : 1); + constexpr auto MaxLinearWGSize = + PropertiesT::template get_property(); + + static_assert( + LinearSize < MaxLinearWGSize.value, + "work_group_size must not exceed max_linear_work_group_size"); + } + } +} + /// Note: it is important that this function *does not* depend on kernel /// name or kernel type, because then it will be instantiated for every /// kernel, even though body of those instantiated functions could be almost @@ -370,6 +416,7 @@ template >> constexpr KernelPropertyHolderStructTy extractKernelProperties(PropertiesT Props) { + validateKernelProperties(); static_assert( !PropertiesT::template has_property< sycl::ext::intel::experimental::fp_control_key>() || diff --git a/sycl/include/sycl/detail/range_rounding.hpp b/sycl/include/sycl/detail/range_rounding.hpp index 37dfe4f603205..7dcc381b75fdb 100644 --- a/sycl/include/sycl/detail/range_rounding.hpp +++ b/sycl/include/sycl/detail/range_rounding.hpp @@ -13,7 +13,7 @@ #include #include #include -#include +#include #include #include #include diff --git a/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp b/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp index f54a46b6fa922..9e87b226519ed 100644 --- a/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp +++ b/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp @@ -10,7 +10,7 @@ #include #include -#include +#include #include #include diff --git a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp index ea8ec1020d54a..7f2e815afc4b7 100644 --- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp @@ -10,7 +10,7 @@ #include #include -#include +#include #include #include diff --git a/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp b/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp index f0edcb0c88706..ca5d2bddcb334 100644 --- a/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp @@ -8,7 +8,7 @@ #pragma once -#include +#include #include #include diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp index ecf8c970e2214..39b1a0921a678 100644 --- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp +++ b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp @@ -11,7 +11,7 @@ #include // for address_space #include // for make_error_code #include // for device_image... -#include // for properties_t +#include // for properties_t #include // for multi_ptr #include // for decorated_gl... diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp index bcec3d2aeced0..f009197c15cc4 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp @@ -13,7 +13,7 @@ #include #include #include -#include +#include #include #include #include diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr_properties.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr_properties.hpp index 45a3fe63b6fda..5308d5c5a92d3 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr_properties.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr_properties.hpp @@ -10,7 +10,7 @@ #pragma once #include -#include // for properties_t +#include // for properties_t #include #include // for false_type, con... diff --git a/sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp b/sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp index 9e0d84afb660f..8bfe2ac56df15 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp @@ -8,7 +8,7 @@ #pragma once -#include +#include #include namespace sycl { diff --git a/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp index 6b138ecd72669..cc3b3c981bfb0 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp @@ -8,7 +8,7 @@ #pragma once -#include // for properties_t +#include // for properties_t #include // for false_type, con... #include // for declval diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index b364e94090360..4bcf9eb5647a4 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -15,7 +15,7 @@ #include #include #include -#include +#include #include #include #include diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph/dynamic.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph/dynamic.hpp index 9bfc71e378ec0..52eb0370c4c06 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph/dynamic.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph/dynamic.hpp @@ -14,7 +14,7 @@ #include // for __SYCL_EXPORT #include // for kernel_param_kind_t #include // for work_group_memory -#include // for empty_properties_t +#include // for empty_properties_t #include // for function #include // for shared_ptr diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp index e6941dd4d19c0..2199663635a55 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp @@ -14,7 +14,7 @@ #include // for min #include // for sycl_category, exception #include // for bfloat16 -#include +#include #include // for memory_scope #include // for range #include // for span diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp index 4bba7b980b5c1..b4cfe057b2a8d 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp @@ -13,7 +13,7 @@ #include #include #include -#include +#include #include #include diff --git a/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp index 8218a744eb1d4..8a6342970934f 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp @@ -10,7 +10,7 @@ #include #include -#include +#include #include #include diff --git a/sycl/include/sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp b/sycl/include/sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp index 50f3c9e0841f1..f5f72d3c0da4b 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/syclbin_kernel_bundle.hpp @@ -9,7 +9,7 @@ #pragma once #include -#include +#include #include #include diff --git a/sycl/include/sycl/ext/oneapi/experimental/syclbin_properties.hpp b/sycl/include/sycl/ext/oneapi/experimental/syclbin_properties.hpp index 3b74faeff9c8e..c6b1e78839df4 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/syclbin_properties.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/syclbin_properties.hpp @@ -8,7 +8,7 @@ #pragma once -#include +#include #include namespace sycl { diff --git a/sycl/include/sycl/ext/oneapi/experimental/use_root_sync_prop.hpp b/sycl/include/sycl/ext/oneapi/experimental/use_root_sync_prop.hpp index 3fbb4b9586d15..76424f2efb8a5 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/use_root_sync_prop.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/use_root_sync_prop.hpp @@ -11,7 +11,7 @@ #pragma once -#include +#include namespace sycl { inline namespace _V1 { diff --git a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp index 4ac20d852ec81..7fd51920f020b 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp @@ -1,6 +1,6 @@ #pragma once -#include +#include #include #include #include diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 099e2c92a2c4f..60089555575db 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include diff --git a/sycl/include/sycl/ext/oneapi/free_function_kernel_properties.hpp b/sycl/include/sycl/ext/oneapi/free_function_kernel_properties.hpp new file mode 100644 index 0000000000000..348478ac61272 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/free_function_kernel_properties.hpp @@ -0,0 +1,389 @@ +//==--- free_function_kernel_properties.hpp - SYCL free-function kernels --==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include + +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +template +using remove_cvref_t = std::remove_cv_t>; + +template struct FunctionPropertyMetaInfo; + +template struct FunctionPropertyAllNonZero { + static constexpr bool value = true; +}; +template struct FunctionPropertyAllNonZero { + static constexpr bool value = + X > 0 && FunctionPropertyAllNonZero::value; +}; + +inline constexpr size_t FunctionPropertyDecimalBase = 10; + +template struct FunctionPropertySizeList {}; +template struct FunctionPropertyCharList {}; + +template struct FunctionPropertyCharsToStr { + static constexpr const char value[] = {Chars..., '\0'}; +}; + +template +struct FunctionPropertySizeListToStrHelper; + +template +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList, + FunctionPropertyCharList, Chars...> + : FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList, + FunctionPropertyCharList, + '0' + (Value % FunctionPropertyDecimalBase), Chars...> {}; + +template +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList<0, Values...>, + FunctionPropertyCharList, Chars...> + : FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList, + FunctionPropertyCharList> {}; + +template +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList<0, Values...>, + FunctionPropertyCharList> + : FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList, + FunctionPropertyCharList> {}; + +template +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList<0>, FunctionPropertyCharList, + Chars...> : FunctionPropertyCharsToStr {}; + +template +struct FunctionPropertySizeListToStrHelper< + FunctionPropertySizeList<0>, FunctionPropertyCharList> + : FunctionPropertyCharsToStr {}; + +template <> +struct FunctionPropertySizeListToStrHelper, + FunctionPropertyCharList<>> + : FunctionPropertyCharsToStr<> {}; + +template +struct FunctionPropertySizeListToStr + : FunctionPropertySizeListToStrHelper, + FunctionPropertyCharList<>> {}; + +} // namespace detail + +struct nd_range_kernel_key + : detail::compile_time_property_key { + template + using value_t = + property_value>; +}; + +struct single_task_kernel_key + : detail::compile_time_property_key { + using value_t = property_value; +}; + +template +struct property_value> + : detail::property_base>, + detail::PropKind::NDRangeKernel, + nd_range_kernel_key> { + static_assert(Dims >= 1 && Dims <= 3, + "nd_range_kernel must use dimension 1, 2, or 3."); + + using value_t = int; + static constexpr int dimensions = Dims; +}; + +template <> +struct property_value + : detail::property_base, + detail::PropKind::SingleTaskKernel, + single_task_kernel_key> {}; + +template +inline constexpr nd_range_kernel_key::value_t nd_range_kernel; + +inline constexpr single_task_kernel_key::value_t single_task_kernel; + +struct work_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct work_group_size_hint_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct sub_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value>; +}; + +struct max_work_group_size_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +struct max_linear_work_group_size_key + : detail::compile_time_property_key< + detail::PropKind::MaxLinearWorkGroupSize> { + template + using value_t = property_value>; +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::WorkGroupSize, work_group_size_key> { + static_assert( + sizeof...(Dims) + 1 <= 3, + "work_group_size property currently only supports up to three values."); + static_assert(detail::FunctionPropertyAllNonZero::value, + "work_group_size property must only contain non-zero values."); + + static constexpr size_t dimensions = sizeof...(Dims) + 1; + + constexpr size_t operator[](int Dim) const { + constexpr size_t Values[] = {Dim0, Dims...}; + return Values[Dim]; + } +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::WorkGroupSizeHint, work_group_size_hint_key> { + static_assert(sizeof...(Dims) + 1 <= 3, + "work_group_size_hint property currently only supports up to " + "three values."); + static_assert( + detail::FunctionPropertyAllNonZero::value, + "work_group_size_hint property must only contain non-zero values."); + + constexpr size_t operator[](int Dim) const { + constexpr size_t Values[] = {Dim0, Dims...}; + return Values[Dim]; + } +}; + +template +struct property_value> + : detail::property_base< + property_value>, + detail::PropKind::SubGroupSize, sub_group_size_key> { + static_assert(Size != 0, + "sub_group_size property must contain a non-zero value."); + + using value_t = std::integral_constant; + static constexpr uint32_t value = Size; +}; + +template +struct property_value, + std::integral_constant...> + : detail::property_base< + property_value, + std::integral_constant...>, + detail::PropKind::MaxWorkGroupSize, max_work_group_size_key> { + static_assert( + sizeof...(Dims) + 1 <= 3, + "max_work_group_size currently only supports up to three values."); + static_assert(detail::FunctionPropertyAllNonZero::value, + "max_work_group_size must only contain non-zero values."); + + static constexpr size_t dimensions = sizeof...(Dims) + 1; + + constexpr size_t operator[](int Dim) const { + constexpr size_t Values[] = {Dim0, Dims...}; + return Values[Dim]; + } +}; + +template +struct property_value> + : detail::property_base< + property_value>, + detail::PropKind::MaxLinearWorkGroupSize, + max_linear_work_group_size_key> { + static_assert(Size != 0, + "max_linear_work_group_size must contain a non-zero value."); + + using value_t = std::integral_constant; + static constexpr size_t value = Size; +}; + +template +inline constexpr work_group_size_key::value_t work_group_size; + +template +inline constexpr work_group_size_hint_key::value_t + work_group_size_hint; + +template +inline constexpr sub_group_size_key::value_t sub_group_size; + +template +inline constexpr max_work_group_size_key::value_t + max_work_group_size; + +template +inline constexpr max_linear_work_group_size_key::value_t + max_linear_work_group_size; + +namespace detail { + +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-nd-range-kernel"; + static constexpr int value = Dims; +}; + +template <> struct FunctionPropertyMetaInfo { + static constexpr const char *name = "sycl-single-task-kernel"; + static constexpr int value = 0; +}; + +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct HasCompileTimeEffect> + : std::true_type {}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-work-group-size-hint"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo< + work_group_size_hint_key::value_t> { + static constexpr const char *name = "sycl-work-group-size-hint"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-sub-group-size"; + static constexpr uint32_t value = Size; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-sub-group-size"; + static constexpr uint32_t value = Size; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-max-work-group-size"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; +template +struct FunctionPropertyMetaInfo< + max_work_group_size_key::value_t> { + static constexpr const char *name = "sycl-max-work-group-size"; + static constexpr const char *value = + FunctionPropertySizeListToStr::value; +}; + +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-max-linear-work-group-size"; + static constexpr size_t value = Size; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-max-linear-work-group-size"; + static constexpr size_t value = Size; +}; + +} // namespace detail +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl + +#ifdef __SYCL_DEVICE_ONLY__ +#define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) \ + [[__sycl_detail__::add_ir_attributes_function( \ + sycl::ext::oneapi::experimental::detail::FunctionPropertyMetaInfo< \ + sycl::ext::oneapi::experimental::detail::remove_cvref_t< \ + decltype(PROP)>>::name, \ + sycl::ext::oneapi::experimental::detail::FunctionPropertyMetaInfo< \ + sycl::ext::oneapi::experimental::detail::remove_cvref_t< \ + decltype(PROP)>>::value)]] +#else +#define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) +#endif \ No newline at end of file diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties.hpp new file mode 100644 index 0000000000000..06a01cb51288f --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/kernel_properties.hpp @@ -0,0 +1,184 @@ +//==----------- kernel_properties.hpp - SYCL kernel properties ------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { + +struct properties_tag {}; + +struct device_has_key + : detail::compile_time_property_key { + template + using value_t = property_value...>; +}; + +template +struct property_value...> + : detail::property_base< + property_value...>, + detail::PropKind::DeviceHas, device_has_key> { + static constexpr std::array value{Aspects...}; +}; + +template +inline constexpr device_has_key::value_t device_has; + +struct work_group_progress_key + : detail::compile_time_property_key { + template + using value_t = property_value< + work_group_progress_key, + std::integral_constant, + std::integral_constant>; +}; + +struct sub_group_progress_key + : detail::compile_time_property_key { + template + using value_t = property_value< + sub_group_progress_key, + std::integral_constant, + std::integral_constant>; +}; + +struct work_item_progress_key + : detail::compile_time_property_key { + template + using value_t = property_value< + work_item_progress_key, + std::integral_constant, + std::integral_constant>; +}; + +template +struct property_value< + work_group_progress_key, + std::integral_constant, + std::integral_constant> + : detail::property_base< + property_value< + work_group_progress_key, + std::integral_constant, + std::integral_constant>, + detail::PropKind::WorkGroupProgress, work_group_progress_key> { + static constexpr forward_progress_guarantee guarantee = Guarantee; + static constexpr execution_scope coordinationScope = CoordinationScope; +}; + +template +struct property_value< + sub_group_progress_key, + std::integral_constant, + std::integral_constant> + : detail::property_base< + property_value< + sub_group_progress_key, + std::integral_constant, + std::integral_constant>, + detail::PropKind::SubGroupProgress, sub_group_progress_key> { + static constexpr forward_progress_guarantee guarantee = Guarantee; + static constexpr execution_scope coordinationScope = CoordinationScope; +}; + +template +struct property_value< + work_item_progress_key, + std::integral_constant, + std::integral_constant> + : detail::property_base< + property_value< + work_item_progress_key, + std::integral_constant, + std::integral_constant>, + detail::PropKind::WorkItemProgress, work_item_progress_key> { + static constexpr forward_progress_guarantee guarantee = Guarantee; + static constexpr execution_scope coordinationScope = CoordinationScope; +}; + +template +inline constexpr work_group_progress_key::value_t + work_group_progress; + +template +inline constexpr sub_group_progress_key::value_t + sub_group_progress; + +template +inline constexpr work_item_progress_key::value_t + work_item_progress; + +namespace detail { +template +struct HasCompileTimeEffect> + : std::true_type {}; +template +struct PropertyMetaInfo> { + static constexpr const char *name = "sycl-device-has"; + static constexpr const char *value = + SizeListToStr(Aspects)...>::value; +}; +template +struct FunctionPropertyMetaInfo> { + static constexpr const char *name = "sycl-device-has"; + static constexpr const char *value = + SizeListToStr(Aspects)...>::value; +}; + +template +struct HasKernelPropertiesGetMethod : std::false_type {}; + +template +struct HasKernelPropertiesGetMethod().get( + std::declval()))>> + : std::true_type { + using properties_t = + decltype(std::declval().get(std::declval())); +}; + +template +auto RetrieveGetMethodPropertiesOrEmpty(RestT &&...Rest) { + auto Identity = [](const auto &x) -> decltype(auto) { return x; }; + const auto &KernelObj = (Identity(Rest), ...); + if constexpr (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< + decltype(KernelObj)>::value) { + return KernelObj.get(ext::oneapi::experimental::properties_tag{}); + } else { + return ext::oneapi::experimental::empty_properties_t{}; + } +} + +} // namespace detail +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl \ No newline at end of file diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp index bd2d7b8c41172..de7babba8c99f 100644 --- a/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp @@ -1,4 +1,4 @@ -//==--- function_properties.hpp - SYCL standalone function properties -----==// +//==--- function_properties.hpp - compatibility forwarding header ---------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -8,396 +8,6 @@ #pragma once -// This header is the lightweight split entry point for free-function kernel -// annotations and standalone compile-time kernel properties that do not need -// the umbrella property's property-list machinery. Keep these public -// property_value definitions here so standalone users and umbrella users -// observe the same decltype(...) while avoiding the heavier machinery on this -// path. - -#include -#include -#include - -#include - -namespace sycl { -inline namespace _V1 { -namespace ext::oneapi::experimental { -namespace detail { - -template -using remove_cvref_t = std::remove_cv_t>; - -template struct FunctionPropertyMetaInfo; -template struct ConflictingProperties; - -template struct FunctionPropertyAllNonZero { - static constexpr bool value = true; -}; -template struct FunctionPropertyAllNonZero { - static constexpr bool value = - X > 0 && FunctionPropertyAllNonZero::value; -}; - -inline constexpr size_t FunctionPropertyDecimalBase = 10; - -template struct FunctionPropertySizeList {}; -template struct FunctionPropertyCharList {}; - -template struct FunctionPropertyCharsToStr { - static constexpr const char value[] = {Chars..., '\0'}; -}; - -template -struct FunctionPropertySizeListToStrHelper; - -template -struct FunctionPropertySizeListToStrHelper< - FunctionPropertySizeList, - FunctionPropertyCharList, Chars...> - : FunctionPropertySizeListToStrHelper< - FunctionPropertySizeList, - FunctionPropertyCharList, - '0' + (Value % FunctionPropertyDecimalBase), Chars...> {}; - -template -struct FunctionPropertySizeListToStrHelper< - FunctionPropertySizeList<0, Values...>, - FunctionPropertyCharList, Chars...> - : FunctionPropertySizeListToStrHelper< - FunctionPropertySizeList, - FunctionPropertyCharList> {}; - -template -struct FunctionPropertySizeListToStrHelper< - FunctionPropertySizeList<0, Values...>, - FunctionPropertyCharList> - : FunctionPropertySizeListToStrHelper< - FunctionPropertySizeList, - FunctionPropertyCharList> {}; - -template -struct FunctionPropertySizeListToStrHelper< - FunctionPropertySizeList<0>, FunctionPropertyCharList, - Chars...> : FunctionPropertyCharsToStr {}; - -template -struct FunctionPropertySizeListToStrHelper< - FunctionPropertySizeList<0>, FunctionPropertyCharList> - : FunctionPropertyCharsToStr {}; - -template <> -struct FunctionPropertySizeListToStrHelper, - FunctionPropertyCharList<>> - : FunctionPropertyCharsToStr<> {}; - -template -struct FunctionPropertySizeListToStr - : FunctionPropertySizeListToStrHelper, - FunctionPropertyCharList<>> {}; - -} // namespace detail - -struct nd_range_kernel_key - : detail::compile_time_property_key { - template - using value_t = - property_value>; -}; - -struct single_task_kernel_key - : detail::compile_time_property_key { - using value_t = property_value; -}; - -template -struct property_value> - : detail::property_base>, - detail::PropKind::NDRangeKernel, - nd_range_kernel_key> { - static_assert(Dims >= 1 && Dims <= 3, - "nd_range_kernel must use dimension 1, 2, or 3."); - - using value_t = int; - static constexpr int dimensions = Dims; -}; - -template <> -struct property_value - : detail::property_base, - detail::PropKind::SingleTaskKernel, - single_task_kernel_key> {}; - -template -inline constexpr nd_range_kernel_key::value_t nd_range_kernel; - -inline constexpr single_task_kernel_key::value_t single_task_kernel; - -struct work_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct work_group_size_hint_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct sub_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value>; -}; - -struct max_work_group_size_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -struct max_linear_work_group_size_key - : detail::compile_time_property_key< - detail::PropKind::MaxLinearWorkGroupSize> { - template - using value_t = property_value>; -}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::WorkGroupSize, work_group_size_key> { - static_assert( - sizeof...(Dims) + 1 <= 3, - "work_group_size property currently only supports up to three values."); - static_assert(detail::FunctionPropertyAllNonZero::value, - "work_group_size property must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - constexpr size_t Values[] = {Dim0, Dims...}; - return Values[Dim]; - } - -private: - constexpr size_t size() const { return sizeof...(Dims) + 1; } - - template friend struct detail::ConflictingProperties; -}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::WorkGroupSizeHint, work_group_size_hint_key> { - static_assert(sizeof...(Dims) + 1 <= 3, - "work_group_size_hint property currently only supports up to " - "three values."); - static_assert( - detail::FunctionPropertyAllNonZero::value, - "work_group_size_hint property must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - constexpr size_t Values[] = {Dim0, Dims...}; - return Values[Dim]; - } -}; - -template -struct property_value> - : detail::property_base< - property_value>, - detail::PropKind::SubGroupSize, sub_group_size_key> { - static_assert(Size != 0, - "sub_group_size property must contain a non-zero value."); - - using value_t = std::integral_constant; - static constexpr uint32_t value = Size; -}; - -template -struct property_value, - std::integral_constant...> - : detail::property_base< - property_value, - std::integral_constant...>, - detail::PropKind::MaxWorkGroupSize, max_work_group_size_key> { - static_assert( - sizeof...(Dims) + 1 <= 3, - "max_work_group_size currently only supports up to three values."); - static_assert(detail::FunctionPropertyAllNonZero::value, - "max_work_group_size must only contain non-zero values."); - - constexpr size_t operator[](int Dim) const { - constexpr size_t Values[] = {Dim0, Dims...}; - return Values[Dim]; - } - -private: - constexpr size_t size() const { return sizeof...(Dims) + 1; } - - template friend struct detail::ConflictingProperties; -}; - -template -struct property_value> - : detail::property_base< - property_value>, - detail::PropKind::MaxLinearWorkGroupSize, - max_linear_work_group_size_key> { - static_assert(Size != 0, - "max_linear_work_group_size must contain a non-zero value."); - - using value_t = std::integral_constant; - static constexpr size_t value = Size; -}; - -template -inline constexpr work_group_size_key::value_t work_group_size; - -template -inline constexpr work_group_size_hint_key::value_t - work_group_size_hint; - -template -inline constexpr sub_group_size_key::value_t sub_group_size; - -template -inline constexpr max_work_group_size_key::value_t - max_work_group_size; - -template -inline constexpr max_linear_work_group_size_key::value_t - max_linear_work_group_size; - -namespace detail { - -template -struct FunctionPropertyMetaInfo> { - static constexpr const char *name = "sycl-nd-range-kernel"; - static constexpr int value = Dims; -}; - -template <> struct FunctionPropertyMetaInfo { - static constexpr const char *name = "sycl-single-task-kernel"; - static constexpr int value = 0; -}; - -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct HasCompileTimeEffect> - : std::true_type {}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-work-group-size"; - static constexpr const char *value = - FunctionPropertySizeListToStr::value; -}; -template -struct FunctionPropertyMetaInfo> { - static constexpr const char *name = "sycl-work-group-size"; - static constexpr const char *value = - FunctionPropertySizeListToStr::value; -}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-work-group-size-hint"; - static constexpr const char *value = - FunctionPropertySizeListToStr::value; -}; -template -struct FunctionPropertyMetaInfo< - work_group_size_hint_key::value_t> { - static constexpr const char *name = "sycl-work-group-size-hint"; - static constexpr const char *value = - FunctionPropertySizeListToStr::value; -}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-sub-group-size"; - static constexpr uint32_t value = Size; -}; -template -struct FunctionPropertyMetaInfo> { - static constexpr const char *name = "sycl-sub-group-size"; - static constexpr uint32_t value = Size; -}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-max-work-group-size"; - static constexpr const char *value = - FunctionPropertySizeListToStr::value; -}; -template -struct FunctionPropertyMetaInfo< - max_work_group_size_key::value_t> { - static constexpr const char *name = "sycl-max-work-group-size"; - static constexpr const char *value = - FunctionPropertySizeListToStr::value; -}; - -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-max-linear-work-group-size"; - static constexpr size_t value = Size; -}; -template -struct FunctionPropertyMetaInfo> { - static constexpr const char *name = "sycl-max-linear-work-group-size"; - static constexpr size_t value = Size; -}; - -} // namespace detail -} // namespace ext::oneapi::experimental -} // namespace _V1 -} // namespace sycl - -#ifdef __SYCL_DEVICE_ONLY__ -#define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) \ - [[__sycl_detail__::add_ir_attributes_function( \ - sycl::ext::oneapi::experimental::detail::FunctionPropertyMetaInfo< \ - sycl::ext::oneapi::experimental::detail::remove_cvref_t< \ - decltype(PROP)>>::name, \ - sycl::ext::oneapi::experimental::detail::FunctionPropertyMetaInfo< \ - sycl::ext::oneapi::experimental::detail::remove_cvref_t< \ - decltype(PROP)>>::value)]] -#else -#define SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(PROP) -#endif +// Compatibility shim for the legacy kernel_properties/* include path. +// Delete this header when support for that include path is removed. +#include diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp index b26d68a4d3be1..11f58071799c1 100644 --- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp @@ -1,4 +1,4 @@ -//==------- properties.hpp - SYCL properties associated with kernels -------==// +//==------- properties.hpp - compatibility forwarding header --------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -8,231 +8,6 @@ #pragma once -#include // for array -#include // for size_t -#include // for uint32_t -#include // for aspect -#include // for forward_progress_guarantee enum -#include -#include -#include // for true_type -#include // for declval -namespace sycl { -inline namespace _V1 { -namespace ext::oneapi::experimental { - -struct properties_tag {}; - -struct device_has_key - : detail::compile_time_property_key { - template - using value_t = property_value...>; -}; - -template -struct property_value...> - : detail::property_base< - property_value...>, - detail::PropKind::DeviceHas, device_has_key> { - static constexpr std::array value{Aspects...}; -}; - -template -inline constexpr device_has_key::value_t device_has; - -struct work_group_progress_key - : detail::compile_time_property_key { - template - using value_t = property_value< - work_group_progress_key, - std::integral_constant, - std::integral_constant>; -}; - -struct sub_group_progress_key - : detail::compile_time_property_key { - template - using value_t = property_value< - sub_group_progress_key, - std::integral_constant, - std::integral_constant>; -}; - -struct work_item_progress_key - : detail::compile_time_property_key { - template - using value_t = property_value< - work_item_progress_key, - std::integral_constant, - std::integral_constant>; -}; - -template -struct property_value< - work_group_progress_key, - std::integral_constant, - std::integral_constant> - : detail::property_base< - property_value< - work_group_progress_key, - std::integral_constant, - std::integral_constant>, - detail::PropKind::WorkGroupProgress, work_group_progress_key> { - static constexpr forward_progress_guarantee guarantee = Guarantee; - static constexpr execution_scope coordinationScope = CoordinationScope; -}; - -template -struct property_value< - sub_group_progress_key, - std::integral_constant, - std::integral_constant> - : detail::property_base< - property_value< - sub_group_progress_key, - std::integral_constant, - std::integral_constant>, - detail::PropKind::SubGroupProgress, sub_group_progress_key> { - static constexpr forward_progress_guarantee guarantee = Guarantee; - static constexpr execution_scope coordinationScope = CoordinationScope; -}; - -template -struct property_value< - work_item_progress_key, - std::integral_constant, - std::integral_constant> - : detail::property_base< - property_value< - work_item_progress_key, - std::integral_constant, - std::integral_constant>, - detail::PropKind::WorkItemProgress, work_item_progress_key> { - static constexpr forward_progress_guarantee guarantee = Guarantee; - static constexpr execution_scope coordinationScope = CoordinationScope; -}; - -template -inline constexpr work_group_progress_key::value_t - work_group_progress; - -template -inline constexpr sub_group_progress_key::value_t - sub_group_progress; - -template -inline constexpr work_item_progress_key::value_t - work_item_progress; - -namespace detail { -template -struct HasCompileTimeEffect> - : std::true_type {}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-device-has"; - static constexpr const char *value = - SizeListToStr(Aspects)...>::value; -}; -template -struct FunctionPropertyMetaInfo> { - static constexpr const char *name = "sycl-device-has"; - static constexpr const char *value = - SizeListToStr(Aspects)...>::value; -}; - -template -struct HasKernelPropertiesGetMethod : std::false_type {}; - -template -struct HasKernelPropertiesGetMethod().get( - std::declval()))>> - : std::true_type { - using properties_t = - decltype(std::declval().get(std::declval())); -}; - -// If work_group_size and max_work_group_size coexist, check that the -// dimensionality matches and that the required work-group size doesn't -// trivially exceed the maximum size. -template -struct ConflictingProperties { - static constexpr bool value = []() constexpr { - if constexpr (Properties::template has_property()) { - constexpr auto wg_size = - Properties::template get_property(); - constexpr auto max_wg_size = - Properties::template get_property(); - static_assert( - wg_size.size() == max_wg_size.size(), - "work_group_size and max_work_group_size dimensionality must match"); - if constexpr (wg_size.size() == max_wg_size.size()) { - constexpr auto Dims = wg_size.size(); - static_assert(Dims < 1 || wg_size[0] <= max_wg_size[0], - "work_group_size must not exceed max_work_group_size"); - static_assert(Dims < 2 || wg_size[1] <= max_wg_size[1], - "work_group_size must not exceed max_work_group_size"); - static_assert(Dims < 3 || wg_size[2] <= max_wg_size[2], - "work_group_size must not exceed max_work_group_size"); - } - } - return false; - }(); -}; - -// If work_group_size and max_linear_work_group_size coexist, check that the -// required linear work-group size doesn't trivially exceed the maximum size. -template -struct ConflictingProperties { - static constexpr bool value = []() constexpr { - if constexpr (Properties::template has_property()) { - constexpr auto wg_size = - Properties::template get_property(); - constexpr auto dims = wg_size.size(); - constexpr auto linear_size = wg_size[0] * (dims > 1 ? wg_size[1] : 1) * - (dims > 2 ? wg_size[2] : 1); - constexpr auto max_linear_wg_size = - Properties::template get_property(); - static_assert( - linear_size < max_linear_wg_size.value, - "work_group_size must not exceed max_linear_work_group_size"); - } - return false; - }(); -}; - -// If the kernel (last element in the parameter pack) has a get(properties_tag) -// method, return the property list specified by this getter. Otherwise, return -// an empty properety list. -template -auto RetrieveGetMethodPropertiesOrEmpty(RestT &&...Rest) { - // Note: the following trivial identity lambda is used to avoid the issue - // that line "const auto &KernelObj = (Rest, ...);" may result in a "left - // operand of comma operator has no effect" error for certain compiler(s) - auto Identity = [](const auto &x) -> decltype(auto) { return x; }; - const auto &KernelObj = (Identity(Rest), ...); - if constexpr (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod< - decltype(KernelObj)>::value) { - return KernelObj.get(ext::oneapi::experimental::properties_tag{}); - } else { - return ext::oneapi::experimental::empty_properties_t{}; - } -} - -} // namespace detail -} // namespace ext::oneapi::experimental -} // namespace _V1 -} // namespace sycl +// Compatibility shim for the legacy kernel_properties/* include path. +// Delete this header when support for that include path is removed. +#include diff --git a/sycl/include/sycl/ext/oneapi/properties.hpp b/sycl/include/sycl/ext/oneapi/properties.hpp new file mode 100644 index 0000000000000..158e75e49b682 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/properties.hpp @@ -0,0 +1,421 @@ +//==---------------- properties.hpp - SYCL oneAPI properties --------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +#include +#include +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace detail {} +namespace ext::oneapi::experimental { + +template class __SYCL_EBO properties; + +namespace detail { +using namespace sycl::detail; + +// Keep a distinct name for the local helper to avoid colliding with the +// imported sycl::detail::nth_type_t alias. +#if __has_builtin(__type_pack_element) +template +using properties_nth_type_t = __type_pack_element; +#else +template struct nth_type { + using type = typename nth_type::type; +}; + +template struct nth_type<0, T, Ts...> { + using type = T; +}; + +template +using properties_nth_type_t = typename nth_type::type; +#endif + +// NOTE: Meta-function to implement CTAD rules isn't allowed to return +// `properties` and it's impossible to return a pack as well. As +// such, we're forced to have an extra level of `detail::properties_type_list` +// for the purpose of providing CTAD rules. +template struct properties_type_list; + +// This is used in a separate `properties` specialization to report friendlier +// errors. +template struct invalid_properties_type_list {}; + +// Helper for reconstructing a properties type. This assumes that +// PropertyValueTs is sorted and contains only valid properties. +// +// It also allows us to hide details of `properties` implementation from the +// code that uses/defines them (with the exception of ESIMD which is extremely +// hacky in its own esimd::properties piggybacking on these ones). +template +using properties_t = + properties>; + +template +inline constexpr bool properties_are_unique = []() constexpr -> bool { + if constexpr (sizeof...(property_tys) == 0) { + return true; + } else { + const std::array kinds = {PropertyID::value...}; + auto N = kinds.size(); + for (std::size_t i = 0; i < N; ++i) + for (std::size_t j = i + 1; j < N; ++j) + if (kinds[i] == kinds[j]) + return false; + + return true; + } +}(); + +template +inline constexpr bool properties_are_sorted = []() constexpr -> bool { + if constexpr (sizeof...(property_tys) == 0) { + return true; + } else { + const std::array kinds = {PropertyID::value...}; + // std::is_sorted isn't constexpr until C++20. + for (std::size_t idx = 1; idx < kinds.size(); ++idx) + if (kinds[idx - 1] >= kinds[idx]) + return false; + return true; + } +}(); + +template +constexpr bool properties_are_valid_for_ctad = []() constexpr { + // Need `if constexpr` to avoid hard error in "unique" check when querying + // property kind if `property_tys` isn't a property. + if constexpr (!((is_property_value_v && ...))) { + return false; + } else if constexpr (!detail::properties_are_unique) { + return false; + } else { + return true; + } +}(); + +template struct properties_sorter { + // Not using "auto" due to MSVC bug in v19.36 and older. v19.37 and later is + // able to compile "auto" just fine. See https://godbolt.org/z/eW3rjjs7n. + static constexpr std::array sorted_indices = + []() constexpr { + int idx = 0; + int N = sizeof...(property_tys); + // std::sort isn't constexpr until C++20. Also, it's possible there will + // be a compiler builtin to sort types, in which case we should start + // using that. + std::array to_sort{ + std::pair{PropertyID::value, idx++}...}; + auto swap_pair = [](auto &x, auto &y) constexpr { + auto tmp_first = x.first; + auto tmp_second = x.second; + x.first = y.first; + x.second = y.second; + y.first = tmp_first; + y.second = tmp_second; + }; + for (int i = 0; i < N; ++i) + for (int j = i; j < N; ++j) + if (to_sort[j].first < to_sort[i].first) + swap_pair(to_sort[i], to_sort[j]); + + std::array sorted_indices{}; + for (int i = 0; i < N; ++i) + sorted_indices[i] = to_sort[i].second; + + return sorted_indices; + }(); + + template struct helper; + template + struct helper> { + using type = properties_type_list< + properties_nth_type_t...>; + }; + + using type = typename helper< + std::make_integer_sequence>::type; +}; +// Specialization to avoid zero-size array creation. +template <> struct properties_sorter<> { + using type = properties_type_list<>; +}; + +} // namespace detail + +// Empty property list. +template <> class __SYCL_EBO properties> { + template + static constexpr bool empty_properties_list_contains = false; + +public: + template static constexpr bool has_property() { + return false; + } + + // Never exists for empty property list, provide this for a better error + // message: + template + static std::enable_if_t> get_property() {} +}; + +// Base implementation to provide nice user error in case of mis-use. Without it +// an error "base class '' specified more than once as a direct base +// class" is reported prior to static_assert's error. +template +class __SYCL_EBO + properties> { +public: + properties(property_tys...) { + if constexpr (!((is_property_value_v && ...))) { + static_assert(((is_property_value_v && ...)), + "Non-property argument!"); + } else { + // This is a separate specialization to report an error, we can afford + // doing extra work to provide nice error message without sacrificing + // compile time on non-exceptional path. Let's find *a* pair of properties + // that failed the check. Note that there might be multiple duplicate + // names, we're only reporting one instance. Once user addresses that, the + // next pair will be reported. + static constexpr auto conflict = []() constexpr { + const std::array kinds = {detail::PropertyID::value...}; + auto N = kinds.size(); + for (int i = 0; i < N; ++i) + for (int j = i + 1; j < N; ++j) + if (kinds[i] == kinds[j]) + return std::pair{i, j}; + }(); + using first_type = + detail::properties_nth_type_t; + using second_type = + detail::properties_nth_type_t; + if constexpr (std::is_same_v) { + static_assert(!std::is_same_v, + "Duplicate properties in property list."); + } else { + static_assert( + detail::PropertyToKind::Kind != + detail::PropertyToKind::Kind, + "Property Kind collision between different property keys!"); + } + } + } + + template static constexpr bool has_property() { + return false; + } +}; + +template +class __SYCL_EBO properties> + : private property_tys... { + static_assert(detail::properties_are_sorted, + "Properties must be sorted!"); + using property_tys::get_property_impl...; + + template friend class __SYCL_EBO properties; + + template static constexpr bool is_valid_ctor_arg() { + return ((std::is_same_v || ...)); + } + + template + static constexpr bool can_be_constructed_from() { + return std::is_default_constructible_v || + ((false || ... || std::is_same_v)); + } + + // It's possible it shouldn't be that complicated, but clang doesn't accept + // simpler version: https://godbolt.org/z/oPff4h738, reported upstream at + // https://github.com/llvm/llvm-project/issues/115547. Note that if the + // `decltype(...)` is "inlined" then it has no issues with it, but that's too + // verbose. + struct helper : property_tys... { + using property_tys::get_property_impl...; + }; + template + using prop_t = decltype(std::declval().get_property_impl( + detail::property_key_tag{})); + +public: + template < + typename... unsorted_property_tys, + typename = std::enable_if_t< + ((is_valid_ctor_arg() && ...))>, + typename = std::enable_if_t< + ((can_be_constructed_from() && + ...))>, + typename = std::enable_if_t< + detail::properties_are_unique>> + constexpr properties(unsorted_property_tys... props); + + template static constexpr bool has_property() { + return std::is_base_of_v, + properties>; + } + + template + static constexpr auto + get_property() -> std::enable_if_t>, + prop_t> { + return prop_t{}; + } + + template + constexpr auto get_property(int = 0) const + -> std::enable_if_t>, + prop_t> { + return get_property_impl(detail::property_key_tag{}); + } +}; + +template +template +constexpr properties>::properties( + unsorted_property_tys... props) + : unsorted_property_tys(props)... { + static_assert(((!detail::ConflictingProperties::value && + ...)), + "Conflicting properties in property list."); +} + +template >> +properties(unsorted_property_tys... props) + -> properties< + typename detail::properties_sorter::type>; + +template >> +properties(unsorted_property_tys... props) + -> properties< + detail::invalid_properties_type_list>; + +using empty_properties_t = decltype(properties{}); + +namespace detail { + +template