[NFCI][SYCL] Introduce oneapi::experimental::detail::property_base (#…

…15993)

This lays some ground base for subsequent refactoring PRs by isolating
simple but noisy diff into a separate change.

We're anticipating at least two changes that would be built on top of
this:
* Change implementation of `has_property`/`get_property` to be
inheritance-based vs type-lists/boost.
* Bigger design-wise refactoring of the compile time properties, e.g.
removal of `property_value` class template.

sycl/include/sycl/ext/intel/esimd/memory_properties.hpp

@@ -289,22 +289,34 @@ namespace ext::oneapi::experimental {
 
 template <__ESIMD_NS::cache_hint Hint>
 struct property_value<__ESIMD_NS::cache_hint_L1_key,
-                      std::integral_constant<__ESIMD_NS::cache_hint, Hint>> {
-  using key_t = __ESIMD_NS::cache_hint_L1_key;
+                      std::integral_constant<__ESIMD_NS::cache_hint, Hint>>
+    : detail::property_base<
+          property_value<__ESIMD_NS::cache_hint_L1_key,
+                         std::integral_constant<__ESIMD_NS::cache_hint, Hint>>,
+          oneapi::experimental::detail::PropKind::ESIMDL1CacheHint,
+          __ESIMD_NS::cache_hint_L1_key> {
   static constexpr __ESIMD_NS::cache_level level = __ESIMD_NS::cache_level::L1;
   static constexpr __ESIMD_NS::cache_hint hint = Hint;
 };
 template <__ESIMD_NS::cache_hint Hint>
 struct property_value<__ESIMD_NS::cache_hint_L2_key,
-                      std::integral_constant<__ESIMD_NS::cache_hint, Hint>> {
-  using key_t = __ESIMD_NS::cache_hint_L2_key;
+                      std::integral_constant<__ESIMD_NS::cache_hint, Hint>>
+    : detail::property_base<
+          property_value<__ESIMD_NS::cache_hint_L2_key,
+                         std::integral_constant<__ESIMD_NS::cache_hint, Hint>>,
+          oneapi::experimental::detail::PropKind::ESIMDL2CacheHint,
+          __ESIMD_NS::cache_hint_L2_key> {
   static constexpr __ESIMD_NS::cache_level level = __ESIMD_NS::cache_level::L2;
   static constexpr __ESIMD_NS::cache_hint hint = Hint;
 };
 template <__ESIMD_NS::cache_hint Hint>
 struct property_value<__ESIMD_NS::cache_hint_L3_key,
-                      std::integral_constant<__ESIMD_NS::cache_hint, Hint>> {
-  using key_t = __ESIMD_NS::cache_hint_L3_key;
+                      std::integral_constant<__ESIMD_NS::cache_hint, Hint>>
+    : detail::property_base<
+          property_value<__ESIMD_NS::cache_hint_L3_key,
+                         std::integral_constant<__ESIMD_NS::cache_hint, Hint>>,
+          oneapi::experimental::detail::PropKind::ESIMDL3CacheHint,
+          __ESIMD_NS::cache_hint_L3_key> {
   static constexpr __ESIMD_NS::cache_level level = __ESIMD_NS::cache_level::L3;
   static constexpr __ESIMD_NS::cache_hint hint = Hint;
 };

sycl/include/sycl/ext/intel/experimental/kernel_execution_properties.hpp

@@ -26,8 +26,9 @@ inline constexpr cache_config_enum large_slm =
 inline constexpr cache_config_enum large_data =
     cache_config_enum::large_data;
 
-struct cache_config : oneapi::experimental::detail::run_time_property_key<
-                          oneapi::experimental::detail::PropKind::CacheConfig> {
+struct cache_config
+    : oneapi::experimental::detail::run_time_property_key<
+          cache_config, oneapi::experimental::detail::PropKind::CacheConfig> {
   cache_config(cache_config_enum v) : value(v) {}
   cache_config_enum value;
 };

sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp

@@ -19,6 +19,7 @@ namespace cuda {
 template <int Dim>
 struct cluster_size
     : ::sycl::ext::oneapi::experimental::detail::run_time_property_key<
+          cluster_size<Dim>,
           ::sycl::ext::oneapi::experimental::detail::ClusterLaunch> {
   cluster_size(const range<Dim> &size) : size(size) {}
   sycl::range<Dim> get_cluster_size() { return size; }

sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp

@@ -60,13 +60,15 @@ struct device_has_key
                                  std::integral_constant<aspect, Aspects>...>;
 };
 
-struct nd_range_kernel_key {
+struct nd_range_kernel_key
+    : detail::compile_time_property_key<detail::PropKind::NDRangeKernel> {
   template <int Dims>
   using value_t =
       property_value<nd_range_kernel_key, std::integral_constant<int, Dims>>;
 };
 
-struct single_task_kernel_key {
+struct single_task_kernel_key
+    : detail::compile_time_property_key<detail::PropKind::SingleTaskKernel> {
   using value_t = property_value<single_task_kernel_key>;
 };
 
@@ -87,15 +89,18 @@ struct max_linear_work_group_size_key
 
 template <size_t Dim0, size_t... Dims>
 struct property_value<work_group_size_key, std::integral_constant<size_t, Dim0>,
-                      std::integral_constant<size_t, Dims>...> {
+                      std::integral_constant<size_t, Dims>...>
+    : detail::property_base<
+          property_value<work_group_size_key,
+                         std::integral_constant<size_t, Dim0>,
+                         std::integral_constant<size_t, Dims>...>,
+          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<Dim0, Dims...>::value,
                 "work_group_size property must only contain non-zero values.");
 
-  using key_t = work_group_size_key;
-
   constexpr size_t operator[](int Dim) const {
     return std::array<size_t, sizeof...(Dims) + 1>{Dim0, Dims...}[Dim];
   }
@@ -104,75 +109,94 @@ struct property_value<work_group_size_key, std::integral_constant<size_t, Dim0>,
 template <size_t Dim0, size_t... Dims>
 struct property_value<work_group_size_hint_key,
                       std::integral_constant<size_t, Dim0>,
-                      std::integral_constant<size_t, Dims>...> {
+                      std::integral_constant<size_t, Dims>...>
+    : detail::property_base<
+          property_value<work_group_size_hint_key,
+                         std::integral_constant<size_t, Dim0>,
+                         std::integral_constant<size_t, Dims>...>,
+          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<Dim0, Dims...>::value,
       "work_group_size_hint property must only contain non-zero values.");
 
-  using key_t = work_group_size_hint_key;
-
   constexpr size_t operator[](int Dim) const {
     return std::array<size_t, sizeof...(Dims) + 1>{Dim0, Dims...}[Dim];
   }
 };
 
 template <uint32_t Size>
 struct property_value<sub_group_size_key,
-                      std::integral_constant<uint32_t, Size>> {
+                      std::integral_constant<uint32_t, Size>>
+    : detail::property_base<
+          property_value<sub_group_size_key,
+                         std::integral_constant<uint32_t, Size>>,
+          detail::PropKind::SubGroupSize, sub_group_size_key> {
   static_assert(Size != 0,
                 "sub_group_size_key property must contain a non-zero value.");
 
-  using key_t = sub_group_size_key;
   using value_t = std::integral_constant<uint32_t, Size>;
   static constexpr uint32_t value = Size;
 };
 
 template <aspect... Aspects>
 struct property_value<device_has_key,
-                      std::integral_constant<aspect, Aspects>...> {
-  using key_t = device_has_key;
+                      std::integral_constant<aspect, Aspects>...>
+    : detail::property_base<
+          property_value<device_has_key,
+                         std::integral_constant<aspect, Aspects>...>,
+          detail::PropKind::DeviceHas, device_has_key> {
   static constexpr std::array<aspect, sizeof...(Aspects)> value{Aspects...};
 };
 
 template <int Dims>
-struct property_value<nd_range_kernel_key, std::integral_constant<int, Dims>> {
+struct property_value<nd_range_kernel_key, std::integral_constant<int, Dims>>
+    : detail::property_base<property_value<nd_range_kernel_key,
+                                           std::integral_constant<int, Dims>>,
+                            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 key_t = nd_range_kernel_key;
   using value_t = int;
   static constexpr int dimensions = Dims;
 };
 
-template <> struct property_value<single_task_kernel_key> {
-  using key_t = single_task_kernel_key;
-};
+template <>
+struct property_value<single_task_kernel_key>
+    : detail::property_base<property_value<single_task_kernel_key>,
+                            detail::PropKind::SingleTaskKernel,
+                            single_task_kernel_key> {};
 
 template <size_t Dim0, size_t... Dims>
 struct property_value<max_work_group_size_key,
                       std::integral_constant<size_t, Dim0>,
-                      std::integral_constant<size_t, Dims>...> {
+                      std::integral_constant<size_t, Dims>...>
+    : detail::property_base<
+          property_value<max_work_group_size_key,
+                         std::integral_constant<size_t, Dim0>,
+                         std::integral_constant<size_t, Dims>...>,
+          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<Dim0, Dims...>::value,
       "max_work_group_size property must only contain non-zero values.");
 
-  using key_t = max_work_group_size_key;
-
   constexpr size_t operator[](int Dim) const {
     return std::array<size_t, sizeof...(Dims) + 1>{Dim0, Dims...}[Dim];
   }
 };
 
-template <> struct property_value<max_linear_work_group_size_key> {
-  using key_t = max_linear_work_group_size_key;
-};
+template <>
+struct property_value<max_linear_work_group_size_key>
+    : detail::property_base<property_value<max_linear_work_group_size_key>,
+                            detail::PropKind::MaxLinearWorkGroupSize,
+                            max_linear_work_group_size_key> {};
 
 template <size_t Dim0, size_t... Dims>
 inline constexpr work_group_size_key::value_t<Dim0, Dims...> work_group_size;
@@ -235,8 +259,13 @@ template <forward_progress_guarantee Guarantee,
 struct property_value<
     work_group_progress_key,
     std::integral_constant<forward_progress_guarantee, Guarantee>,
-    std::integral_constant<execution_scope, CoordinationScope>> {
-  using key_t = work_group_progress_key;
+    std::integral_constant<execution_scope, CoordinationScope>>
+    : detail::property_base<
+          property_value<
+              work_group_progress_key,
+              std::integral_constant<forward_progress_guarantee, Guarantee>,
+              std::integral_constant<execution_scope, CoordinationScope>>,
+          detail::PropKind::WorkGroupProgress, work_group_progress_key> {
   static constexpr forward_progress_guarantee guarantee = Guarantee;
   static constexpr execution_scope coordinationScope = CoordinationScope;
 };
@@ -246,8 +275,13 @@ template <forward_progress_guarantee Guarantee,
 struct property_value<
     sub_group_progress_key,
     std::integral_constant<forward_progress_guarantee, Guarantee>,
-    std::integral_constant<execution_scope, CoordinationScope>> {
-  using key_t = work_group_progress_key;
+    std::integral_constant<execution_scope, CoordinationScope>>
+    : detail::property_base<
+          property_value<
+              sub_group_progress_key,
+              std::integral_constant<forward_progress_guarantee, Guarantee>,
+              std::integral_constant<execution_scope, CoordinationScope>>,
+          detail::PropKind::SubGroupProgress, sub_group_progress_key> {
   static constexpr forward_progress_guarantee guarantee = Guarantee;
   static constexpr execution_scope coordinationScope = CoordinationScope;
 };
@@ -257,8 +291,13 @@ template <forward_progress_guarantee Guarantee,
 struct property_value<
     work_item_progress_key,
     std::integral_constant<forward_progress_guarantee, Guarantee>,
-    std::integral_constant<execution_scope, CoordinationScope>> {
-  using key_t = work_group_progress_key;
+    std::integral_constant<execution_scope, CoordinationScope>>
+    : detail::property_base<
+          property_value<
+              work_item_progress_key,
+              std::integral_constant<forward_progress_guarantee, Guarantee>,
+              std::integral_constant<execution_scope, CoordinationScope>>,
+          detail::PropKind::WorkItemProgress, work_item_progress_key> {
   static constexpr forward_progress_guarantee guarantee = Guarantee;
   static constexpr execution_scope coordinationScope = CoordinationScope;
 };

sycl/include/sycl/ext/oneapi/latency_control/properties.hpp

@@ -58,8 +58,15 @@ struct property_value<
     intel::experimental::latency_constraint_key,
     std::integral_constant<int, Target>,
     std::integral_constant<intel::experimental::latency_control_type, Type>,
-    std::integral_constant<int, Cycle>> {
-  using key_t = intel::experimental::latency_constraint_key;
+    std::integral_constant<int, Cycle>>
+    : detail::property_base<
+          property_value<intel::experimental::latency_constraint_key,
+                         std::integral_constant<int, Target>,
+                         std::integral_constant<
+                             intel::experimental::latency_control_type, Type>,
+                         std::integral_constant<int, Cycle>>,
+          oneapi::experimental::detail::PropKind::LatencyConstraint,
+          intel::experimental::latency_constraint_key> {
   static constexpr int target = Target;
   static constexpr intel::experimental::latency_control_type type = Type;
   static constexpr int cycle = Cycle;

sycl/include/sycl/ext/oneapi/properties/properties.hpp

@@ -200,6 +200,10 @@ template <typename PropertiesT> class properties {
     static_assert(NumContainedProps == sizeof...(PropertyValueTs),
                   "One or more property argument is not a property in the "
                   "property list.");
+    // We're in process of refactoring properties infrastructure, make sure that
+    // any newly added properties use `detail::property_base`!
+    static_assert(
+        (std::is_base_of_v<detail::property_tag, PropertyValueTs> && ...));
   }
 
   template <typename PropertyT>

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -94,7 +94,7 @@ enum PropKind : uint32_t {
 namespace sycl::ext::oneapi::experimental {
 
 // (2.)
-struct foo : detail::run_time_property_key<PropKind::Foo> {
+struct foo : detail::run_time_property_key<foo, PropKind::Foo> {
   foo(int v) : value(v) {}
   int value;
 };
@@ -215,10 +215,35 @@ enum PropKind : uint32_t {
   PropKindSize = 79,
 };
 
+template <typename PropertyT> struct PropertyToKind {
+  static constexpr PropKind Kind = PropertyT::Kind;
+};
+
+struct property_tag {};
+template <typename property_t, PropKind Kind,
+          typename property_key_t = property_t>
+struct property_base : property_tag {
+  using key_t = property_key_t;
+#if !defined(_MSC_VER)
+  // Temporary, to ensure new code matches previous behavior and to catch any
+  // silly copy-paste mistakes. MSVC can't compile it, but linux-only is enough
+  // for this temporary check.
+  static_assert([]() constexpr {
+    if constexpr (std::is_same_v<property_t, key_t>)
+      // key_t is incomplete at this point for runtime properties.
+      return true;
+    else
+      return Kind == PropertyToKind<key_t>::Kind;
+  }());
+#endif
+};
+
 struct property_key_base_tag {};
 struct compile_time_property_key_base_tag : property_key_base_tag {};
 
-template <PropKind Kind_> struct run_time_property_key : property_key_base_tag {
+template <typename property_t, PropKind Kind_>
+struct run_time_property_key : property_key_base_tag,
+                               property_base<property_t, Kind_> {
 protected:
   static constexpr PropKind Kind = Kind_;
 
@@ -235,12 +260,6 @@ struct compile_time_property_key : compile_time_property_key_base_tag {
   friend struct PropertyToKind;
 };
 
-// This trait must be specialized for all properties and must have a unique
-// constexpr PropKind member named Kind.
-template <typename PropertyT> struct PropertyToKind {
-  static constexpr PropKind Kind = PropertyT::Kind;
-};
-
 // Get unique ID for property.
 template <typename PropertyT> struct PropertyID {
   static constexpr int value =

sycl/include/sycl/ext/oneapi/properties/property_value.hpp

@@ -41,9 +41,11 @@ struct PropertyValueBase<T> : public detail::SingleNontypePropertyValueBase<T> {
 } // namespace detail
 
 template <typename PropertyT, typename... Ts>
-struct property_value : public detail::PropertyValueBase<Ts...> {
-  using key_t = PropertyT;
-};
+struct property_value
+    : public detail::PropertyValueBase<Ts...>,
+      public detail::property_base<property_value<PropertyT, Ts...>,
+                                   detail::PropertyToKind<PropertyT>::Kind,
+                                   PropertyT> {};
 
 template <typename PropertyT, typename... A, typename... B>
 constexpr std::enable_if_t<detail::IsCompileTimeProperty<PropertyT>::value,