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..d547fcdeaac88 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,27 @@ supports. feature-test macro always has this value. |=== +=== Headers + +The APIs defined in this extension are provided by either of the following header files: + +* `` +* `` + +In addition, the following lightweight header provides a subset of the APIs from this extension: + +* `` + +This lightweight header provides `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, and the +properties `nd_range_kernel` and `single_task_kernel`. See +link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[ +sycl_ext_oneapi_kernel_properties] for other APIs that are provided by this +header. + +[_Note:_ The lightweight header is intended for cases where fast compilation +time is a priority. +_{endnote}_] + === Defining a free function kernel A free function kernel is a normal C++ function definition, where the function @@ -804,6 +825,8 @@ sycl_ext_oneapi_kernel_properties] by applying the properties to the function declaration as illustrated below. ``` +#include + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size<32>)) void iota(float start, float *ptr) { 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..c530d65d6fb17 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc @@ -9,6 +9,7 @@ :encoding: utf-8 :lang: en :dpcpp: pass:[DPC++] +:endnote: —{nbsp}end{nbsp}note :blank: pass:[ +] @@ -115,6 +116,28 @@ supports. feature-test macro always has this value. |=== +=== Headers + +The APIs defined in this extension are provided by either of the following header files: + +* `` +* `` + +In addition, the following lightweight header provides a subset of the APIs that +is suitable for applications that define kernels using the +link:../experimental/sycl_ext_oneapi_free_function_kernels.asciidoc[ +sycl_ext_oneapi_free_function_kernels] extension: + +* `` + +This lightweight header provides `SYCL_EXT_ONEAPI_FUNCTION_PROPERTY`, +and the properties `work_group_size`, `work_group_size_hint`, +`sub_group_size`, `max_work_group_size`, and `max_linear_work_group_size`. + +[_Note:_ The lightweight header is intended for cases where +kernels are defined using the free-function kernel syntax and fast +compilation time is a priority._{endnote}_] + === Kernel Properties The kernel properties below correspond to kernel attributes defined in diff --git a/sycl/include/sycl/detail/kernel_launch_helper.hpp b/sycl/include/sycl/detail/kernel_launch_helper.hpp index 909f6f2f49ae8..f9f1f31462bee 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,48 @@ 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< + max_work_group_size_key>()) { + 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 +412,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 9b13f6e3ed123..7fd51920f020b 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/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..09f7b6e99d4ae --- /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..fe4bd30269580 --- /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 new file mode 100644 index 0000000000000..de7babba8c99f --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/function_properties.hpp @@ -0,0 +1,13 @@ +//==--- 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. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +// 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 defb1589679b9..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,471 +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 // 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 - using value_t = property_value...>; -}; - -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...> - : detail::property_base< - property_value...>, - detail::PropKind::DeviceHas, device_has_key> { - static constexpr std::array 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 - 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 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 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 - -#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 +// 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..f00d0ae9577e9 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/properties.hpp @@ -0,0 +1,419 @@ +//==---------------- 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