diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp index 8362e83f445b..6548f3d3c367 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/annotated_usm/alloc_util.hpp @@ -45,23 +45,12 @@ using HasUsmKind = HasProperty; template using HasBufferLocation = HasProperty; -// Get the value of a property from a property list -template -struct GetPropertyValueFromPropList {}; - template struct GetPropertyValueFromPropList> { - using prop_val_t = std::conditional_t< - detail::ContainsProperty>::value, - typename detail::FindCompileTimePropertyValueType< - PropKey, std::tuple>::type, - DefaultPropVal>; - static constexpr ConstType value = - detail::PropertyMetaInfo>::value; -}; + detail::properties_t> + : GetPropertyValueFromPropList> {}; // Get the value of alignment from a property list // If alignment is not present in the property list, set to default value 0 diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp index 2c9a31cf05ed..70055c6680b7 100644 --- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp @@ -9,13 +9,12 @@ #pragma once #include // for array +#include #include // for size_t #include // for uint32_T #include // for aspect #include // for forward_progress_guarantee enum -#include // for PropKind -#include // for SizeListToStr -#include // for property_value +#include #include // for true_type #include // for declval namespace sycl { @@ -351,6 +350,80 @@ struct HasKernelPropertiesGetMethod().get(std::declval())); }; +// Trait for property compile-time meta names and values. +template struct WGSizePropertyMetaInfo { + static constexpr std::array WGSize = {}; + static constexpr size_t LinearSize = 0; +}; + +template +struct WGSizePropertyMetaInfo> { + static constexpr std::array WGSize = {Dim0, + Dims...}; + static constexpr size_t LinearSize = (Dim0 * ... * Dims); +}; + +template +struct WGSizePropertyMetaInfo> { + static constexpr std::array WGSize = {Dim0, + Dims...}; + static constexpr size_t LinearSize = (Dim0 * ... * Dims); +}; + +// Get the value of a work-group size related property from a property list +template +struct GetWGPropertyFromPropList {}; + +template +struct GetWGPropertyFromPropList> { + using prop_val_t = std::conditional_t< + ContainsProperty>::value, + typename FindCompileTimePropertyValueType< + PropKey, std::tuple>::type, + void>; + static constexpr auto WGSize = + WGSizePropertyMetaInfo>::WGSize; + static constexpr size_t LinearSize = + WGSizePropertyMetaInfo>::LinearSize; +}; + +// 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 + : std::false_type { + using WGSizeVal = GetWGPropertyFromPropList; + using MaxWGSizeVal = + GetWGPropertyFromPropList; + // If work_group_size_key doesn't exist in the list of properties, WGSize is + // an empty array and so Dims == 0. + static constexpr size_t Dims = WGSizeVal::WGSize.size(); + static_assert( + Dims == 0 || Dims == MaxWGSizeVal::WGSize.size(), + "work_group_size and max_work_group_size dimensionality must match"); + static_assert(Dims < 1 || WGSizeVal::WGSize[0] <= MaxWGSizeVal::WGSize[0], + "work_group_size must not exceed max_work_group_size"); + static_assert(Dims < 2 || WGSizeVal::WGSize[1] <= MaxWGSizeVal::WGSize[1], + "work_group_size must not exceed max_work_group_size"); + static_assert(Dims < 3 || WGSizeVal::WGSize[2] <= MaxWGSizeVal::WGSize[2], + "work_group_size must not exceed max_work_group_size"); +}; + +// 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 + : std::false_type { + using WGSizeVal = GetWGPropertyFromPropList; + using MaxLinearWGSizeVal = + GetPropertyValueFromPropList; + static_assert(WGSizeVal::WGSize.empty() || + WGSizeVal::LinearSize <= MaxLinearWGSizeVal::value, + "work_group_size must not exceed max_linear_work_group_size"); +}; + } // namespace detail } // namespace ext::oneapi::experimental } // namespace _V1 diff --git a/sycl/include/sycl/ext/oneapi/properties/properties.hpp b/sycl/include/sycl/ext/oneapi/properties/properties.hpp index ccd44f441a7d..1c93e00dbe88 100644 --- a/sycl/include/sycl/ext/oneapi/properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/properties.hpp @@ -150,6 +150,24 @@ struct ExtractProperties +struct GetPropertyValueFromPropList {}; + +template +struct GetPropertyValueFromPropList> { + using prop_val_t = std::conditional_t< + ContainsProperty>::value, + typename FindCompileTimePropertyValueType< + PropKey, std::tuple>::type, + DefaultPropVal>; + static constexpr ConstType value = + PropertyMetaInfo>::value; +}; + } // namespace detail template class properties { diff --git a/sycl/test/extensions/properties/properties_kernel_negative.cpp b/sycl/test/extensions/properties/properties_kernel_negative.cpp index 4b2f722bdd7f..52cbba23e50e 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative.cpp @@ -317,9 +317,68 @@ void check_sub_group_size() { KernelFunctorWithSGSize<2>{}); } +void check_max_work_group_size() { + sycl::queue Q; + + // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size and max_work_group_size dimensionality must match}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size<2, 2>, + sycl::ext::oneapi::experimental::max_work_group_size<1>}, + []() {}); + + // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_work_group_size}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size<2>, + sycl::ext::oneapi::experimental::max_work_group_size<1>}, + []() {}); + + // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_work_group_size}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size<2, 2>, + sycl::ext::oneapi::experimental::max_work_group_size<2, 1>}, + []() {}); + + // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_work_group_size}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size<2, 2, 2>, + sycl::ext::oneapi::experimental::max_work_group_size<2, 2, 1>}, + []() {}); +} + +void check_max_linear_work_group_size() { + sycl::queue Q; + + // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_linear_work_group_size}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size<2>, + sycl::ext::oneapi::experimental::max_linear_work_group_size<1>}, + []() {}); + + // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_linear_work_group_size}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size<2, 4>, + sycl::ext::oneapi::experimental::max_linear_work_group_size<7>}, + []() {}); + + // expected-error-re@sycl/ext/oneapi/kernel_properties/properties.hpp:* {{static assertion failed due to requirement {{.+}}: work_group_size must not exceed max_linear_work_group_size}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size<2, 4, 2>, + sycl::ext::oneapi::experimental::max_linear_work_group_size<15>}, + []() {}); +} + int main() { check_work_group_size(); check_work_group_size_hint(); check_sub_group_size(); + check_max_work_group_size(); + check_max_linear_work_group_size(); return 0; }