From ccc75b24f079fda3e73799b19147c9f555ed54e2 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris Date: Thu, 23 Jul 2026 06:38:37 -0700 Subject: [PATCH] [SYCL] Remove duplicate FunctionPropertyMetaInfo trait Commit bba4f4d0e36f (#21738) split out the lightweight free_function_kernel_properties.hpp and introduced a parallel detail::FunctionPropertyMetaInfo trait alongside PropertyMetaInfo. This forced every property usable on both the property-list and free-function paths to carry two structurally identical specializations, and caused the grf_size breakage fixed in #22713. The duplication was unnecessary: PropertyMetaInfo (property.hpp) and the shared SizeListToStr/AllNonZero helpers (property_utils.hpp) are already on the lightweight include path, so reusing them is compile-time neutral Drop FunctionPropertyMetaInfo entirely and point the macro at PropertyMetaInfo, remove the duplicate specializations in free_function_kernel_properties.hpp, grf_size_properties.hpp, kernel_properties.hpp and virtual_functions.hpp, and move AllNonZero into property_utils.hpp next to SizeListToStr. Measured compilation times between this PR and HEAD for `free_function_kernel_properties.hpp` remains the same. Fixes #22729 --- .../experimental/grf_size_properties.hpp | 15 --- .../oneapi/experimental/virtual_functions.hpp | 22 ---- .../free_function_kernel_properties.hpp | 124 ++---------------- .../sycl/ext/oneapi/kernel_properties.hpp | 6 - .../ext/oneapi/properties/property_utils.hpp | 8 ++ 5 files changed, 20 insertions(+), 155 deletions(-) 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 6e49a1cc84a6f..395555596fcd6 100644 --- a/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp @@ -53,21 +53,6 @@ struct PropertyMetaInfo< static constexpr unsigned int value = 0; }; -template -struct FunctionPropertyMetaInfo< - sycl::ext::intel::experimental::grf_size_key::value_t> { - static_assert(Size == 128 || Size == 256 || Size == 512, - "Unsupported GRF size"); - static constexpr const char *name = "sycl-grf-size"; - static constexpr unsigned int value = Size; -}; -template <> -struct FunctionPropertyMetaInfo< - sycl::ext::intel::experimental::grf_size_automatic_key::value_t> { - static constexpr const char *name = "sycl-grf-size"; - static constexpr unsigned int value = 0; -}; - template struct ConflictingProperties diff --git a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp index 7fd51920f020b..b70e8506ea1f1 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp @@ -50,17 +50,6 @@ 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. @@ -117,17 +106,6 @@ 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/free_function_kernel_properties.hpp b/sycl/include/sycl/ext/oneapi/free_function_kernel_properties.hpp index 09f7b6e99d4ae..3a38d299fd675 100644 --- a/sycl/include/sycl/ext/oneapi/free_function_kernel_properties.hpp +++ b/sycl/include/sycl/ext/oneapi/free_function_kernel_properties.hpp @@ -10,6 +10,7 @@ #include #include +#include #include #include @@ -22,74 +23,6 @@ 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 @@ -175,7 +108,7 @@ struct property_value, static_assert( sizeof...(Dims) + 1 <= 3, "work_group_size property currently only supports up to three values."); - static_assert(detail::FunctionPropertyAllNonZero::value, + static_assert(detail::AllNonZero::value, "work_group_size property must only contain non-zero values."); static constexpr size_t dimensions = sizeof...(Dims) + 1; @@ -199,7 +132,7 @@ struct property_value::value, + detail::AllNonZero::value, "work_group_size_hint property must only contain non-zero values."); constexpr size_t operator[](int Dim) const { @@ -234,7 +167,7 @@ struct property_value::value, + static_assert(detail::AllNonZero::value, "max_work_group_size must only contain non-zero values."); static constexpr size_t dimensions = sizeof...(Dims) + 1; @@ -281,12 +214,12 @@ inline constexpr max_linear_work_group_size_key::value_t namespace detail { template -struct FunctionPropertyMetaInfo> { +struct PropertyMetaInfo> { static constexpr const char *name = "sycl-nd-range-kernel"; static constexpr int value = Dims; }; -template <> struct FunctionPropertyMetaInfo { +template <> struct PropertyMetaInfo { static constexpr const char *name = "sycl-single-task-kernel"; static constexpr int value = 0; }; @@ -310,28 +243,13 @@ struct HasCompileTimeEffect> 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; + 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 = - 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; + static constexpr const char *value = SizeListToStr::value; }; template @@ -339,24 +257,11 @@ 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; + static constexpr const char *value = SizeListToStr::value; }; template @@ -364,11 +269,6 @@ 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 @@ -378,12 +278,12 @@ struct FunctionPropertyMetaInfo> { #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::PropertyMetaInfo< \ sycl::ext::oneapi::experimental::detail::remove_cvref_t< \ decltype(PROP)>>::name, \ - sycl::ext::oneapi::experimental::detail::FunctionPropertyMetaInfo< \ + sycl::ext::oneapi::experimental::detail::PropertyMetaInfo< \ 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 +#endif diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties.hpp index fe4bd30269580..30e099546ace5 100644 --- a/sycl/include/sycl/ext/oneapi/kernel_properties.hpp +++ b/sycl/include/sycl/ext/oneapi/kernel_properties.hpp @@ -147,12 +147,6 @@ struct PropertyMetaInfo> { 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 {}; diff --git a/sycl/include/sycl/ext/oneapi/properties/property_utils.hpp b/sycl/include/sycl/ext/oneapi/properties/property_utils.hpp index 27d8a09e400c9..8f772032c49c6 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property_utils.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property_utils.hpp @@ -23,6 +23,14 @@ namespace detail { // Property value tooling //****************************************************************************** +// 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; +}; + // Simple helpers for containing primitive types as template arguments. template struct SizeList {}; template struct CharList {};