From dc06285c9f6ec42026adc861e0c9d2c1398e0bd5 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Mon, 26 Feb 2024 15:49:06 -0800
Subject: [PATCH 1/5] [NFCI][SYCL] Simplify property_key creation

---
 .../include/sycl/detail/kernel_properties.hpp |  16 +--
 .../ext/intel/esimd/memory_properties.hpp     |  40 ++----
 .../experimental/cache_control_properties.hpp |  37 ++----
 .../fp_control_kernel_properties.hpp          |  15 +--
 .../fpga_annotated_properties.hpp             |  91 ++++----------
 .../experimental/fpga_kernel_properties.hpp   |  42 ++-----
 .../experimental/fpga_mem/properties.hpp      | 115 ++++--------------
 .../experimental/grf_size_properties.hpp      |  34 +-----
 .../kernel_execution_properties.hpp           |  17 +--
 .../intel/experimental/pipe_properties.hpp    |  84 +++----------
 .../ext/oneapi/device_global/properties.hpp   |  37 ++----
 .../annotated_ptr_properties.hpp              |  12 +-
 .../properties.hpp                            |  12 +-
 .../ext/oneapi/experimental/root_group.hpp    |  12 +-
 .../oneapi/kernel_properties/properties.hpp   |  38 ++----
 .../ext/oneapi/latency_control/properties.hpp |  32 +----
 .../sycl/ext/oneapi/properties/property.hpp   |  70 +++++------
 sycl/include/sycl/kernel_bundle.hpp           |  43 +------
 .../mock_compile_time_properties.hpp          |  64 +++-------
 19 files changed, 183 insertions(+), 628 deletions(-)

diff --git a/sycl/include/sycl/detail/kernel_properties.hpp b/sycl/include/sycl/detail/kernel_properties.hpp
index 581a2c8c5dfbf..61ecbf583d945 100644
--- a/sycl/include/sycl/detail/kernel_properties.hpp
+++ b/sycl/include/sycl/detail/kernel_properties.hpp
@@ -21,7 +21,9 @@ enum class register_alloc_mode_enum : uint32_t {
   large = 2,
 };
 
-struct register_alloc_mode_key {
+struct register_alloc_mode_key
+    : ext::oneapi::experimental::detail::compile_time_property_key<
+          ext::oneapi::experimental::detail::PropKind::RegisterAllocMode> {
   template <register_alloc_mode_enum Mode>
   using value_t = sycl::ext::oneapi::experimental::property_value<
       register_alloc_mode_key,
@@ -36,19 +38,7 @@ inline constexpr register_alloc_mode_key::value_t<Mode> register_alloc_mode
 } // namespace detail
 
 namespace ext::oneapi::experimental {
-template <>
-struct is_property_key<sycl::detail::register_alloc_mode_key> : std::true_type {
-};
-
 namespace detail {
-template <> struct PropertyToKind<sycl::detail::register_alloc_mode_key> {
-  static constexpr PropKind Kind = PropKind::RegisterAllocMode;
-};
-
-template <>
-struct IsCompileTimeProperty<sycl::detail::register_alloc_mode_key>
-    : std::true_type {};
-
 template <sycl::detail::register_alloc_mode_enum Mode>
 struct PropertyMetaInfo<sycl::detail::register_alloc_mode_key::value_t<Mode>> {
   static constexpr const char *name = "sycl-register-alloc-mode";
diff --git a/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp b/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp
index e2c9f17182fc1..4da1a64225bed 100644
--- a/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp
+++ b/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp
@@ -59,17 +59,23 @@ template <int K> inline constexpr alignment_key::value_t<K> alignment;
 /// L2 cache hint property must be used for the old/experimental LSC L3 cache
 /// hints.
 /// L3 cache property is reserved for future devices.
-struct cache_hint_L1_key {
+struct cache_hint_L1_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::ESIMDL1CacheHint> {
   template <cache_hint Hint>
   using value_t = ext::oneapi::experimental::property_value<
       cache_hint_L1_key, std::integral_constant<cache_hint, Hint>>;
 };
-struct cache_hint_L2_key {
+struct cache_hint_L2_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::ESIMDL2CacheHint> {
   template <cache_hint Hint>
   using value_t = ext::oneapi::experimental::property_value<
       cache_hint_L2_key, std::integral_constant<cache_hint, Hint>>;
 };
-struct cache_hint_L3_key {
+struct cache_hint_L3_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::ESIMDL3CacheHint> {
   template <cache_hint Hint>
   using value_t = ext::oneapi::experimental::property_value<
       cache_hint_L3_key, std::integral_constant<cache_hint, Hint>>;
@@ -151,45 +157,17 @@ struct property_value<__ESIMD_NS::cache_hint_L3_key,
   static constexpr __ESIMD_NS::cache_hint hint = Hint;
 };
 
-template <>
-struct is_property_key<sycl::ext::intel::esimd::cache_hint_L1_key>
-    : std::true_type {};
-template <>
-struct is_property_key<sycl::ext::intel::esimd::cache_hint_L2_key>
-    : std::true_type {};
-template <>
-struct is_property_key<sycl::ext::intel::esimd::cache_hint_L3_key>
-    : std::true_type {};
-
 // Declare that esimd::properties is a property_list.
 template <typename... PropertyValueTs>
 struct is_property_list<__ESIMD_NS::properties<std::tuple<PropertyValueTs...>>>
     : is_property_list<properties<std::tuple<PropertyValueTs...>>> {};
 
 namespace detail {
-template <> struct PropertyToKind<sycl::ext::intel::esimd::cache_hint_L1_key> {
-  static constexpr PropKind Kind = PropKind::ESIMDL1CacheHint;
-};
-template <> struct PropertyToKind<sycl::ext::intel::esimd::cache_hint_L2_key> {
-  static constexpr PropKind Kind = PropKind::ESIMDL2CacheHint;
-};
-template <> struct PropertyToKind<sycl::ext::intel::esimd::cache_hint_L3_key> {
-  static constexpr PropKind Kind = PropKind::ESIMDL3CacheHint;
-};
-
-template <>
-struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L1_key> : std::true_type {};
-template <>
-struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L2_key> : std::true_type {};
-template <>
-struct IsCompileTimeProperty<__ESIMD_NS::cache_hint_L3_key> : std::true_type {};
-
 // We do not override the class ConflictingProperties for cache_hint properties
 // because that mechanism would only allow to verify few obvious restrictions
 // without the knowledge of the context in which the cache_hint properties are
 // used (load, store, prefetch, atomic). Thus the function
 // __ESIMD_DNS::check_cache_hint() is used to verify correctness of properties.
-
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1
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 4b2e4869f972c..ea8ec1020d54a 100755
--- a/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/cache_control_properties.hpp
@@ -72,17 +72,23 @@ template <typename PropertyT, typename... Ts>
 using property_value =
     sycl::ext::oneapi::experimental::property_value<PropertyT, Ts...>;
 
-struct read_hint_key {
+struct read_hint_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::CacheControlReadHint> {
   template <typename... Cs>
   using value_t = property_value<read_hint_key, Cs...>;
 };
 
-struct read_assertion_key {
+struct read_assertion_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::CacheControlReadAssertion> {
   template <typename... Cs>
   using value_t = property_value<read_assertion_key, Cs...>;
 };
 
-struct write_hint_key {
+struct write_hint_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::CacheControlWrite> {
   template <typename... Cs>
   using value_t = property_value<write_hint_key, Cs...>;
 };
@@ -104,21 +110,14 @@ namespace experimental {
 
 template <typename T, typename PropertyListT> class annotated_ptr;
 
-template <>
-struct is_property_key<intel::experimental::read_hint_key> : std::true_type {};
 template <typename T, typename PropertyListT>
 struct is_property_key_of<intel::experimental::read_hint_key,
                           annotated_ptr<T, PropertyListT>> : std::true_type {};
 
-template <>
-struct is_property_key<intel::experimental::read_assertion_key>
-    : std::true_type {};
 template <typename T, typename PropertyListT>
 struct is_property_key_of<intel::experimental::read_assertion_key,
                           annotated_ptr<T, PropertyListT>> : std::true_type {};
 
-template <>
-struct is_property_key<intel::experimental::write_hint_key> : std::true_type {};
 template <typename T, typename PropertyListT>
 struct is_property_key_of<intel::experimental::write_hint_key,
                           annotated_ptr<T, PropertyListT>> : std::true_type {};
@@ -178,12 +177,6 @@ template <cache_mode M> static constexpr int checkWriteHint() {
   return 0;
 }
 
-template <> struct PropertyToKind<intel::experimental::read_hint_key> {
-  static constexpr PropKind Kind = PropKind::CacheControlReadHint;
-};
-template <>
-struct IsCompileTimeProperty<intel::experimental::read_hint_key>
-    : std::true_type {};
 template <typename... Cs>
 struct PropertyMetaInfo<intel::experimental::read_hint_key::value_t<Cs...>> {
   static constexpr const char *name = "sycl-cache-read-hint";
@@ -196,12 +189,6 @@ struct PropertyMetaInfo<intel::experimental::read_hint_key::value_t<Cs...>> {
        ((Cs::encoding) | ...));
 };
 
-template <> struct PropertyToKind<intel::experimental::read_assertion_key> {
-  static constexpr PropKind Kind = PropKind::CacheControlReadAssertion;
-};
-template <>
-struct IsCompileTimeProperty<intel::experimental::read_assertion_key>
-    : std::true_type {};
 template <typename... Cs>
 struct PropertyMetaInfo<
     intel::experimental::read_assertion_key::value_t<Cs...>> {
@@ -215,12 +202,6 @@ struct PropertyMetaInfo<
        ((Cs::encoding) | ...));
 };
 
-template <> struct PropertyToKind<intel::experimental::write_hint_key> {
-  static constexpr PropKind Kind = PropKind::CacheControlWrite;
-};
-template <>
-struct IsCompileTimeProperty<intel::experimental::write_hint_key>
-    : std::true_type {};
 template <typename... Cs>
 struct PropertyMetaInfo<intel::experimental::write_hint_key::value_t<Cs...>> {
   static constexpr const char *name = "sycl-cache-write-hint";
diff --git a/sycl/include/sycl/ext/intel/experimental/fp_control_kernel_properties.hpp b/sycl/include/sycl/ext/intel/experimental/fp_control_kernel_properties.hpp
index 069a6fc16fb5e..544535d098e1a 100644
--- a/sycl/include/sycl/ext/intel/experimental/fp_control_kernel_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fp_control_kernel_properties.hpp
@@ -85,7 +85,9 @@ constexpr fp_mode setDefaultValuesIfNeeded(fp_mode mode) {
 }
 } // namespace detail
 
-struct fp_control_key {
+struct fp_control_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::FloatingPointControls> {
   template <fp_mode option>
   using value_t = ext::oneapi::experimental::property_value<
       fp_control_key, std::integral_constant<fp_mode, option>>;
@@ -97,9 +99,6 @@ inline constexpr fp_control_key::value_t<option> fp_control;
 } // namespace ext::intel::experimental
 
 namespace ext::oneapi::experimental {
-template <>
-struct is_property_key<intel::experimental::fp_control_key> : std::true_type {};
-
 template <typename T, typename PropertyListT>
 struct is_property_key_of<
     intel::experimental::fp_control_key,
@@ -107,14 +106,6 @@ struct is_property_key_of<
 };
 
 namespace detail {
-template <> struct PropertyToKind<intel::experimental::fp_control_key> {
-  static constexpr PropKind Kind = FloatingPointControls;
-};
-
-template <>
-struct IsCompileTimeProperty<intel::experimental::fp_control_key>
-    : std::true_type {};
-
 template <intel::experimental::fp_mode FPMode>
 struct PropertyMetaInfo<intel::experimental::fp_control_key::value_t<FPMode>> {
   static_assert(intel::experimental::detail::checkMutuallyExclusive(FPMode),
diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp
index 7bc04324d6b0e..ea9f531c764b9 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp
@@ -31,54 +31,68 @@ using property_value =
 //===----------------------------------------------------------------------===//
 //        FPGA properties of annotated_arg/annotated_ptr
 //===----------------------------------------------------------------------===//
-struct register_map_key {
+struct register_map_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::RegisterMap> {
   using value_t = property_value<register_map_key>;
 };
 
-struct conduit_key {
+struct conduit_key : oneapi::experimental::detail::compile_time_property_key<
+                         oneapi::experimental::detail::PropKind::Conduit> {
   using value_t = property_value<conduit_key>;
 };
 
-struct stable_key {
+struct stable_key : oneapi::experimental::detail::compile_time_property_key<
+                        oneapi::experimental::detail::PropKind::Stable> {
   using value_t = property_value<stable_key>;
 };
 
-struct buffer_location_key {
+struct buffer_location_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::BufferLocation> {
   template <int K>
   using value_t =
       property_value<buffer_location_key, std::integral_constant<int, K>>;
 };
 
-struct awidth_key {
+struct awidth_key : oneapi::experimental::detail::compile_time_property_key<
+                        oneapi::experimental::detail::PropKind::AddrWidth> {
   template <int K>
   using value_t = property_value<awidth_key, std::integral_constant<int, K>>;
 };
 
-struct dwidth_key {
+struct dwidth_key : oneapi::experimental::detail::compile_time_property_key<
+                        oneapi::experimental::detail::PropKind::DataWidth> {
   template <int K>
   using value_t = property_value<dwidth_key, std::integral_constant<int, K>>;
 };
 
-struct latency_key {
+struct latency_key : oneapi::experimental::detail::compile_time_property_key<
+                         oneapi::experimental::detail::PropKind::Latency> {
   template <int K>
   using value_t = property_value<latency_key, std::integral_constant<int, K>>;
 };
 
 enum class read_write_mode_enum : std::uint16_t { read, write, read_write };
 
-struct read_write_mode_key {
+struct read_write_mode_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::RWMode> {
   template <read_write_mode_enum Mode>
   using value_t =
       property_value<read_write_mode_key,
                      std::integral_constant<read_write_mode_enum, Mode>>;
 };
 
-struct maxburst_key {
+struct maxburst_key : oneapi::experimental::detail::compile_time_property_key<
+                          oneapi::experimental::detail::PropKind::MaxBurst> {
   template <int K>
   using value_t = property_value<maxburst_key, std::integral_constant<int, K>>;
 };
 
-struct wait_request_key {
+struct wait_request_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::WaitRequest> {
   template <int K>
   using value_t =
       property_value<wait_request_key, std::integral_constant<int, K>>;
@@ -131,18 +145,6 @@ using maxburst_key = intel::experimental::maxburst_key;
 using wait_request_key = intel::experimental::wait_request_key;
 using read_write_mode_enum = intel::experimental::read_write_mode_enum;
 
-template <> struct is_property_key<register_map_key> : std::true_type {};
-template <> struct is_property_key<conduit_key> : std::true_type {};
-template <> struct is_property_key<stable_key> : std::true_type {};
-
-template <> struct is_property_key<buffer_location_key> : std::true_type {};
-template <> struct is_property_key<awidth_key> : std::true_type {};
-template <> struct is_property_key<dwidth_key> : std::true_type {};
-template <> struct is_property_key<latency_key> : std::true_type {};
-template <> struct is_property_key<read_write_mode_key> : std::true_type {};
-template <> struct is_property_key<maxburst_key> : std::true_type {};
-template <> struct is_property_key<wait_request_key> : std::true_type {};
-
 template <typename T, typename PropertyListT>
 struct is_property_key_of<register_map_key, annotated_arg<T, PropertyListT>>
     : std::true_type {};
@@ -224,51 +226,6 @@ struct is_property_key_of<wait_request_key, annotated_ptr<T, PropertyListT>>
     : std::true_type {};
 
 namespace detail {
-template <> struct PropertyToKind<register_map_key> {
-  static constexpr PropKind Kind = PropKind::RegisterMap;
-};
-template <> struct PropertyToKind<conduit_key> {
-  static constexpr PropKind Kind = PropKind::Conduit;
-};
-template <> struct PropertyToKind<stable_key> {
-  static constexpr PropKind Kind = PropKind::Stable;
-};
-template <> struct PropertyToKind<buffer_location_key> {
-  static constexpr PropKind Kind = PropKind::BufferLocation;
-};
-template <> struct PropertyToKind<awidth_key> {
-  static constexpr PropKind Kind = PropKind::AddrWidth;
-};
-template <> struct PropertyToKind<dwidth_key> {
-  static constexpr PropKind Kind = PropKind::DataWidth;
-};
-template <> struct PropertyToKind<latency_key> {
-  static constexpr PropKind Kind = PropKind::Latency;
-};
-template <> struct PropertyToKind<read_write_mode_key> {
-  static constexpr PropKind Kind = PropKind::RWMode;
-};
-template <> struct PropertyToKind<maxburst_key> {
-  static constexpr PropKind Kind = PropKind::MaxBurst;
-};
-template <> struct PropertyToKind<wait_request_key> {
-  static constexpr PropKind Kind = PropKind::WaitRequest;
-};
-
-template <> struct IsCompileTimeProperty<register_map_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<conduit_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<stable_key> : std::true_type {};
-
-template <>
-struct IsCompileTimeProperty<buffer_location_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<awidth_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<dwidth_key> : std::true_type {};
-template <>
-struct IsCompileTimeProperty<read_write_mode_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<latency_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<maxburst_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<wait_request_key> : std::true_type {};
-
 template <> struct PropertyMetaInfo<register_map_key::value_t> {
   static constexpr const char *name = "sycl-register-map";
   static constexpr std::nullptr_t value = nullptr;
diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
index 1e5b8a7790c2f..0db22e9407a5e 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
@@ -28,21 +28,26 @@ enum class register_map_interface_options_enum : std::uint16_t {
   wait_for_done_write,
 };
 
-struct streaming_interface_key {
+struct streaming_interface_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::StreamingInterface> {
   template <streaming_interface_options_enum option>
   using value_t = ext::oneapi::experimental::property_value<
       streaming_interface_key,
       std::integral_constant<streaming_interface_options_enum, option>>;
 };
 
-struct register_map_interface_key {
+struct register_map_interface_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::RegisterMapInterface> {
   template <register_map_interface_options_enum option>
   using value_t = ext::oneapi::experimental::property_value<
       register_map_interface_key,
       std::integral_constant<register_map_interface_options_enum, option>>;
 };
 
-struct pipelined_key {
+struct pipelined_key : oneapi::experimental::detail::compile_time_property_key<
+                           oneapi::experimental::detail::PropKind::Pipelined> {
   template <int pipeline_directive_or_initiation_interval>
   using value_t = ext::oneapi::experimental::property_value<
       pipelined_key,
@@ -82,15 +87,6 @@ inline constexpr pipelined_key::value_t<
 } // namespace ext::intel::experimental
 
 namespace ext::oneapi::experimental {
-template <>
-struct is_property_key<intel::experimental::streaming_interface_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::register_map_interface_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::pipelined_key> : std::true_type {};
-
 template <typename T, typename PropertyListT>
 struct is_property_key_of<
     intel::experimental::streaming_interface_key,
@@ -108,28 +104,6 @@ struct is_property_key_of<
     : std::true_type {};
 
 namespace detail {
-template <>
-struct PropertyToKind<intel::experimental::streaming_interface_key> {
-  static constexpr PropKind Kind = StreamingInterface;
-};
-template <>
-struct PropertyToKind<intel::experimental::register_map_interface_key> {
-  static constexpr PropKind Kind = RegisterMapInterface;
-};
-template <> struct PropertyToKind<intel::experimental::pipelined_key> {
-  static constexpr PropKind Kind = Pipelined;
-};
-
-template <>
-struct IsCompileTimeProperty<intel::experimental::streaming_interface_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::register_map_interface_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::pipelined_key>
-    : std::true_type {};
-
 template <intel::experimental::streaming_interface_options_enum Stall_Free>
 struct PropertyMetaInfo<
     intel::experimental::streaming_interface_key::value_t<Stall_Free>> {
diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_mem/properties.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_mem/properties.hpp
index 516439a467666..95273529c8073 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_mem/properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_mem/properties.hpp
@@ -32,58 +32,72 @@ using property_value =
 // Property definitions
 enum class resource_enum : std::uint16_t { mlab, block_ram };
 
-struct resource_key {
+struct resource_key : oneapi::experimental::detail::compile_time_property_key<
+                          oneapi::experimental::detail::PropKind::Resource> {
   template <resource_enum Resource>
   using value_t =
       property_value<resource_key,
                      std::integral_constant<resource_enum, Resource>>;
 };
 
-struct num_banks_key {
+struct num_banks_key : oneapi::experimental::detail::compile_time_property_key<
+                           oneapi::experimental::detail::PropKind::NumBanks> {
   template <size_t Elements>
   using value_t =
       property_value<num_banks_key, std::integral_constant<size_t, Elements>>;
 };
 
-struct stride_size_key {
+struct stride_size_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::StrideSize> {
   template <size_t Elements>
   using value_t =
       property_value<stride_size_key, std::integral_constant<size_t, Elements>>;
 };
 
-struct word_size_key {
+struct word_size_key : oneapi::experimental::detail::compile_time_property_key<
+                           oneapi::experimental::detail::PropKind::WordSize> {
   template <size_t Elements>
   using value_t =
       property_value<word_size_key, std::integral_constant<size_t, Elements>>;
 };
 
-struct bi_directional_ports_key {
+struct bi_directional_ports_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::BiDirectionalPorts> {
   template <bool Enable>
   using value_t =
       property_value<bi_directional_ports_key, std::bool_constant<Enable>>;
 };
 
-struct clock_2x_key {
+struct clock_2x_key : oneapi::experimental::detail::compile_time_property_key<
+                          oneapi::experimental::detail::PropKind::Clock2x> {
   template <bool Enable>
   using value_t = property_value<clock_2x_key, std::bool_constant<Enable>>;
 };
 
 enum class ram_stitching_enum : std::uint16_t { min_ram, max_fmax };
 
-struct ram_stitching_key {
+struct ram_stitching_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::RAMStitching> {
   template <ram_stitching_enum RamStitching>
   using value_t =
       property_value<ram_stitching_key,
                      std::integral_constant<ram_stitching_enum, RamStitching>>;
 };
 
-struct max_private_copies_key {
+struct max_private_copies_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::MaxPrivateCopies> {
   template <size_t N>
   using value_t =
       property_value<max_private_copies_key, std::integral_constant<size_t, N>>;
 };
 
-struct num_replicates_key {
+struct num_replicates_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::NumReplicates> {
   template <size_t N>
   using value_t =
       property_value<num_replicates_key, std::integral_constant<size_t, N>>;
@@ -129,30 +143,6 @@ inline constexpr num_replicates_key::value_t<N> num_replicates;
 
 namespace oneapi::experimental {
 
-template <>
-struct is_property_key<intel::experimental::resource_key> : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::num_banks_key> : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::stride_size_key> : std::true_type {
-};
-template <>
-struct is_property_key<intel::experimental::word_size_key> : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::bi_directional_ports_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::clock_2x_key> : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::ram_stitching_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::max_private_copies_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::num_replicates_key>
-    : std::true_type {};
-
 // Associate properties with fpga_mem
 template <typename T, typename PropertyListT>
 struct is_property_key_of<intel::experimental::resource_key,
@@ -192,65 +182,6 @@ struct is_property_key_of<intel::experimental::num_replicates_key,
     : std::true_type {};
 
 namespace detail {
-// Map Property to a PropKind enum
-template <> struct PropertyToKind<intel::experimental::resource_key> {
-  static constexpr PropKind Kind = PropKind::Resource;
-};
-template <> struct PropertyToKind<intel::experimental::num_banks_key> {
-  static constexpr PropKind Kind = PropKind::NumBanks;
-};
-template <> struct PropertyToKind<intel::experimental::stride_size_key> {
-  static constexpr PropKind Kind = PropKind::StrideSize;
-};
-template <> struct PropertyToKind<intel::experimental::word_size_key> {
-  static constexpr PropKind Kind = PropKind::WordSize;
-};
-template <>
-struct PropertyToKind<intel::experimental::bi_directional_ports_key> {
-  static constexpr PropKind Kind = PropKind::BiDirectionalPorts;
-};
-template <> struct PropertyToKind<intel::experimental::clock_2x_key> {
-  static constexpr PropKind Kind = PropKind::Clock2x;
-};
-template <> struct PropertyToKind<intel::experimental::ram_stitching_key> {
-  static constexpr PropKind Kind = PropKind::RAMStitching;
-};
-template <> struct PropertyToKind<intel::experimental::max_private_copies_key> {
-  static constexpr PropKind Kind = PropKind::MaxPrivateCopies;
-};
-template <> struct PropertyToKind<intel::experimental::num_replicates_key> {
-  static constexpr PropKind Kind = PropKind::NumReplicates;
-};
-
-// Mark the properties as compile-time
-template <>
-struct IsCompileTimeProperty<intel::experimental::resource_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::num_banks_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::stride_size_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::word_size_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::bi_directional_ports_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::clock_2x_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::ram_stitching_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::max_private_copies_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::num_replicates_key>
-    : std::true_type {};
-
 // Map Property to MetaInfo
 template <intel::experimental::resource_enum Value>
 struct PropertyMetaInfo<intel::experimental::resource_key::value_t<Value>> {
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 b50b445e9baff..a7c60aae50ad2 100644
--- a/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/grf_size_properties.hpp
@@ -17,13 +17,16 @@
 namespace sycl {
 inline namespace _V1 {
 namespace ext::intel::experimental {
-struct grf_size_key {
+struct grf_size_key : oneapi::experimental::detail::compile_time_property_key<
+                          oneapi::experimental::detail::PropKind::GRFSize> {
   template <unsigned int Size>
   using value_t = oneapi::experimental::property_value<
       grf_size_key, std::integral_constant<unsigned int, Size>>;
 };
 
-struct grf_size_automatic_key {
+struct grf_size_automatic_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::GRFSizeAutomatic> {
   using value_t = oneapi::experimental::property_value<grf_size_automatic_key>;
 };
 
@@ -34,33 +37,6 @@ inline constexpr grf_size_automatic_key::value_t grf_size_automatic;
 
 } // namespace ext::intel::experimental
 namespace ext::oneapi::experimental {
-template <>
-struct is_property_key<sycl::ext::intel::experimental::grf_size_key>
-    : std::true_type {};
-
-template <>
-struct is_property_key<sycl::ext::intel::experimental::grf_size_automatic_key>
-    : std::true_type {};
-
-namespace detail {
-template <>
-struct PropertyToKind<sycl::ext::intel::experimental::grf_size_key> {
-  static constexpr PropKind Kind = PropKind::GRFSize;
-};
-
-template <>
-struct IsCompileTimeProperty<sycl::ext::intel::experimental::grf_size_key>
-    : std::true_type {};
-
-template <>
-struct PropertyToKind<sycl::ext::intel::experimental::grf_size_automatic_key> {
-  static constexpr PropKind Kind = PropKind::GRFSizeAutomatic;
-};
-
-template <>
-struct IsCompileTimeProperty<
-    sycl::ext::intel::experimental::grf_size_automatic_key> : std::true_type {};
-
 template <unsigned int Size>
 struct PropertyMetaInfo<
     sycl::ext::intel::experimental::grf_size_key::value_t<Size>> {
diff --git a/sycl/include/sycl/ext/intel/experimental/kernel_execution_properties.hpp b/sycl/include/sycl/ext/intel/experimental/kernel_execution_properties.hpp
index 42e8bdc88c035..e65a39c072110 100644
--- a/sycl/include/sycl/ext/intel/experimental/kernel_execution_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/kernel_execution_properties.hpp
@@ -26,7 +26,8 @@ inline constexpr cache_config_enum large_slm =
 inline constexpr cache_config_enum large_data =
     cache_config_enum::large_data;
 
-struct cache_config {
+struct cache_config : oneapi::experimental::detail::run_time_property_key<
+                          oneapi::experimental::detail::PropKind::CacheConfig> {
   cache_config(cache_config_enum v) : value(v) {}
   cache_config_enum value;
 };
@@ -45,26 +46,12 @@ inline bool operator!=(const cache_config &lhs,
 } // namespace ext::intel::experimental
 
 namespace ext::oneapi::experimental {
-template <>
-struct is_property_key<intel::experimental::cache_config_key>
-    : std::true_type {};
-
 template <typename T, typename PropertyListT>
 struct is_property_key_of<
     intel::experimental::cache_config_key,
     intel::experimental::gpu_kernel_attribute<T, PropertyListT>>
     : std::true_type {};
 
-namespace detail {
-template <> struct PropertyToKind<intel::experimental::cache_config_key> {
-  static constexpr PropKind Kind = PropKind::CacheConfig;
-};
-
-template <>
-struct IsRuntimeProperty<intel::experimental::cache_config_key>
-    : std::true_type {};
-
-} // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp
index 95685881da18c..a1ac195d8af67 100644
--- a/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp
@@ -17,31 +17,36 @@
 
 namespace sycl {
 inline namespace _V1 {
-namespace ext {
-namespace intel {
-namespace experimental {
+namespace ext::intel::experimental {
 
-struct ready_latency_key {
+struct ready_latency_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::ReadyLatency> {
   template <int Latency>
   using value_t = oneapi::experimental::property_value<
       ready_latency_key, std::integral_constant<int, Latency>>;
 };
 
-struct bits_per_symbol_key {
+struct bits_per_symbol_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::BitsPerSymbol> {
   template <int Bits>
   using value_t =
       oneapi::experimental::property_value<bits_per_symbol_key,
                                            std::integral_constant<int, Bits>>;
 };
 
-struct uses_valid_key {
+struct uses_valid_key : oneapi::experimental::detail::compile_time_property_key<
+                            oneapi::experimental::detail::PropKind::UsesValid> {
   template <bool Valid>
   using value_t =
       oneapi::experimental::property_value<uses_valid_key,
                                            std::bool_constant<Valid>>;
 };
 
-struct first_symbol_in_high_order_bits_key {
+struct first_symbol_in_high_order_bits_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::FirstSymbolInHigherOrderBit> {
   template <bool HighOrder>
   using value_t =
       oneapi::experimental::property_value<first_symbol_in_high_order_bits_key,
@@ -55,7 +60,9 @@ enum class protocol_name : std::uint16_t {
   avalon_mm_uses_ready = 3
 };
 
-struct protocol_key {
+struct protocol_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::PipeProtocol> {
   template <protocol_name Protocol>
   using value_t = oneapi::experimental::property_value<
       protocol_key, std::integral_constant<protocol_name, Protocol>>;
@@ -92,65 +99,6 @@ inline constexpr protocol_key::value_t<protocol_name::avalon_mm>
 inline constexpr protocol_key::value_t<protocol_name::avalon_mm_uses_ready>
     protocol_avalon_mm_uses_ready;
 
-} // namespace experimental
-} // namespace intel
-
-namespace oneapi {
-namespace experimental {
-
-template <>
-struct is_property_key<intel::experimental::ready_latency_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::bits_per_symbol_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::uses_valid_key> : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::first_symbol_in_high_order_bits_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::protocol_key> : std::true_type {};
-
-namespace detail {
-template <> struct PropertyToKind<intel::experimental::ready_latency_key> {
-  static constexpr PropKind Kind = PropKind::ReadyLatency;
-};
-template <> struct PropertyToKind<intel::experimental::bits_per_symbol_key> {
-  static constexpr PropKind Kind = PropKind::BitsPerSymbol;
-};
-template <> struct PropertyToKind<intel::experimental::uses_valid_key> {
-  static constexpr PropKind Kind = PropKind::UsesValid;
-};
-template <>
-struct PropertyToKind<
-    intel::experimental::first_symbol_in_high_order_bits_key> {
-  static constexpr PropKind Kind = PropKind::FirstSymbolInHigherOrderBit;
-};
-template <> struct PropertyToKind<intel::experimental::protocol_key> {
-  static constexpr PropKind Kind = PropKind::PipeProtocol;
-};
-
-template <>
-struct IsCompileTimeProperty<intel::experimental::ready_latency_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::bits_per_symbol_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::uses_valid_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<
-    intel::experimental::first_symbol_in_high_order_bits_key> : std::true_type {
-};
-template <>
-struct IsCompileTimeProperty<intel::experimental::protocol_key>
-    : std::true_type {};
-
-} // namespace detail
-} // namespace experimental
-} // namespace oneapi
-} // namespace ext
+} // namespace ext::intel::experimental
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/device_global/properties.hpp b/sycl/include/sycl/ext/oneapi/device_global/properties.hpp
index 06fa051cdf463..23bad9af0d88f 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/properties.hpp
@@ -21,13 +21,15 @@ namespace ext::oneapi::experimental {
 
 template <typename T, typename PropertyListT> class device_global;
 
-struct device_image_scope_key {
+struct device_image_scope_key
+    : detail::compile_time_property_key<detail::PropKind::DeviceImageScope> {
   using value_t = property_value<device_image_scope_key>;
 };
 
 enum class host_access_enum : std::uint16_t { read, write, read_write, none };
 
-struct host_access_key {
+struct host_access_key
+    : detail::compile_time_property_key<detail::PropKind::HostAccess> {
   template <host_access_enum Access>
   using value_t =
       property_value<host_access_key,
@@ -36,14 +38,16 @@ struct host_access_key {
 
 enum class init_mode_enum : std::uint16_t { reprogram, reset };
 
-struct init_mode_key {
+struct init_mode_key
+    : detail::compile_time_property_key<detail::PropKind::InitMode> {
   template <init_mode_enum Trigger>
   using value_t =
       property_value<init_mode_key,
                      std::integral_constant<init_mode_enum, Trigger>>;
 };
 
-struct implement_in_csr_key {
+struct implement_in_csr_key
+    : detail::compile_time_property_key<detail::PropKind::ImplementInCSR> {
   template <bool Enable>
   using value_t =
       property_value<implement_in_csr_key, std::bool_constant<Enable>>;
@@ -73,11 +77,6 @@ inline constexpr implement_in_csr_key::value_t<Enable> implement_in_csr;
 inline constexpr implement_in_csr_key::value_t<true> implement_in_csr_on;
 inline constexpr implement_in_csr_key::value_t<false> implement_in_csr_off;
 
-template <> struct is_property_key<device_image_scope_key> : std::true_type {};
-template <> struct is_property_key<host_access_key> : std::true_type {};
-template <> struct is_property_key<init_mode_key> : std::true_type {};
-template <> struct is_property_key<implement_in_csr_key> : std::true_type {};
-
 template <typename T, typename PropertyListT>
 struct is_property_key_of<device_image_scope_key,
                           device_global<T, PropertyListT>> : std::true_type {};
@@ -92,26 +91,6 @@ struct is_property_key_of<implement_in_csr_key, device_global<T, PropertyListT>>
     : std::true_type {};
 
 namespace detail {
-template <> struct PropertyToKind<device_image_scope_key> {
-  static constexpr PropKind Kind = PropKind::DeviceImageScope;
-};
-template <> struct PropertyToKind<host_access_key> {
-  static constexpr PropKind Kind = PropKind::HostAccess;
-};
-template <> struct PropertyToKind<init_mode_key> {
-  static constexpr PropKind Kind = PropKind::InitMode;
-};
-template <> struct PropertyToKind<implement_in_csr_key> {
-  static constexpr PropKind Kind = PropKind::ImplementInCSR;
-};
-
-template <>
-struct IsCompileTimeProperty<device_image_scope_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<host_access_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<init_mode_key> : std::true_type {};
-template <>
-struct IsCompileTimeProperty<implement_in_csr_key> : std::true_type {};
-
 template <> struct PropertyMetaInfo<device_image_scope_key::value_t> {
   static constexpr const char *name = "sycl-device-image-scope";
   static constexpr std::nullptr_t value = nullptr;
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 8e126a282a68e..fcdd05f3a497e 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
@@ -24,7 +24,8 @@ namespace experimental {
 
 template <typename T, typename PropertyListT> class annotated_ptr;
 
-struct usm_kind_key {
+struct usm_kind_key
+    : detail::compile_time_property_key<detail::PropKind::UsmKind> {
   template <sycl::usm::alloc Kind>
   using value_t =
       property_value<usm_kind_key,
@@ -39,8 +40,6 @@ inline constexpr usm_kind_key::value_t<sycl::usm::alloc::host> usm_kind_host;
 inline constexpr usm_kind_key::value_t<sycl::usm::alloc::shared>
     usm_kind_shared;
 
-template <> struct is_property_key<usm_kind_key> : std::true_type {};
-
 template <typename T, sycl::usm::alloc Kind>
 struct is_valid_property<T, usm_kind_key::value_t<Kind>>
     : std::bool_constant<std::is_pointer<T>::value> {};
@@ -50,13 +49,6 @@ struct is_property_key_of<usm_kind_key, annotated_ptr<T, PropertyListT>>
     : std::true_type {};
 
 namespace detail {
-
-template <> struct PropertyToKind<usm_kind_key> {
-  static constexpr PropKind Kind = PropKind::UsmKind;
-};
-
-template <> struct IsCompileTimeProperty<usm_kind_key> : std::true_type {};
-
 template <sycl::usm::alloc Kind>
 struct PropertyMetaInfo<usm_kind_key::value_t<Kind>> {
   static constexpr const char *name = "sycl-usm-kind";
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 1a73da026df8f..dd540a14232b3 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
@@ -66,15 +66,14 @@ struct propagateToPtrAnnotation<property_value<PropKeyT, PropValuesTs...>>
 //===----------------------------------------------------------------------===//
 //        Common properties of annotated_arg/annotated_ptr
 //===----------------------------------------------------------------------===//
-struct alignment_key {
+struct alignment_key
+    : detail::compile_time_property_key<detail::PropKind::Alignment> {
   template <int K>
   using value_t = property_value<alignment_key, std::integral_constant<int, K>>;
 };
 
 template <int K> inline constexpr alignment_key::value_t<K> alignment;
 
-template <> struct is_property_key<alignment_key> : std::true_type {};
-
 template <typename T, int W>
 struct is_valid_property<T, alignment_key::value_t<W>>
     : std::bool_constant<std::is_pointer<T>::value> {};
@@ -90,13 +89,6 @@ struct is_property_key_of<alignment_key, annotated_arg<T, PropertyListT>>
 template <> struct propagateToPtrAnnotation<alignment_key> : std::true_type {};
 
 namespace detail {
-
-template <> struct PropertyToKind<alignment_key> {
-  static constexpr PropKind Kind = PropKind::Alignment;
-};
-
-template <> struct IsCompileTimeProperty<alignment_key> : std::true_type {};
-
 template <int N> struct PropertyMetaInfo<alignment_key::value_t<N>> {
   static constexpr const char *name = "sycl-alignment";
   static constexpr int value = N;
diff --git a/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp
index f21099737b97f..9b6fd04a7e5a9 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp
@@ -27,21 +27,13 @@ struct max_num_work_group_sync {
 };
 } // namespace info::kernel_queue_specific
 
-struct use_root_sync_key {
+struct use_root_sync_key
+    : detail::compile_time_property_key<detail::PropKind::UseRootSync> {
   using value_t = property_value<use_root_sync_key>;
 };
 
 inline constexpr use_root_sync_key::value_t use_root_sync;
 
-template <> struct is_property_key<use_root_sync_key> : std::true_type {};
-
-template <> struct detail::PropertyToKind<use_root_sync_key> {
-  static constexpr PropKind Kind = PropKind::UseRootSync;
-};
-
-template <>
-struct detail::IsCompileTimeProperty<use_root_sync_key> : std::true_type {};
-
 template <int Dimensions> class root_group {
 public:
   using id_type = id<Dimensions>;
diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
index dc53678d2618c..c3e2340af718d 100644
--- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
@@ -34,25 +34,28 @@ template <size_t X, size_t... Xs> struct AllNonZero<X, Xs...> {
 
 struct properties_tag {};
 
-struct work_group_size_key {
+struct work_group_size_key
+    : detail::compile_time_property_key<detail::PropKind::WorkGroupSize> {
   template <size_t... Dims>
   using value_t = property_value<work_group_size_key,
                                  std::integral_constant<size_t, Dims>...>;
 };
 
-struct work_group_size_hint_key {
+struct work_group_size_hint_key
+    : detail::compile_time_property_key<detail::PropKind::WorkGroupSizeHint> {
   template <size_t... Dims>
   using value_t = property_value<work_group_size_hint_key,
                                  std::integral_constant<size_t, Dims>...>;
 };
 
-struct sub_group_size_key {
+struct sub_group_size_key
+    : detail::compile_time_property_key<detail::PropKind::SubGroupSize> {
   template <uint32_t Size>
   using value_t = property_value<sub_group_size_key,
                                  std::integral_constant<uint32_t, Size>>;
 };
 
-struct device_has_key {
+struct device_has_key : detail::compile_time_property_key<detail::PropKind::DeviceHas> {
   template <aspect... Aspects>
   using value_t = property_value<device_has_key,
                                  std::integral_constant<aspect, Aspects>...>;
@@ -123,34 +126,7 @@ inline constexpr sub_group_size_key::value_t<Size> sub_group_size;
 template <aspect... Aspects>
 inline constexpr device_has_key::value_t<Aspects...> device_has;
 
-template <> struct is_property_key<work_group_size_key> : std::true_type {};
-template <>
-struct is_property_key<work_group_size_hint_key> : std::true_type {};
-template <> struct is_property_key<sub_group_size_key> : std::true_type {};
-template <> struct is_property_key<device_has_key> : std::true_type {};
-
 namespace detail {
-template <> struct PropertyToKind<work_group_size_key> {
-  static constexpr PropKind Kind = PropKind::WorkGroupSize;
-};
-template <> struct PropertyToKind<work_group_size_hint_key> {
-  static constexpr PropKind Kind = PropKind::WorkGroupSizeHint;
-};
-template <> struct PropertyToKind<sub_group_size_key> {
-  static constexpr PropKind Kind = PropKind::SubGroupSize;
-};
-template <> struct PropertyToKind<device_has_key> {
-  static constexpr PropKind Kind = PropKind::DeviceHas;
-};
-
-template <>
-struct IsCompileTimeProperty<work_group_size_key> : std::true_type {};
-template <>
-struct IsCompileTimeProperty<work_group_size_hint_key> : std::true_type {};
-template <>
-struct IsCompileTimeProperty<sub_group_size_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<device_has_key> : std::true_type {};
-
 template <size_t Dim0, size_t... Dims>
 struct PropertyMetaInfo<work_group_size_key::value_t<Dim0, Dims...>> {
   static constexpr const char *name = "sycl-work-group-size";
diff --git a/sycl/include/sycl/ext/oneapi/latency_control/properties.hpp b/sycl/include/sycl/ext/oneapi/latency_control/properties.hpp
index 630a2ae0ae248..0716720a3ab36 100644
--- a/sycl/include/sycl/ext/oneapi/latency_control/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/latency_control/properties.hpp
@@ -24,14 +24,18 @@ enum class latency_control_type {
   min
 };
 
-struct latency_anchor_id_key {
+struct latency_anchor_id_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::LatencyAnchorID> {
   template <int Anchor>
   using value_t =
       oneapi::experimental::property_value<latency_anchor_id_key,
                                            std::integral_constant<int, Anchor>>;
 };
 
-struct latency_constraint_key {
+struct latency_constraint_key
+    : oneapi::experimental::detail::compile_time_property_key<
+          oneapi::experimental::detail::PropKind::LatencyConstraint> {
   template <int Target, latency_control_type Type, int Cycle>
   using value_t = oneapi::experimental::property_value<
       latency_constraint_key, std::integral_constant<int, Target>,
@@ -61,30 +65,6 @@ struct property_value<
   static constexpr int cycle = Cycle;
 };
 
-template <>
-struct is_property_key<intel::experimental::latency_anchor_id_key>
-    : std::true_type {};
-template <>
-struct is_property_key<intel::experimental::latency_constraint_key>
-    : std::true_type {};
-
-namespace detail {
-
-template <> struct PropertyToKind<intel::experimental::latency_anchor_id_key> {
-  static constexpr PropKind Kind = PropKind::LatencyAnchorID;
-};
-template <> struct PropertyToKind<intel::experimental::latency_constraint_key> {
-  static constexpr PropKind Kind = PropKind::LatencyConstraint;
-};
-
-template <>
-struct IsCompileTimeProperty<intel::experimental::latency_anchor_id_key>
-    : std::true_type {};
-template <>
-struct IsCompileTimeProperty<intel::experimental::latency_constraint_key>
-    : std::true_type {};
-
-} // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index 72cd83b7c435a..6748e2c8a3f23 100644
--- a/sycl/include/sycl/ext/oneapi/properties/property.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -54,7 +54,7 @@ enum PropKind : uint32_t {
 namespace sycl::ext::oneapi::experimental {
 
 // (2.)
-struct bar_key {
+struct bar_key : detail::compile_time_property_key<PropKind::Bar> {
   using value_t = property_value<bar_key>;
 };
 
@@ -62,21 +62,11 @@ struct bar_key {
 inline constexpr bar_key::value_t bar;
 
 // (4.)
-template <> struct is_property_key<bar_key> : std::true_type {};
 // Replace SYCL_OBJ with the SYCL object to support the property.
 template <> struct is_property_key_of<bar_key, SYCL_OBJ> : std::true_type {};
 
 namespace detail {
-
 // (5.)
-template <> struct PropertyToKind<bar_key> {
-  static constexpr PropKind Kind = PropKind::Bar;
-};
-
-// (6.)
-template <> struct IsCompileTimeProperty<bar_key> : std::true_type {};
-
-// (7.)
 template <> struct PropertyMetaInfo<bar_key::value_t> {
   static constexpr const char *name = "sycl-bar";
   static constexpr int value = 5;
@@ -90,21 +80,14 @@ template <> struct PropertyMetaInfo<bar_key::value_t> {
 //  1. Add a new enumerator to `sycl::ext::oneapi::detail::PropKind`
 //     representing the new property. Increment
 //     `sycl::ext::oneapi::experimental::detail::PropKind::PropKindSize`
-//  2. Define property class.
+//  2. Define property class, inheriting from `detail::run_time_property_key`.
 //  3. Declare the property key as an alias to the property class. The name of
 //     the key class must be the property name suffixed by `_key`, i.e. for a
 //     property `foo` the class should be named `foo_key`.
 //  4. Overload the `==` and `!=` operators for the new property class. The
 //     comparison should compare all data members of the property class.
-//  5. Specialize `sycl::ext::oneapi::experimental::is_property_key` and
-//     `sycl::ext::oneapi::experimental::is_property_key_of` for the property
-//     class.
-//  6. Specialize `sycl::ext::oneapi::detail::PropertyToKind` for the new
-//     property class. The specialization should have a `Kind` member with the
-//     value equal to the enumerator added in 1.
-//  7. Specialize `sycl::ext::oneapi::experimental::detail::IsRuntimeProperty`
-//     for the new property class. This specialization should derive from
-//     `std::true_type`.
+//  5. Specialize `sycl::ext::oneapi::experimental::is_property_key_of` for the
+//     property class.
 /******************************* EXAMPLE ***************************************
 ------------- sycl/include/sycl/ext/oneapi/properties/property.hpp -------------
 // (1.)
@@ -134,21 +117,9 @@ inline bool operator!=(const foo &lhs, const foo &rhs) {
 }
 
 // (5.)
-template <> struct is_property_key<foo> : std::true_type {};
 // Replace SYCL_OBJ with the SYCL object to support the property.
 template <> struct is_property_key_of<foo, SYCL_OBJ> : std::true_type {};
 
-namespace detail {
-
-// (6.)
-template <> struct PropertyToKind<foo> {
-  static constexpr PropKind Kind = PropKind::Foo;
-};
-
-// (7.)
-template <> struct IsRuntimeProperty<foo> : std::true_type {};
-
-} // namespace detail
 } // namespace sycl::ext::oneapi::experimental
 *******************************************************************************/
 
@@ -225,9 +196,23 @@ enum PropKind : uint32_t {
   PropKindSize = 54,
 };
 
+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 {
+  static constexpr PropKind Kind = Kind_;
+};
+
+template <PropKind Kind_>
+struct compile_time_property_key : compile_time_property_key_base_tag {
+  static constexpr PropKind Kind = Kind_;
+};
+
 // This trait must be specialized for all properties and must have a unique
 // constexpr PropKind member named Kind.
-template <typename PropertyT> struct PropertyToKind {};
+template <typename PropertyT> struct PropertyToKind {
+  static constexpr PropKind Kind = PropertyT::Kind;
+};
 
 // Get unique ID for property.
 template <typename PropertyT> struct PropertyID {
@@ -236,10 +221,18 @@ template <typename PropertyT> struct PropertyID {
 };
 
 // Trait for identifying runtime properties.
-template <typename PropertyT> struct IsRuntimeProperty : std::false_type {};
+template <typename PropertyT>
+struct IsRuntimeProperty
+    : std::bool_constant<
+          std::is_base_of_v<property_key_base_tag, PropertyT> &&
+          !std::is_base_of_v<compile_time_property_key_base_tag, PropertyT>> {};
 
 // Trait for identifying compile-time properties.
-template <typename PropertyT> struct IsCompileTimeProperty : std::false_type {};
+template <typename PropertyT>
+struct IsCompileTimeProperty
+    : std::bool_constant<
+          std::is_base_of_v<property_key_base_tag, PropertyT> &&
+          std::is_base_of_v<compile_time_property_key_base_tag, PropertyT>> {};
 
 // Trait for property compile-time meta names and values.
 template <typename PropertyT> struct PropertyMetaInfo {
@@ -251,7 +244,10 @@ template <typename PropertyT> struct PropertyMetaInfo {
 
 } // namespace detail
 
-template <typename> struct is_property_key : std::false_type {};
+template <typename T>
+struct is_property_key
+    : std::bool_constant<std::is_base_of_v<detail::property_key_base_tag, T>> {
+};
 template <typename, typename> struct is_property_key_of : std::false_type {};
 
 } // namespace experimental
diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index 91caddfec75df..d48cf27b3d531 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -816,70 +816,33 @@ namespace ext::oneapi::experimental {
 /////////////////////////
 // PropertyT syclex::build_options
 /////////////////////////
-struct build_options {
+struct build_options
+    : detail::run_time_property_key<detail::PropKind::BuildOptions> {
   std::vector<std::string> opts;
   build_options(const std::string &optsArg) : opts{optsArg} {}
   build_options(const std::vector<std::string> &optsArg) : opts(optsArg) {}
 };
 using build_options_key = build_options;
 
-template <> struct is_property_key<build_options_key> : std::true_type {};
-
 template <>
 struct is_property_key_of<build_options_key,
                           sycl::kernel_bundle<bundle_state::ext_oneapi_source>>
     : std::true_type {};
 
-namespace detail {
-
-template <>
-struct PropertyToKind<sycl::ext::oneapi::experimental::build_options_key> {
-  static constexpr PropKind Kind = PropKind::BuildOptions;
-};
-
-template <>
-struct IsRuntimeProperty<sycl::ext::oneapi::experimental::build_options_key>
-    : std::true_type {};
-
-template <>
-struct IsCompileTimeProperty<sycl::ext::oneapi::experimental::build_options_key>
-    : std::false_type {};
-
-} // namespace detail
-
 /////////////////////////
 // PropertyT syclex::save_log
 /////////////////////////
-struct save_log {
+struct save_log : detail::run_time_property_key<detail::PropKind::BuildLog> {
   std::string *log;
   save_log(std::string *logArg) : log(logArg) {}
 };
 using save_log_key = save_log;
 
-template <> struct is_property_key<save_log_key> : std::true_type {};
-
 template <>
 struct is_property_key_of<save_log_key,
                           sycl::kernel_bundle<bundle_state::ext_oneapi_source>>
     : std::true_type {};
 
-namespace detail {
-
-template <>
-struct PropertyToKind<sycl::ext::oneapi::experimental::save_log_key> {
-  static constexpr PropKind Kind = PropKind::BuildLog;
-};
-
-template <>
-struct IsRuntimeProperty<sycl::ext::oneapi::experimental::save_log_key>
-    : std::true_type {};
-
-template <>
-struct IsCompileTimeProperty<sycl::ext::oneapi::experimental::save_log_key>
-    : std::false_type {};
-
-} // namespace detail
-
 /////////////////////////
 // syclex::is_source_kernel_bundle_supported
 /////////////////////////
diff --git a/sycl/test/extensions/properties/mock_compile_time_properties.hpp b/sycl/test/extensions/properties/mock_compile_time_properties.hpp
index 5e3bfed70e1d9..7af3615038d4b 100644
--- a/sycl/test/extensions/properties/mock_compile_time_properties.hpp
+++ b/sycl/test/extensions/properties/mock_compile_time_properties.hpp
@@ -17,20 +17,28 @@ namespace ext {
 namespace oneapi {
 namespace experimental {
 
-struct bar_key {
+struct bar_key
+    : detail::compile_time_property_key<static_cast<detail::PropKind>(
+          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 0)> {
   using value_t = property_value<bar_key>;
 };
 
-struct baz_key {
+struct baz_key
+    : detail::compile_time_property_key<static_cast<detail::PropKind>(
+          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 1)> {
   template <int K>
   using value_t = property_value<baz_key, std::integral_constant<int, K>>;
 };
 
-struct boo_key {
+struct boo_key
+    : detail::compile_time_property_key<static_cast<detail::PropKind>(
+          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 2)> {
   template <typename... Ts> using value_t = property_value<boo_key, Ts...>;
 };
 
-struct foo {
+struct foo
+    : detail::run_time_property_key<static_cast<detail::PropKind>(
+          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 3)> {
   constexpr foo(int v) : value(v) {}
   int value;
 };
@@ -40,7 +48,9 @@ inline bool operator==(const foo &lhs, const foo &rhs) {
 }
 inline bool operator!=(const foo &lhs, const foo &rhs) { return !(lhs == rhs); }
 
-struct foz {
+struct foz
+    : detail::run_time_property_key<static_cast<detail::PropKind>(
+          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 4)> {
   constexpr foz(float v1, bool v2) : value1(v1), value2(v2) {}
   // Define copy constructor to make foz non-trivially copyable
   constexpr foz(const foz &f) {
@@ -56,7 +66,9 @@ inline bool operator==(const foz &lhs, const foz &rhs) {
 }
 inline bool operator!=(const foz &lhs, const foz &rhs) { return !(lhs == rhs); }
 
-struct fir {
+struct fir
+    : detail::run_time_property_key<static_cast<detail::PropKind>(
+          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 5)> {
   // Intentionally not constexpr to test for properties that cannot be constexpr
   fir(float v1, bool v2) : value1(v1), value2(v2) {}
   // Define copy constructor to make foz non-trivially copyable
@@ -81,13 +93,6 @@ using foo_key = foo;
 using foz_key = foz;
 using fir_key = fir;
 
-template <> struct is_property_key<bar_key> : std::true_type {};
-template <> struct is_property_key<baz_key> : std::true_type {};
-template <> struct is_property_key<boo_key> : std::true_type {};
-template <> struct is_property_key<foo_key> : std::true_type {};
-template <> struct is_property_key<foz_key> : std::true_type {};
-template <> struct is_property_key<fir_key> : std::true_type {};
-
 template <typename syclObjectT>
 struct is_property_key_of<bar_key, syclObjectT> : std::true_type {};
 template <typename syclObjectT>
@@ -102,39 +107,6 @@ template <typename syclObjectT>
 struct is_property_key_of<fir_key, syclObjectT> : std::true_type {};
 
 namespace detail {
-template <> struct PropertyToKind<bar_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 0);
-};
-template <> struct PropertyToKind<baz_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 1);
-};
-template <> struct PropertyToKind<foo_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 2);
-};
-template <> struct PropertyToKind<boo_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 3);
-};
-template <> struct PropertyToKind<foz_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4);
-};
-template <> struct PropertyToKind<fir_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4);
-};
-
-template <> struct IsCompileTimeProperty<bar_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<baz_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<boo_key> : std::true_type {};
-
-template <> struct IsRuntimeProperty<foo_key> : std::true_type {};
-template <> struct IsRuntimeProperty<foz_key> : std::true_type {};
-template <> struct IsRuntimeProperty<fir_key> : std::true_type {};
-
 template <typename Properties>
 struct ConflictingProperties<boo_key, Properties>
     : ContainsProperty<fir_key, Properties> {};

From 5ab3e4ffddc395e17affe5ffbde1c88c1917bd2a Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Tue, 27 Feb 2024 14:44:13 -0800
Subject: [PATCH 2/5] Update the example

---
 sycl/include/sycl/ext/oneapi/properties/property.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index 6748e2c8a3f23..8067f7c52c9b3 100644
--- a/sycl/include/sycl/ext/oneapi/properties/property.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -100,7 +100,7 @@ enum PropKind : uint32_t {
 namespace sycl::ext::oneapi::experimental {
 
 // (2.)
-struct foo {
+struct foo : detail::run_time_property_key<PropKind::Foo> {
   foo(int v) : value(v) {}
   int value;
 };

From 1aa7447d8c65eec9e799fb78bf2b1995483c787a Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Tue, 27 Feb 2024 15:08:21 -0800
Subject: [PATCH 3/5] Make Kind protected

---
 sycl/include/sycl/ext/oneapi/properties/property.hpp | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index 8067f7c52c9b3..586d7baf4eba0 100644
--- a/sycl/include/sycl/ext/oneapi/properties/property.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -200,12 +200,20 @@ 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 {
+protected:
   static constexpr PropKind Kind = Kind_;
+
+  template <typename T>
+  friend struct PropertyToKind;
 };
 
 template <PropKind Kind_>
 struct compile_time_property_key : compile_time_property_key_base_tag {
+protected:
   static constexpr PropKind Kind = Kind_;
+
+  template <typename T>
+  friend struct PropertyToKind;
 };
 
 // This trait must be specialized for all properties and must have a unique

From 65c6e14d2d304b962466fab2bc3505ecce918915 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Tue, 27 Feb 2024 15:17:18 -0800
Subject: [PATCH 4/5] Update the last remaining test

---
 .../annotated_usm/fake_properties.hpp         | 295 +++---------------
 1 file changed, 46 insertions(+), 249 deletions(-)

diff --git a/sycl/test/extensions/annotated_usm/fake_properties.hpp b/sycl/test/extensions/annotated_usm/fake_properties.hpp
index 06a026e48c9b5..656ea1a1ad75e 100644
--- a/sycl/test/extensions/annotated_usm/fake_properties.hpp
+++ b/sycl/test/extensions/annotated_usm/fake_properties.hpp
@@ -17,17 +17,22 @@ namespace ext {
 namespace oneapi {
 namespace experimental {
 
+inline constexpr detail::PropKind fakePropKind(int N) {
+  return static_cast<detail::PropKind>(
+      static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + N);
+}
+
 // Compile-time properties
-struct bar_key {
+struct bar_key : detail::compile_time_property_key<fakePropKind(0)> {
   using value_t = property_value<bar_key>;
 };
 
-struct baz_key {
+struct baz_key : detail::compile_time_property_key<fakePropKind(1)> {
   template <int K>
   using value_t = property_value<baz_key, std::integral_constant<int, K>>;
 };
 
-struct boo_key {
+struct boo_key : detail::compile_time_property_key<fakePropKind(2)> {
   template <typename... Ts> using value_t = property_value<boo_key, Ts...>;
 };
 
@@ -35,10 +40,6 @@ inline constexpr bar_key::value_t bar;
 template <int K> inline constexpr baz_key::value_t<K> baz;
 template <typename... Ts> inline constexpr boo_key::value_t<Ts...> boo;
 
-template <> struct is_property_key<bar_key> : std::true_type {};
-template <> struct is_property_key<baz_key> : std::true_type {};
-template <> struct is_property_key<boo_key> : std::true_type {};
-
 template <typename objectT>
 struct is_property_key_of<bar_key, objectT> : std::true_type {};
 template <typename objectT>
@@ -46,72 +47,53 @@ struct is_property_key_of<baz_key, objectT> : std::true_type {};
 template <typename objectT>
 struct is_property_key_of<boo_key, objectT> : std::true_type {};
 
-namespace detail {
-template <> struct PropertyToKind<bar_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 0);
-};
-template <> struct PropertyToKind<baz_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 1);
-};
-
-template <> struct PropertyToKind<boo_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 2);
-};
-
-template <> struct IsCompileTimeProperty<bar_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<baz_key> : std::true_type {};
-template <> struct IsCompileTimeProperty<boo_key> : std::true_type {};
-
-} // namespace detail
-
 // Runtime properties
 enum foo_enum : unsigned { a, b, c };
-struct foo {
+struct foo : detail::run_time_property_key<fakePropKind(3)> {
   constexpr foo(foo_enum v) : value(v) {}
   foo_enum value;
 };
 
-struct foz {
+struct foz  : detail::run_time_property_key<fakePropKind(4)> {
   float value1;
   bool value2;
-};
 
-struct rt_prop1 {};
-struct rt_prop2 {};
-struct rt_prop3 {};
-struct rt_prop4 {};
-struct rt_prop5 {};
-struct rt_prop6 {};
-struct rt_prop7 {};
-struct rt_prop8 {};
-struct rt_prop9 {};
-struct rt_prop10 {};
-struct rt_prop11 {};
-struct rt_prop12 {};
-struct rt_prop13 {};
-struct rt_prop14 {};
-struct rt_prop15 {};
-struct rt_prop16 {};
-struct rt_prop17 {};
-struct rt_prop18 {};
-struct rt_prop19 {};
-struct rt_prop20 {};
-struct rt_prop21 {};
-struct rt_prop22 {};
-struct rt_prop23 {};
-struct rt_prop24 {};
-struct rt_prop25 {};
-struct rt_prop26 {};
-struct rt_prop27 {};
-struct rt_prop28 {};
-struct rt_prop29 {};
-struct rt_prop30 {};
-struct rt_prop31 {};
-struct rt_prop32 {};
-struct rt_prop33 {};
+  foz(float value1, bool value2) : value1(value1), value2(value2) {}
+};
+
+struct rt_prop1 : detail::run_time_property_key<fakePropKind(5)> {};
+struct rt_prop2 : detail::run_time_property_key<fakePropKind(6)> {};
+struct rt_prop3 : detail::run_time_property_key<fakePropKind(7)> {};
+struct rt_prop4 : detail::run_time_property_key<fakePropKind(8)> {};
+struct rt_prop5 : detail::run_time_property_key<fakePropKind(9)> {};
+struct rt_prop6 : detail::run_time_property_key<fakePropKind(10)> {};
+struct rt_prop7 : detail::run_time_property_key<fakePropKind(11)> {};
+struct rt_prop8 : detail::run_time_property_key<fakePropKind(12)> {};
+struct rt_prop9 : detail::run_time_property_key<fakePropKind(13)> {};
+struct rt_prop10 : detail::run_time_property_key<fakePropKind(14)> {};
+struct rt_prop11 : detail::run_time_property_key<fakePropKind(15)> {};
+struct rt_prop12 : detail::run_time_property_key<fakePropKind(16)> {};
+struct rt_prop13 : detail::run_time_property_key<fakePropKind(17)> {};
+struct rt_prop14 : detail::run_time_property_key<fakePropKind(18)> {};
+struct rt_prop15 : detail::run_time_property_key<fakePropKind(19)> {};
+struct rt_prop16 : detail::run_time_property_key<fakePropKind(20)> {};
+struct rt_prop17 : detail::run_time_property_key<fakePropKind(21)> {};
+struct rt_prop18 : detail::run_time_property_key<fakePropKind(22)> {};
+struct rt_prop19 : detail::run_time_property_key<fakePropKind(23)> {};
+struct rt_prop20 : detail::run_time_property_key<fakePropKind(24)> {};
+struct rt_prop21 : detail::run_time_property_key<fakePropKind(25)> {};
+struct rt_prop22 : detail::run_time_property_key<fakePropKind(26)> {};
+struct rt_prop23 : detail::run_time_property_key<fakePropKind(27)> {};
+struct rt_prop24 : detail::run_time_property_key<fakePropKind(28)> {};
+struct rt_prop25 : detail::run_time_property_key<fakePropKind(29)> {};
+struct rt_prop26 : detail::run_time_property_key<fakePropKind(30)> {};
+struct rt_prop27 : detail::run_time_property_key<fakePropKind(31)> {};
+struct rt_prop28 : detail::run_time_property_key<fakePropKind(32)> {};
+struct rt_prop29 : detail::run_time_property_key<fakePropKind(33)> {};
+struct rt_prop30 : detail::run_time_property_key<fakePropKind(34)> {};
+struct rt_prop31 : detail::run_time_property_key<fakePropKind(35)> {};
+struct rt_prop32 : detail::run_time_property_key<fakePropKind(36)> {};
+struct rt_prop33 : detail::run_time_property_key<fakePropKind(37)> {};
 
 using foo_key = foo;
 using foz_key = foz;
@@ -149,9 +131,6 @@ using rt_prop31_key = rt_prop31;
 using rt_prop32_key = rt_prop32;
 using rt_prop33_key = rt_prop33;
 
-template <> struct is_property_key<foo_key> : std::true_type {};
-template <> struct is_property_key<foz_key> : std::true_type {};
-
 template <> struct is_property_key<rt_prop1_key> : std::true_type {};
 template <> struct is_property_key<rt_prop2_key> : std::true_type {};
 template <> struct is_property_key<rt_prop3_key> : std::true_type {};
@@ -258,188 +237,6 @@ struct is_property_key_of<rt_prop32_key, objectT> : std::true_type {};
 template <typename objectT>
 struct is_property_key_of<rt_prop33_key, objectT> : std::true_type {};
 
-namespace detail {
-template <> struct PropertyToKind<foo_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 3);
-};
-
-template <> struct PropertyToKind<foz_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4);
-};
-
-template <> struct PropertyToKind<rt_prop1_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 1);
-};
-template <> struct PropertyToKind<rt_prop2_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 2);
-};
-template <> struct PropertyToKind<rt_prop3_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 3);
-};
-template <> struct PropertyToKind<rt_prop4_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 4);
-};
-template <> struct PropertyToKind<rt_prop5_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 5);
-};
-template <> struct PropertyToKind<rt_prop6_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 6);
-};
-template <> struct PropertyToKind<rt_prop7_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 7);
-};
-template <> struct PropertyToKind<rt_prop8_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 8);
-};
-template <> struct PropertyToKind<rt_prop9_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 9);
-};
-template <> struct PropertyToKind<rt_prop10_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 10);
-};
-template <> struct PropertyToKind<rt_prop11_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 11);
-};
-template <> struct PropertyToKind<rt_prop12_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 12);
-};
-template <> struct PropertyToKind<rt_prop13_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 13);
-};
-template <> struct PropertyToKind<rt_prop14_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 14);
-};
-template <> struct PropertyToKind<rt_prop15_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 15);
-};
-template <> struct PropertyToKind<rt_prop16_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 16);
-};
-template <> struct PropertyToKind<rt_prop17_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 17);
-};
-template <> struct PropertyToKind<rt_prop18_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 18);
-};
-template <> struct PropertyToKind<rt_prop19_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 19);
-};
-template <> struct PropertyToKind<rt_prop20_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 20);
-};
-template <> struct PropertyToKind<rt_prop21_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 21);
-};
-template <> struct PropertyToKind<rt_prop22_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 22);
-};
-template <> struct PropertyToKind<rt_prop23_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 23);
-};
-template <> struct PropertyToKind<rt_prop24_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 24);
-};
-template <> struct PropertyToKind<rt_prop25_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 25);
-};
-template <> struct PropertyToKind<rt_prop26_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 26);
-};
-template <> struct PropertyToKind<rt_prop27_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 27);
-};
-template <> struct PropertyToKind<rt_prop28_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 28);
-};
-template <> struct PropertyToKind<rt_prop29_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 29);
-};
-template <> struct PropertyToKind<rt_prop30_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 30);
-};
-template <> struct PropertyToKind<rt_prop31_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 31);
-};
-template <> struct PropertyToKind<rt_prop32_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 32);
-};
-template <> struct PropertyToKind<rt_prop33_key> {
-  static constexpr PropKind Kind = static_cast<enum PropKind>(
-      static_cast<std::uint32_t>(PropKind::PropKindSize) + 4 + 33);
-};
-
-template <> struct IsRuntimeProperty<foo_key> : std::true_type {};
-template <> struct IsRuntimeProperty<foz_key> : std::true_type {};
-
-template <> struct IsRuntimeProperty<rt_prop1_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop2_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop3_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop4_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop5_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop6_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop7_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop8_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop9_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop10_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop11_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop12_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop13_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop14_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop15_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop16_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop17_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop18_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop19_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop20_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop21_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop22_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop23_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop24_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop25_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop26_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop27_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop28_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop29_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop30_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop31_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop32_key> : std::true_type {};
-template <> struct IsRuntimeProperty<rt_prop33_key> : std::true_type {};
-
-} // namespace detail
 } // namespace experimental
 } // namespace oneapi
 } // namespace ext

From 1995fc56e18d43e03c34dc769c185bed0a52de95 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Tue, 27 Feb 2024 15:19:27 -0800
Subject: [PATCH 5/5] Simplify mock_compile_time_properties.hpp by fakePropKind
 helper

---
 .../mock_compile_time_properties.hpp          | 28 +++++++------------
 1 file changed, 10 insertions(+), 18 deletions(-)

diff --git a/sycl/test/extensions/properties/mock_compile_time_properties.hpp b/sycl/test/extensions/properties/mock_compile_time_properties.hpp
index 7af3615038d4b..5b64951bae24c 100644
--- a/sycl/test/extensions/properties/mock_compile_time_properties.hpp
+++ b/sycl/test/extensions/properties/mock_compile_time_properties.hpp
@@ -16,29 +16,25 @@ namespace sycl {
 namespace ext {
 namespace oneapi {
 namespace experimental {
+inline constexpr detail::PropKind fakePropKind(int N) {
+  return static_cast<detail::PropKind>(
+      static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + N);
+}
 
-struct bar_key
-    : detail::compile_time_property_key<static_cast<detail::PropKind>(
-          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 0)> {
+struct bar_key : detail::compile_time_property_key<fakePropKind(0)> {
   using value_t = property_value<bar_key>;
 };
 
-struct baz_key
-    : detail::compile_time_property_key<static_cast<detail::PropKind>(
-          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 1)> {
+struct baz_key : detail::compile_time_property_key<fakePropKind(1)> {
   template <int K>
   using value_t = property_value<baz_key, std::integral_constant<int, K>>;
 };
 
-struct boo_key
-    : detail::compile_time_property_key<static_cast<detail::PropKind>(
-          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 2)> {
+struct boo_key : detail::compile_time_property_key<fakePropKind(2)> {
   template <typename... Ts> using value_t = property_value<boo_key, Ts...>;
 };
 
-struct foo
-    : detail::run_time_property_key<static_cast<detail::PropKind>(
-          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 3)> {
+struct foo : detail::run_time_property_key<fakePropKind(3)> {
   constexpr foo(int v) : value(v) {}
   int value;
 };
@@ -48,9 +44,7 @@ inline bool operator==(const foo &lhs, const foo &rhs) {
 }
 inline bool operator!=(const foo &lhs, const foo &rhs) { return !(lhs == rhs); }
 
-struct foz
-    : detail::run_time_property_key<static_cast<detail::PropKind>(
-          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 4)> {
+struct foz : detail::run_time_property_key<fakePropKind(4)> {
   constexpr foz(float v1, bool v2) : value1(v1), value2(v2) {}
   // Define copy constructor to make foz non-trivially copyable
   constexpr foz(const foz &f) {
@@ -66,9 +60,7 @@ inline bool operator==(const foz &lhs, const foz &rhs) {
 }
 inline bool operator!=(const foz &lhs, const foz &rhs) { return !(lhs == rhs); }
 
-struct fir
-    : detail::run_time_property_key<static_cast<detail::PropKind>(
-          static_cast<std::uint32_t>(detail::PropKind::PropKindSize) + 5)> {
+struct fir : detail::run_time_property_key<fakePropKind(5)> {
   // Intentionally not constexpr to test for properties that cannot be constexpr
   fir(float v1, bool v2) : value1(v1), value2(v2) {}
   // Define copy constructor to make foz non-trivially copyable